
    Wh                    h   d dl mZ d dlm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 ]&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ddZ eeed          Z eeed          ZdS )    )annotations)partialN)
xla_client   )custom_call)z.cudajax_cuda12_pluginz._prngjaxlib)package_ffiCUDA)platformapi_version)z.rocmjax_rocm60_pluginROCMFr   strlengthint | ir.Value | Noneoutput_shapeir.Value | Noneforward_compatibility_modeboolc           
        ~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}i }t          |t                    rd}n|J ||g}| d}t          |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_threefry2x32_ffi)r   result_typesoperandsbackend_configoperand_layoutsresult_layoutsresult_shapes)lenirRankedTensorTypetypeelement_typeIntegerTypeget_unsignedshape	itertoolschaintuplerange
isinstanceintr   results)prngr   keysdatar   r   r   typdimsxndimslayoutr!   r   opaquer#   custom_call_targets                    U/var/www/html/movieo_spanner_bot/venv/lib/python3.11/site-packages/jaxlib/gpu_prng.py_threefry2x32_loweringr>   >   s    !	Ta	Ta

d1gl
+
+
8
.
%
%b
)
)* * *,0GL* * * 	Q#		S	!	!	'$?4&& ( (a6S===163-====
d))%uqy"b))**&HqL/1gtAwQa1(& 1MM###!<0M"555	:%X\!
# 
# 
# $++    cuhip)NNF)r   r   r   r   r   r   r   r   )
__future__r   	functoolsr   	importlibr,   jaxlib.mlir.irmlirr%   r	   r   hlo_helpersr   cuda_module_nameimport_module
_cuda_prngImportErrorregistrationsitems_name_valuer   register_custom_call_targetrocm_module_name	_hip_prngr>   cuda_threefry2x32rocm_threefry2x32 r?   r=   <module>rV      s   # " " " " "                              $ $ $ $ $ $ 7 
 

((###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
 <@;?>C*+ *+ *+ *+ *+Z G2JEE G2IuEE   s#   AAA"B99CC