
    Spf                       d dl mZ d dlZd dlmZ d dlZd dlZd dlZd dlmc m	Z	 d dl
mZ ddlmZ ddlmZ dD ]&Z	  ej        e d	d
          Z n# e$ r dZY #w xY werHe                                                                D ]!\  ZZdev rdnd Z ej        eede           "dD ]&Z	  ej        e d	d
          Z n# e$ r dZY #w xY werHe                                                                D ]!\  ZZdev rdnd Z ej        eede           "d Z	 	 	 dddZ eeed          Z eeed          ZdS )     )annotationsN)partial)
xla_client   )custom_call)GpuLibNotLinkedError)z.cudajax_cuda12_pluginz._prngjaxlib)package_ffiCUDA)platformapi_version)z.rocmjax_rocm60_pluginROCMc                B    t          j        t          j        | d          S )Nr   )	functoolsreduceoperatormul)xss    O/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jaxlib/gpu_prng.py<lambda>r   @   s    9#HL"a88     Tr   strlengthint | ir.Value | Noneoutput_shapeir.Value | Noneforward_compatibility_modeboolc           
     P   |r| st                      t          |          dk    s
J |            t          |          dk    s
J |            t          j        |d         j                  j        t          j                            d          k    sJ |d         j                    |d         j        }t          j        |          j        }t          j
        ||          D ]}	|	j        |k    sJ |	j        |f            t          |          }
t          t          |
dz
  dd                    }|gdz  }|d         |d         |d         |d         g}|r|t          |          }i }t          |t                    r|r|                     |          }d}n|J |r|                     d          }t          j        |j                  j        t          j                            d          k    s
J |            t          j        |j                  j        dgk    s(J |t          j        |j                  j        f            |                    |           |                    d	           ||g}|r| d
n| d}t'          ||rdnd||g||||gdz  |          j        S )zThreeFry2x32 kernel for GPU.

  In presence of dynamic shapes, `length` is an `ir.Value` and `output_shape`
  is a 1D tensor describing the shape of the two outputs.
     r       r      N@   )r   _threefry2x32_threefry2x32_ffi)r   result_typesoperandsbackend_configoperand_layoutsresult_layoutsresult_shapes)r   lenirRankedTensorTypetypeelement_typeIntegerTypeget_unsignedshape	itertoolschaintuplerange_prod
isinstanceintthreefry2x32_descriptorget_signlessappendr   results)prngr   keysdatar   r   r    typdimsxndimslayoutr-   r+   opaquer/   custom_call_targets                    r   _threefry2x32_loweringrM   C   s      ! !

 
  	Ta	Ta

d1gl
+
+
8
.
%
%b
)
)* * *,0GL* * * 	Q#		S	!	!	'$?4&& ( (a6S===163-====
d))%uqy"b))**&HqL/1gtAwQa1( FN4[[F& 1! 4++F33fMM###! 	#++B//f!&+..;n))"--. . .06. . .!&+..4c  R0==CD   oofT"""!<0M 
$*    ))) 
 
2911:%X\!
# 
# 
# $++r   cuhip)NNT)r   r   r   r   r   r   r    r!   ) 
__future__r   r   r   	importlibr8   r   jaxlib.mlir.irmlirr1   r
   r   hlo_helpersr   gpu_common_utilsr   cuda_module_nameimport_module
_cuda_prngImportErrorregistrationsitems_name_valuer   register_custom_call_targetrocm_module_name	_hip_prngr<   rM   cuda_threefry2x32rocm_threefry2x32 r   r   <module>rd      s   # " " " " "                                   $ $ $ $ $ $ 2 2 2 2 2 26 
 

((###X  J 
E 
   JJJ
  D!//117799 D DmeV!!AK*J*5&67BD D D D D 7 
 

'	'###X  I 
E 
   III
  D ..006688 D DmeV!!AK*J*5&67BD D D D D 	98 <@;?>B>+ >+ >+ >+ >+B G2JEE G2IuEE   s#   AAA0CCC