
    9i]                       S SK Jr  S SKrS SKrS SKrS SKJrJrJrJ	r	  S SK
r
S SKrS SKrS SKJr  S SKrS SKJr  S SKJr  S SKJr  SSKJr  S S	KJr  S
SKJr  S
SKJr  \ " S S5      5       r  " S S5      r! " S S5      r"\" SS9 " S S5      5       r#S r$S r%S r&S r'S r(\RR                  " \'\RT                  /S9r+\RR                  " \'\RX                  /S9r-\RR                  " \(\R\                  /S9r/ " S S5      r0 " S S 5      r1S! r2S" r3S# r4 " S$ S%5      r5 " S& S'\55      r6 " S( S)\55      r7S* r8S+ r9S, r:S- r;S. r<\1" 5       r=\" \=5      r>S/ r?S0 r@ " S1 S25      rA " S3 S4\R                  5      rC " S5 S65      rD " S7 S85      rEg)9    )annotationsN)TupleListDictCallable)	dataclass)TritonSemantic)TensorDescriptor   )InterpreterError)partial   )interpreter)irc                  p    \ rS rSr% SrS\S'   S\S'   \R                  " \S9r	S\S	'   S
 r
S rS rS rSrg)TensorHandle   z
data: numpy array
dtype: triton type, either pointer_type or scalar_type.
we don't store block_type here because the shape information is already available in the data field
attr: a dictionary of attributes
znp.arraydataztl.dtypedtype)default_factoryr   attrc                H    [        U R                  R                  5       5      $ N)boolr   allselfs    Z/var/www/html/land-doc-ocr/venv/lib/python3.13/site-packages/triton/runtime/interpreter.py__bool__TensorHandle.__bool__#   s    DIIMMO$$    c                ~    U R                   n[        US5      (       a  UR                  n[        US5      (       a  M  U$ )N
element_ty)r   hasattrr#   )r   r   s     r   get_element_tyTensorHandle.get_element_ty&   s7    

e\**$$E e\**r!   c                ^    [        U R                  R                  5       U R                  5      $ r   )r   r   copyr   r   s    r   cloneTensorHandle.clone,   s    DIINN,djj99r!   c                     X R                   U'   g r   )r   )r   keyvalues      r   set_attrTensorHandle.set_attr/   s    		#r!    N)__name__
__module____qualname____firstlineno____doc____annotations__dataclassesfielddictr   r   r%   r)   r.   __static_attributes__r0   r!   r   r   r      s<     NO""48D$8%:r!   r   c                       \ rS rSrS rS rSrg)BlockPointerHandle3   c                L    Xl         X l        X0l        X@l        XPl        X`l        g r   )baseshapestridesoffsetsblock_shapeorder)r   r?   r@   rA   rB   rC   rD   s          r   __init__BlockPointerHandle.__init__5   s!    	
&
r!   c                d   U R                   R                  5       nUR                  S-  n[        R                  " U R                   R
                  U R                  5      n[        R                  " U R                  [        S9n[        [        U R                  5      5       H  nS/[        U R                  5      -  nU R                  U   Xv'   U R                  U   R
                  [        R                  " U R                  U   5      -   R                  U5      nXCU-  U R                  U   R
                  -  R                  [        R                   5      -   nXa;   d  M  XXU R"                  U   R
                  :  -  US:  -  nM     [%        X@R                   R&                  R(                  5      nXE4$ )N   r   r   r   )r?   r%   primitive_bitwidthnpbroadcast_tor   rC   onesr   rangelenrB   arangereshaperA   astypeuint64r@   r   r   scalar)	r   boundary_checkdtype_ttn_bytesptrsmasksdim
bcast_dimsoffs	            r   materialize_pointers'BlockPointerHandle.materialize_pointers=   sO   99++---2tyy~~t/?/?@((5T--./Cs4#3#344J"..s3JO<<$))BIId6F6Fs6K,LLUUV`aCS=4<<+<+A+AAII"))TTD$tzz#';';!;<qI 0 D))//"8"89{r!   )r?   rC   rB   rD   r@   rA   N)r1   r2   r3   r4   rE   r]   r:   r0   r!   r   r<   r<   3   s    r!   r<   c                  2    \ rS rSr  SS jrS rSS jrSrg)	TensorDescHandleM   c                `    Xl         [        U5      U l        X l        X0l        X@l        XPl        g r   )r?   rO   ndimr@   rA   rC   padding)r   r?   r@   rA   rC   rd   s         r   rE   TensorDescHandle.__init__O   s'    	J	
&r!   c                
   U R                   R                  R                  5       S-  S:X  d   S5       e[        U R                  5      U R
                  :X  d   e[        U R                  5      U R
                  :X  d   eU R
                  S:  d   S5       eU R                  S S  H+  nUR                  R                  5       S-  S:X  a  M&   S5       e   U R                  S   R                  R                  5       S:X  d   S5       eg )	N   r   zbase must be 16-byte alignedr   z"descriptor cannot be 0 dimensionalzstride must be 16-byte alignedzlast dim must be contiguous)r?   r   itemrO   rA   rc   rC   )r   strides     r   validateTensorDescHandle.validateX   s    yy~~""$r)Q.N0NN.4<< DII---4##$		111yyA~CCC~ll3B'F;;##%*a/Q1QQ/ (||B$$))+q0O2OO0r!   c                   [        U5      U R                  :X  d   eU R                  R                  R                  nUR
                  S-  nUS   R                  U-  S-  S:X  d   S5       e[        R                  " U R                  R                  U R                  5      n[        R                  " U R                  [        S9n[        [        U R                  5      5       H  nS/[        U R                  5      -  nU R                  U   Xv'   X   R                  [        R                  " U R                  U   5      -   R                  U5      nXCU-  U R                  U   R                  -  R!                  [        R"                  5      -   nUSU:*  -  XR$                  U   R                  :  -  nM     UR                  [        R"                  :X  d   e['        X@R                  R                  R(                  5      nXE4$ )NrH   rh   rg   r   z*block offset start must be 16-byte alignedrI   r   )rO   rc   r?   r   r#   rJ   r   rK   rL   rC   rM   r   rN   rP   rQ   rA   rR   rS   r@   r   rT   )	r   rB   	scalar_tyitemsizerX   rY   rZ   r[   r\   s	            r   r]   %TensorDescHandle.materialize_pointersb   s   7|tyy(((IIOO..	//14  8+r1Q6d8dd6tyy~~t/?/?@((5T--./Cs4#3#344J"..s3JO<$$ryy1A1A#1F'GGPPQ[\Cc>DLL,=,B,BBJJ299UUDQ#X&#

30D0D*DEE 0 zzRYY&&&D))//"8"89{r!   )r?   rC   rc   rd   r@   rA   N)r?   r   r@   List[TensorHandle]rA   rq   rC   	List[int])rB   rq   )r1   r2   r3   r4   rE   rk   r]   r:   r0   r!   r   r`   r`   M   s    'Pr!   r`   T)frozenc                      \ rS rSr% SrS\S'   SrS\S'   SrS\S	'   SrS
\S'   Sr	S\S'   Sr
S\S'   SrS
\S'   SrS\S'   SrS\S'   SrS
\S'   Srg)InterpreterOptionsu   Nr9   extern_libsFr   debugTsanitize_overflowstrarch)fp8e5fp8e5b16fp8e4nvfp8e4b8fp8e4b15z
Tuple[str]supported_fp8_dtypesr0   !deprecated_fp8_dot_operand_dtypestf32default_dot_input_precision)r   tf32x3ieeeallowed_dot_input_precisionsr   intmax_num_imprecise_acc_defaultr   backend_name)r1   r2   r3   r4   rw   r6   rx   ry   r{   r   r   r   r   r   r   r:   r0   r!   r   ru   ru   u   sl    KE4"t"D#'^*^46%z6'--/I *I)*!3*%L#%r!   ru   c                &   U [         R                  :X  a  [         R                  $ U [         R                  :X  a  [         R                  $ U [         R
                  :X  a  [         R                  $ U [         R                  :X  a  [         R                  $ U $ r   )	rK   uint8int8uint16int16uint32int32rS   int64rI   s    r   _get_signed_np_dtyper      s[    ww		xx		xx		xxLr!   c                   [        U [        R                  5      (       a$  [        R                  " [        R
                  5      $ 0 [        R                  [        R                  " [        5      _[        R                  [        R                  " [        R                  5      _[        R                  [        R                  " [        R                  5      _[        R                  [        R                  " [        R                  5      _[        R                  [        R                  " [        R                  5      _[        R                  [        R                  " [        R                  5      _[        R                  [        R                  " [        R                  5      _[        R                  [        R                  " [        R                  5      _[        R                  [        R                  " [        R                  5      _[        R                   [        R                  " [        R                   5      _[        R"                  [        R                  " [        R"                  5      _[        R
                  [        R                  " [        R
                  5      _[        R$                  [        R                  " [        R                  5      _[        R&                  [        R                  " [        R                  5      _[        R(                  [        R                  " [        R                  5      _[        R*                  [        R                  " [        R                  5      _[        R,                  [        R                  " [        R                  5      _[        R.                  [        R                  " [        R                  5      0En[        U [        R0                  5      (       a[  [        U R2                  [        R                  5      (       a$  [        R                  " [        R
                  5      $ XR2                     $ X   $ r   )
isinstancetlpointer_typerK   r   rS   int1r   float16float32float64r   r   r   r   r   r   r   bfloat16float8e5float8e5b16
float8e4nv
float8e4b8float8e4b15
block_typer#   )tt_dtypenp_typess     r   _get_np_dtyper      s   (BOO,,xx		""
$


BHHRZZ( 	

BHHRZZ( 	

BHHRZZ(	
 	"''" 	"((288$ 	"((288$ 			288BII& 	"((288$ 			288BII& 	"((288$ 			288BII& 	RXXbii(  	RXXbhh'!" 	*#$ 	rxx)%& 	rxx)'( 	*)H, (BMM**h))2??;;88BII&&++,,r!   c                   [        [        SUR                   35      n[        [        SUR                   35      n[        R                  " U R	                  5       US9nXaR                  S-
  -	  S-  nUR                  UR
                  -
  S-
  nUR                  UR
                  -
  S-
  n	USUR
                  -  S-
  -  n
UR                  nUR                  nXaR
                  -	  SU-  S-
  -  R                  [        R                  5      nUS:H  n[        R                  " U5      (       a  [        R                  " U[        R                  S9n[        UR
                  5       H   nU
U-	  S-  nUR
                  U-
  UUS:H  '   M"     U
S:H  nSX   -
  X'   X-
  UUU-  '   X   X   -  SUR
                  -  S-
  -  X'   [        R                  " S[        R                  " X-
  U-   SU	-  S-
  5      5      nUR                  U5      nUR                  U5      nUR                  UR                  :  a  XR
                  UR
                  -
  -	  SUR
                  -  S-
  -  nU[        R                  R                   :X  a*  U
SUR
                  UR
                  -
  S-
  -  -  nUUS:  -   nUR                  U5      nO>U
R                  U5      UR
                  UR
                  -
  -  SUR
                  -  S-
  -  nUS:H  n[        R                  " U5      (       a  XaR
                  -	  SU-  S-
  -  R                  [        R                  5      nUS:g  nUU-  n[        R                  " U[        R                  S9nSU-
  X   U-
  -
  UU'   UU   UU   -	  SUR
                  UU   -
  -  -  UU'   UUR                  S-
  -  UUR
                  -  -  U-  nUR#                  U R$                  5      $ )NuintrI   r   r   )getattrrK   rJ   
frombuffertobytesfp_mantissa_widthexponent_biasrR   r   any
zeros_likerN   maximumminimum_irROUNDING_MODERTNErQ   r@   )inputinput_dtypeoutput_dtyperounding_modeinput_uint_dtypeoutput_unint_dtype	input_binsigninput_exponent_widthoutput_exponent_widthsignificand
bias_inputbias_outputexponentsubnormal_indexbit_posi	bit_indexzero_significand_indexexponent_outputsign_outputsignificand_outputcut_offnon_zero_exponent_indexshiftoutputs                             r   _convert_floatr      s   rT+*H*H)I#JK tL,K,K+L%MNemmo5EFI881<=ED&99K<Y<YY\]](;;l>\>\\_``[%B%B BaGHK**J,,K;;;FZAZ^_@_`hhikiqiqrH!mO	vvo
 --	:{445A%*d2I&1&C&Ca&GGIN# 6 "-!1$%(@$@!=G=U'/9:(3(DH`(`+///14(6$ jjBJJ0E0SWX\qWquvVv$wxO%,,-?@O++01K%%(G(GG).K.KlNlNl.lm,000A57C--222!Q;+H+H<KiKi+ilm+m%noG!3w{!C/667IJ)001CD+==@]@]]_#$(F(F#F!"KM &*O	vvo
 "?"??QJ^E^bcDcdllmomumuv"*a-),CCirxx8"#k/h6OR\6\!]o/A//RV[\kVl/l,0053IIJ/L?+l==AB<999;=OPF>>%++&&r!   c                .    [         R                  " U 5      $ r   )matherfxs    r   _erfr      s    88A;r!   c                6    [        U 5      [        U5      -  S-	  $ )N@   )r   )abs     r   
_umulhi_64r      s     FSVO""r!   )otypesc                  $    \ rS rSr\S 5       rSrg)ExtraFunctions   c                x    [         R                  " UR                  R                  U R                  X5      U5      $ r   )r   tensorbuildercreate_fp_to_fphandle)r   dst_tyfp_downcast_rounding	_semantics       r   _convert_custom_types$ExtraFunctions._convert_custom_types  s+    yy**::5<<fhnoor!   r0   N)r1   r2   r3   r4   staticmethodr   r:   r0   r!   r   r   r      s    p pr!   r   c                     \ rS rSr\R
                  R                  \R
                  R                  \R
                  R                  \R
                  R                  \R
                  R                  \R
                  R                  \R
                  R                  \R
                  R                  0r\R                  R                  \R                  R                  \R                  R                  \R                  R                  \R                  R                   \R                  R                   \R                  R"                  \R                  R"                  \R                  R$                  \R                  R$                  \R                  R&                  \R                  R&                  \R                  R(                  \R                  R(                  \R                  R*                  \R                  R*                  \R                  R,                  \R                  R,                  \R                  R.                  \R                  R.                  0
rSS jrS rS rS rS rS rS rS	 r S
 r!S r"S r#S r$S r%S r&S r'S r(S r)S r*S r+S r,S r-S r.S r/S r0S r1S r2S r3S r4S r5S r6S  r7S! r8S" r9S# r:S$ r;S% r<S& r=S' r>S( r?S) r@S* rAS+ rBS, rCS- rDS. rES/ rFS0 rGS1 rHS2 rIS3 rJS4 rKS5 rLS6 rMS7 rNS8 rOS9 rPS: rQS; rRS< rSS= rTS> rUS? rVS@ rWSA rXSB rYSC rZSD r[SE r\SF r]SG r^SH r_SI r`SJ raSK rbSL rcSM rdSN reSO rfSP rgSQ rhSR riSS rjST rkSU rlSV rmSW rnSX roSY rpSZ rqS[ rrS\ rsS] rtS^ ruS_ rvS` rwSa rxSb rySc rzSd r{Se r|Sf r}\Lr~\LrSg rSh rSi rSj rSk rSl rSm rSn rSo rSp rSq rSr rSs rSt rSu rSv rSw rSx rSy rSz rS{ rS| rS} rS~ rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS r S     SS jjrSS jrSS jrSS jr  SS jrS rSrg)InterpreterBuilderi  c                    S U l         [        5       U l        0 U l        [        R
                  U R                  S'   S U R                  S'   g )Nconvert_custom_typesc                    g)N)r   r   r   r0   )lhsTyperhsTypes     r   <lambda>-InterpreterBuilder.__init__.<locals>.<lambda>   s    Ir!   min_dot_size)r{   ru   optionscodegen_fnsr   r   r   s    r   rE   InterpreterBuilder.__init__  sB    	)+3A3W3W/0+M(r!   c                    XR                   S   :  d  [        S5      eX R                   S   :  d  [        S5      eX0R                   S   :  d  [        S5      eXU4U l        g )Nr   zx >= grid_dim[0]r   zy >= grid_dim[1]r   zz >= grid_dim[2])grid_dim
ValueErrorgrid_idxr   r   yzs       r   set_grid_idxInterpreterBuilder.set_grid_idx"  s^    ==##/00==##/00==##/00q	r!   c                    XU4U l         g r   )r   )r   nxnynzs       r   set_grid_dimInterpreterBuilder.set_grid_dim+  s    r!   c                "    [         R                  $ r   )r   r   r   s    r   get_half_tyInterpreterBuilder.get_half_ty0      zzr!   c                "    [         R                  $ r   )r   r   r   s    r   get_bf16_tyInterpreterBuilder.get_bf16_ty3      {{r!   c                "    [         R                  $ r   )r   r   r   s    r   get_float_tyInterpreterBuilder.get_float_ty6  r  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_double_ty InterpreterBuilder.get_double_ty9  r  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_int1_tyInterpreterBuilder.get_int1_ty<      wwr!   c                "    [         R                  $ r   )r   r   r   s    r   get_int8_tyInterpreterBuilder.get_int8_ty?  r  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_uint8_tyInterpreterBuilder.get_uint8_tyB      xxr!   c                "    [         R                  $ r   )r   r   r   s    r   get_int16_tyInterpreterBuilder.get_int16_tyE  r  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_uint16_ty InterpreterBuilder.get_uint16_tyH      yyr!   c                "    [         R                  $ r   )r   r   r   s    r   get_int32_tyInterpreterBuilder.get_int32_tyK  r  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_uint32_ty InterpreterBuilder.get_uint32_tyN  r$  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_int64_tyInterpreterBuilder.get_int64_tyQ  r  r!   c                "    [         R                  $ r   )r   rS   r   s    r   get_uint64_ty InterpreterBuilder.get_uint64_tyT  r$  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e4nv_ty!InterpreterBuilder.get_fp8e4nv_tyW      }}r!   c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e4b15_ty"InterpreterBuilder.get_fp8e4b15_tyZ      ~~r!   c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e4b8_ty!InterpreterBuilder.get_fp8e4b8_ty]  r4  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e5_tyInterpreterBuilder.get_fp8e5_ty`  r  r!   c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e5b16_ty"InterpreterBuilder.get_fp8e5b16_tyc  r8  r!   c                .    [         R                  " X5      $ r   )r   r   )r   elt_ty
addr_spaces      r   
get_ptr_tyInterpreterBuilder.get_ptr_tyf  s    v22r!   c                .    [         R                  " X5      $ r   )r   r   )r   r   r@   s      r   get_block_tyInterpreterBuilder.get_block_tyi  s    }}U**r!   c                z    [        [        R                  " U/[        R                  S9[        R
                  5      $ NrI   )r   rK   arraybool_r   r   r   r-   s     r   get_int1InterpreterBuilder.get_int1l  s$    BHHeWBHH=rwwGGr!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   	get_uint8InterpreterBuilder.get_uint8o  $    BHHeWBHH=rxxHHr!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   get_int8InterpreterBuilder.get_int8r  s$    BHHeWBGG<bggFFr!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   
get_uint16InterpreterBuilder.get_uint16u  $    BHHeWBII>		JJr!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   	get_int16InterpreterBuilder.get_int16x  rT  r!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   
get_uint32InterpreterBuilder.get_uint32{  r[  r!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   	get_int32InterpreterBuilder.get_int32~  rT  r!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  rS   r   rN  s     r   
get_uint64InterpreterBuilder.get_uint64  r[  r!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   	get_int64InterpreterBuilder.get_int64  rT  r!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   get_fp16InterpreterBuilder.get_fp16  $    BHHeWBJJ?LLr!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   get_fp32InterpreterBuilder.get_fp32  rn  r!   c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rK  )r   rK   rL  r   r   rN  s     r   get_fp64InterpreterBuilder.get_fp64  rn  r!   c                T    [        [        R                  " S/[        U5      S9U5      $ Nr   rI   )r   rK   rL  r   )r   types     r   get_null_value!InterpreterBuilder.get_null_value  s!    BHHaSd0CDdKKr!   c                    U R                   c  [        S5      e[        [        R                  " U R                   U   /[        R
                  S9[        R
                  5      $ )Nzgrid_idx is NonerI   )r   r   r   rK   rL  r   r   r   axiss     r   create_get_program_id(InterpreterBuilder.create_get_program_id  sD    == /00BHHdmmD&9%:"((KRXXVVr!   c                    [        [        R                  " U R                  U   /[        R                  S9[
        R                  5      $ rK  )r   rK   rL  r   r   r   r{  s     r   create_get_num_programs*InterpreterBuilder.create_get_num_programs  s.    BHHdmmD&9%:"((KRXXVVr!   c                    [        [        R                  " UR                  [        S9[
        R                  5      nS nU R                  XXbX45      $ rK  )r   rK   	ones_liker   r   r   r   create_masked_load)r   ptr_0_1is_volatilemaskothers          r   create_loadInterpreterBuilder.create_load  s;    BLL>H&&s%RMMr!   c                    [        [        R                  " UR                  [        S9[
        R                  5      nU R                  XUS S 5      $ rK  )r   rK   r  r   r   r   r   create_masked_store)r   r  valr  r  r  s         r   create_storeInterpreterBuilder.create_store  s8    BLL>H''$dCCr!   c                   UR                  5       n[        U5      nUc)  [        [        R                  " UR
                  US9U5      n[        R                  " UR
                  UR
                  UR
                  U5      n	[        X5      $ rK  )r%   r   r   rK   r   r   _interpreterload)
r   rX   r  r  cache_modifiereviction_policyr  rV   dtype_nprets
             r   r  %InterpreterBuilder.create_masked_load  sg    &&( *= tyy!I8TE		499ejj(KC**r!   c                n    [         R                  " UR                  UR                  UR                  5      $ r   )r  storer   )r   rX   r-   r  r  r  s         r   r  &InterpreterBuilder.create_masked_store  s#    !!$))UZZCCr!   c                   UR                   R                  nUR                  nU[        R                  :X  a  U[        R                  :X  d(  U[        R                  :X  aX  U[        R                  :X  aD  [        UR                  X4S 5      R                  [        U5      5      n[        XRR                  5      $ [        UR                  R                  [        U5      5      UR                  5      $ r   )r   rT   r   r   r   r   r   viewr   r   rR   )r   srcdst_typesrc_element_typedst_element_typer   s         r   	cast_implInterpreterBuilder.cast_impl  s    99++#??+0@BJJ0N

*/?2;;/N!#((,<PTUZZ[hiq[rsDoo66h0G H(//ZZr!   c                $    U R                  X5      $ r   r  r   r  r  s      r   r   InterpreterBuilder.<lambda>      $..2Or!   c                $    U R                  X5      $ r   r  r  s      r   r   r    r  r!   c                $    U R                  X5      $ r   r  r  s      r   r   r    r  r!   c                $    U R                  X5      $ r   r  r  s      r   r   r    r  r!   c                $    U R                  X5      $ r   r  r  s      r   r   r    s    s0Mr!   c                $    U R                  X5      $ r   r  r  s      r   r   r    r  r!   c                $    U R                  X5      $ r   r  )r   r  r  	is_signeds       r   r   r    s    T^^C=Zr!   c                    UR                   R                  nUR                  n[        UR                  XEU5      R	                  [        U5      5      n[        XbR                  5      $ r   )r   rT   r   r   r  r   r   )r   r  r  r   r  r  r   s          r   r   "InterpreterBuilder.create_fp_to_fp  sP    99++#??chh(8MZ__`mnv`wxD//22r!   c                r    [        UR                  R                  [        U5      5      UR                  5      $ r   )r   r   r  r   rT   r  s      r   create_bitcast!InterpreterBuilder.create_bitcast  s%    CHHMM-*ABHOOTTr!   c                x    [        U" UR                  UR                  5      UR                  R                  5      $ r   r   r   r   rT   )r   lhsrhsops       r   	binary_opInterpreterBuilder.binary_op  s(    Bsxx2CII4D4DEEr!   c                B    U R                  X[        R                  5      $ r   r  rK   addr   r  r  s      r   r   r    s    "&&)Ir!   c                B    U R                  X[        R                  5      $ r   r  rK   multiplyr  s      r   r   r        "++)Nr!   c                B    U R                  X[        R                  5      $ r   r  rK   divider  s      r   r   r    s    ")))Lr!   c                B    U R                  X[        R                  5      $ r   r  rK   fmodr  s      r   r   r        "'')Jr!   c                B    U R                  X[        R                  5      $ r   r  rK   subtractr  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r        s(Mr!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    s    "))1Tr!   c                $    U R                  X5      $ r   create_idivr  s      r   r   r        )9)9#)Cr!   c                $    U R                  X5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    s    s(Hr!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   )r  rK   
left_shiftr  s      r   r   r    s    s(Or!   c                B    U R                  X[        R                  5      $ r   )r  rK   right_shiftr  s      r   r   r    s    "..)Qr!   c                B    U R                  X[        R                  5      $ r   r  rK   r   r  s      r   r   r        $..2::*Nr!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r        T^^Cbjj-Qr!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r        DNN3RZZ,Pr!   c                B    U R                  X[        R                  5      $ r   r  rK   r   r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  rK   
less_equalr  s      r   r   r        DNN3R]],Sr!   c                B    U R                  X[        R                  5      $ r   r  rK   lessr  s      r   r   r        DNN3RWW,Mr!   c                B    U R                  X[        R                  5      $ r   r  rK   greater_equalr  s      r   r   r        DNN3REUEU,Vr!   c                B    U R                  X[        R                  5      $ r   r  rK   greaterr  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  rK   equalr  s      r   r   r    s    4>>#BHH+Mr!   c                B    U R                  X[        R                  5      $ r   r  rK   	not_equalr  s      r   r   r    s    4>>#BLL+Qr!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r   r  s      r   r   r        DNN3RXX,Nr!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r        DNN3R\\,Rr!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   r   r  s      r   r   r    r
  r!   c                B    U R                  X[        R                  5      $ r   r  r  s      r   r   r    r  r!   c                B    U R                  X[        R                  5      $ r   )r  rK   bitwise_andr  s      r   r   r        s(Pr!   c                B    U R                  X[        R                  5      $ r   )r  rK   bitwise_xorr  s      r   r   r     r  r!   c                B    U R                  X[        R                  5      $ r   )r  rK   
bitwise_orr  s      r   r   r    s    t~~c'Nr!   c                    [        UR                  [        R                  " UR                  UR                  5      -
  UR                  -  UR                  R
                  5      $ r   )r   r   rK   r  r   rT   r  s      r   r  InterpreterBuilder.create_idiv  sC     SXX#(((CCPRUR[R[RbRbccr!   c                >   [        UR                  R                  5      n[        UR                  R                  5      nUR                  R                  U5      Ul        UR                  R                  U5      Ul        U R	                  X[
        R                  5      $ r   )r   r   r   rR   r  rK   r  )r   r  r  	lhs_dtype	rhs_dtypes        r   create_ashrInterpreterBuilder.create_ashr  sc    (8	(8	88??9-88??9-~~c77r!   c                V   UR                   R                  nU[        R                  :X  d  U[        R                  :X  a>  [        [        UR                   UR                   5      UR                  R                  5      $ [        [        SUR                  S-  S-   35      nUR                   R                  U5      nUR                   R                  U5      n[        R                  " XV5      UR                  S-  -	  n[        UR                  U5      UR                  R                  5      $ )Nr   rH   r   )r   r   rK   r   rS   r   np_umulhi_u64rT   r   ro   rR   r  )r   r  r  r   compute_dtypelhs_datarhs_dataret_datas           r   create_umulhi InterpreterBuilder.create_umulhi  s    BHH 2chh A399CSCSTT#B$u~~/AA/E.F(GHMxx}5Hxx}5H{{865>>A;MNH 6		8H8HIIr!   c                    [        U" UR                  UR                  UR                  5      UR                  R                  5      $ r   r  )r   r  r  r  r  s        r   
ternary_opInterpreterBuilder.ternary_op  s.    Bsxx5::>@R@RSSr!   c                D    U R                  XU[        R                  5      $ r   )r*  rK   clip)r   arglohipropagate_nanss        r   r   r  "  s    doocWY[][b[b>cr!   c                D    U R                  XU[        R                  5      $ r   )r*  rK   where)r   condr  r  s       r   r   r  #  s    CQSQYQY1Zr!   c                    [        UR                  UR                  -  UR                  -   UR                  R                  5      $ r   r  r   s       r   
create_fmaInterpreterBuilder.create_fma%  s,    AFFQVVOaff4aggnnEEr!   c                b    [        U" UR                  5      UR                  R                  5      $ r   r  )r   r.  r  s      r   unary_opInterpreterBuilder.unary_op)  s!    BsxxL#))*:*:;;r!   c                .   UR                   nUR                  S-
  n[        [        SUR                   35      nUR                  R                  U5      nSU-  S-
  nXV-  R                  [        U5      5      n[        XqR                   R                  5      $ )Nr   r   )	r   rJ   r   rK   r   r  r   r   rT   )r   r.  rV   mask_bitwidthnp_uint_dtyper   r  r  s           r   create_fabsInterpreterBuilder.create_fabs,  s    99 33a7d8+F+F*G$HIxx}}]+]"a'{  x!89C!1!122r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   cosr   r.  s     r   r   r  6      4==bff#=r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   exprB  s     r   r   r  7  rC  r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   exp2rB  s     r   r   r  8      DMM#rww$?r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   absrB  s     r   r   r  9  s    DMM#rvv$>r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   floorrB  s     r   r   r  :  s    T]]3%Ar!   c                B    U R                  U[        R                  5      $ r   )r9  rK   ceilrB  s     r   r   r  ;  rH  r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   logrB  s     r   r   r  <  rC  r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   log2rB  s     r   r   r  =  rH  r!   c                B    U R                  U[        R                  5      $ r   r9  rK   sqrtrB  s     r   r   r  >  s    DMM#rww,Gr!   c                B    U R                  U[        R                  5      $ r   rT  rB  s     r   r   r  ?  rH  r!   c                B    U R                  U[        R                  5      $ r   )r9  rK   sinrB  s     r   r   r  @  rC  r!   c                    UR                   R                  [        R                  :X  a  [	        UR                   5      O[        UR                   5      n[        X!R                  R                  5      $ r   )r   r   rK   r   np_erf_fp32np_erf_fp64r   rT   )r   r.  r  s      r   
create_erfInterpreterBuilder.create_erfB  sF    '*xx~~'Ck#((#UXU]U]I^C!1!122r!   c                    [        S[        R                  " UR                  5      -  UR                  R
                  5      $ Nr   )r   rK   rU  r   r   rT   rB  s     r   create_rsqrtInterpreterBuilder.create_rsqrtF  s+    A 113993C3CDDr!   c                t    [        UR                  R                  U5      UR                  R                  5      $ r   )r   r   rQ   r   rT   )r   r.  r@   allow_reorders       r   r   r  J  s(    \#((JZJZ[`Jacfclclcscs=tr!   c                    [        [        R                  " UR                  U5      UR                  R
                  5      $ r   )r   rK   	transposer   r   rT   )r   r.  perms      r   create_transInterpreterBuilder.create_transL  s(    BLL48#)):J:JKKr!   c                   UR                   nUR                   nUR                  R                  S:X  a  UR                  R                  5       (       d9  UR                  R                  S:X  a  UR                  R                  5       (       a  [	        XaR                  [
        R                  S 5      R                  [        R                  5      n[	        XrR                  [
        R                  S 5      R                  [        R                  5      n[        [        R                  " XgUR                   R                  S9UR                   -   UR                  R                  5      $ )NrH   rI   )r   r   rJ   is_floatingr   r   r   r  rK   r   matmulrT   )r   r   r   dinput_precisionmax_num_imprecise_acca_datab_datas           r   
create_dotInterpreterBuilder.create_dotO  s    GG&&!+0C0C0E0EGG&&!+0C0C0E0E#FGGRZZFKKBJJWF#FGGRZZFKKBJJWFBIIfAFFLLIAFFRTUT[T[TbTbccr!   c                x    [        [        R                  " X#[        R                  S9[        R                  5      $ rK  )r   rK   rP   r   r   )r   ret_tystartstops       r   create_make_range$InterpreterBuilder.create_make_rangeX  s"    BIIeBBHHMMr!   c                   Uc;  [        [        R                  " UR                  [        S9[
        R                  5      n[        R                  " UR                  UR                  [        R                  " UR                  5      5      n[        R                  " XSU4S9S   nUS==   [        R                  " UR                  5      R                  5       -  ss'   [        U[
        R                  5      $ )NrI   r   )binsrN   )r   rK   r  r   r   r   r   r3  r   	histogramlogical_notsumr   )r   r   rz  r  r{  s        r   create_histogram#InterpreterBuilder.create_histogram[  s    <TYYd CRWWMDxx		499bmmDII.FGLLD	B1E	!tyy15577Irxx00r!   c                    [        [        R                  " UR                  UR                  US9UR                  R
                  5      $ )Nr|  )r   rK   take_along_axisr   r   rT   )r   r  indicesr|  s       r   create_gather InterpreterBuilder.create_gathere  s3    B..sxxDQSVS\S\ScScddr!   c                    UR                  5       nUR                  n[        SUS-  5      n[        UR                  XRR                  R                  [        R                  5      -  -   UR                  5      $ )Nr   rH   )	r%   rJ   maxr   r   rR   rK   rS   r   )r   r  offsetrV   element_bitwidthelement_bytewidths         r   create_addptr InterpreterBuilder.create_addptrj  sc    %%'#66#3q#89CHH'8;;;M;Mbii;X'XXZ]ZcZcddr!   c                   UR                  U5      u  pxUR                  5       n	[        U	5      n
Uc  S nOU[        R                  R
                  :X  a*  [        [        R                  " UR                  U
S9U	5      nO`U[        R                  R                  :X  a4  [        [        R                  " UR                  [        S5      U
S9U	5      nO[        SU 35      eU R                  XxXXV5      $ )NrI   nanzunsupported padding option )r]   r%   r   r   PADDING_OPTIONPAD_ZEROr   rK   r   r   PAD_NAN	full_likefloatr   r  )r   r  rU   padding_optionr  r  r  rX   rY   rV   r  r  s               r   create_tensor_pointer_load-InterpreterBuilder.create_tensor_pointer_loadq  s    ..~>&&( *!Es11::: tyy!I8TEs11999 diiuX!VX`aE:>:JKLL&&tE?hhr!   c                N    UR                  U5      u  pgU R                  XbXtU5      $ r   r]   r  )r   r  r-   rU   r  r  rX   rY   s           r   create_tensor_pointer_store.InterpreterBuilder.create_tensor_pointer_store  s)    ..~>''UO\\r!   c                    [        [        R                  " UR                  U5      UR                  R
                  5      $ r   )r   rK   expand_dimsr   r   rT   )r   r.  r|  s      r   create_expand_dims%InterpreterBuilder.create_expand_dims  s(    BNN388T:CII<L<LMMr!   c                    [        [        R                  " UR                  U5      UR                  R
                  5      $ r   )r   rK   rL   r   r   rT   )r   r.  r@   s      r   create_broadcast#InterpreterBuilder.create_broadcast  s(    BOOCHHe<cii>N>NOOr!   c                    [        [        R                  " UR                  UR                  /5      UR                  R
                  5      $ r   )r   rK   concatenater   r   rT   r  s      r   
create_catInterpreterBuilder.create_cat  s/    BNNCHHchh+?@#))BRBRSSr!   c                    [        [        R                  " UR                  UR                  /SS9UR                  R
                  5      $ )Nrh   r  )r   rK   stackr   r   rT   r  s      r   create_joinInterpreterBuilder.create_join  s1    BHHchh%9CSYYEUEUVVr!   c                    [        UR                  S   UR                  R                  5      [        UR                  S   UR                  R                  5      4$ )N).r   ).r   r  )r   r  s     r   create_splitInterpreterBuilder.create_split  sE    SXXf-syy/?/?@,sxxX^O_adajajaqaqBrssr!   c           	        UR                   n[        UR                  [        R                  5      (       aS  [        [        R                  " X2R                  S   [        UR                  5      S9UR                  R                  5      $ [        [        R                  " X2R                  [        UR                  5      S9UR                  R                  5      $ rv  )r@   r   r   r   r   r   rK   fullr   r   rT   )r   rt  r.  r@   s       r   create_splatInterpreterBuilder.create_splat  s    cii//xx{-PSPYPYBZ []`]f]f]m]mnnxx}SYY?W XZ]ZcZcZjZjkkr!   c           	         [        [        R                  " SUR                  S   [	        UR
                  5      S9UR
                  R                  5      $ )Nr   r   rI   )r   rK   r  r   r   r   rT   rB  s     r   create_unsplat!InterpreterBuilder.create_unsplat  s:    BGGE388A;mCII>VWY\YbYbYiYijjr!   c                   X@R                   ;  a  [        SU 35      eU R                   U   n[        [        R                  " UR
                  UR
                  UR
                  U5      UR                  R                  5      $ )Nunsupported semantic )ir_sem_to_interpreter_semr   r   r  
atomic_casr   r   rT   )r   r  cmpr  semscopes         r   create_atomic_cas$InterpreterBuilder.create_atomic_cas  si    4444SE:;;,,S1L33CHHchhRUVX[XaXaXhXhiir!   c           	     \   XR                   ;  a  [        SU 35      eXPR                  ;  a  [        SU 35      eU R                   U   nU R                  U   n[        [        R
                  " XR                  UR                  UR                  U5      UR                  R                  5      $ )Nzunsupported rmwOp r  )	ir_rmw_op_to_interpreter_rmw_opr   r  r   r  
atomic_rmwr   r   rT   )r   rmwOpr  r  r  r  r  s          r   create_atomic_rmw$InterpreterBuilder.create_atomic_rmw  s    <<<1%9::4444SE:;;44U;,,S1L33E88SXXtyyZ]^`c`i`i`p`pqqr!   c                    [        S5      e)Nz4extern_elementwise not supported in interpreter modeNotImplementedError)r   libNamelibPathsymbolargListretTypeisPures          r   create_extern_elementwise,InterpreterBuilder.create_extern_elementwise  s    !"XYYr!   c                    [        S5      e)Nz,inline_asm not supported in interpreter moder  )r   	inlineAsmconstraintsvaluesrw  r  packs          r   create_inline_asm$InterpreterBuilder.create_inline_asm  s    !"PQQr!   c                D   SU R                   S    SU R                   S    SU R                   S    S3nU(       a  USU 3-  nU(       a  [        R                  " SS	 0S
9  U H  n[        USUR                   3-   5        M      U(       a  [        R                  " S S
9  g g )N(r   z, r   r   ) r   c                    SU S 3$ )N0x02xr0   r   s    r   r   1InterpreterBuilder.create_print.<locals>.<lambda>  s    b3Lr!   )	formatter)r   rK   set_printoptionsprintr   )r   prefixhexr  isSignedmsgr-   s          r   create_printInterpreterBuilder.create_print  s    
 $--"#2dmmA&6%7r$--:J9K1MQvh<C52H*IJE#!EJJ<(() $/ r!   c                "    U(       d   U 5       eg r   r0   )r   	conditionmessages      r   create_assert InterpreterBuilder.create_assert  s    &WI&yr!   c                     U(       d   S5       eg )NzAssume failedr0   )r   r  s     r   create_assume InterpreterBuilder.create_assume  s    )/)yr!   c                    g r   r0   r   s    r   create_barrier!InterpreterBuilder.create_barrier  s    r!   c                d    U Vs/ s H  owR                  5       PM     nn[        XX8XV5      $ s  snf r   )r)   r<   )	r   r?   r@   rA   rB   rC   rD   r  new_offsetss	            r   create_make_block_ptr(InterpreterBuilder.create_make_block_ptr  s.    4;<G&||~G<!$w[XX =s   -c                   [        UR                  5      [        U5      :w  a  [        S5      eUR                   Vs/ s H  o3R                  5       PM     nn[	        UR
                  UR                  UR                  XAR                  UR                  5      n[        [        U5      5       H1  nUR                  U   =R                  X&   R                  -  sl        M3     U$ s  snf )Nz len(ptr.offsets) != len(offsets))rO   rB   r   r)   r<   r?   r@   rA   rC   rD   rN   r   )r   r  rB   r  r  r  r   s          r   create_advance!InterpreterBuilder.create_advance  s    s{{s7|+?@@47KK@K&||~K@ 399ckk;P_P_adajajks7|$AKKN7:??2 %
	 As   C#c                @    [        XX4U5      nUR                  5         U$ r   )r`   rk   )r   r?   r@   rA   tensor_shaper  rd   descs           r   create_make_tensor_descriptor0InterpreterBuilder.create_make_tensor_descriptor  s    WGLr!   c           	        [        U[        5      (       d   eUR                  U5      u  pVUR                  5       n[	        U5      nUR
                  n	U	[        R                  R                  :X  a*  [        [        R                  " UR                  US9U5      n
O`U	[        R                  R                  :X  a4  [        [        R                  " UR                  [        S5      US9U5      n
O[!        SU	 35      eU R#                  XVXUSS9$ )NrI   r  zunsupported padding F)r  r  r  )r   r`   r]   r%   r   rd   r   r  r  r   rK   r   r   r  r  r  r   r  )r   r  r  r  r  rX   r  rV   r  rd   r  s              r   create_descriptor_load)InterpreterBuilder.create_descriptor_load  s    $ 01111..w7
&&( *,,c((111 tyy!I8TE**222 diiuX!VX`aE3G9=>>&&t57FTY ' [ 	[r!   c                P    UR                  U5      u  pEU R                  XBUS S 5      $ r   r  )r   r  r-   r  rX   r  s         r   create_descriptor_store*InterpreterBuilder.create_descriptor_store  s+    ..w7
''T4FFr!   c                   UR                   R                  R                  n[        U5      n[        R
                  " UR                  R                  S   UR                  S   /US9nS nS n	[        UR                  5       HC  u  p[        U[        R                  5      U/nU R                  XX5      R                  XzS S 24'   ME     [        Xu5      $ )Nr   rh   rI   )r?   r   r#   r   rK   zerosr   r@   rC   	enumerater   r   r   r  )r   r  	x_offsetsy_offsetrw  r   np_dtyperesultr  r  r   x_offsetr  s                r   create_descriptor_gather+InterpreterBuilder.create_descriptor_gather  s    		** '9>>//2D4D4DR4HIQYZ$Y^^4KA#Hbhh7BG66tnfkkFa4L 5 F**r!   c                    [        UR                  5       HV  u  pV[        UR                  U   UR                  5      n[        U[        R
                  5      U/nU R                  XU5        MX     g r   )r  r   r   r   r   r   r  )	r   r  r-   r  r  r   r  slicer  s	            r   create_descriptor_scatter,InterpreterBuilder.create_descriptor_scatter  sT    $Y^^4KA A<E#Hbhh7BG((g> 5r!   c                $   [        U5      nSUR                  ;   a*  [        [        R                  " SSUS9UR
                  5      $ U[        R                  :X  a*  [        [        R                  " SSUS9UR
                  5      $ [        SU 35      e)Nr   r   rh   rI   Tzunsupported type )r   namer   rK   r  rT   rM  	TypeError)r   rw  np_types      r   get_all_ones_value%InterpreterBuilder.get_all_ones_value	  ss    %GLL 2W =t{{KK 4w ?MM/v677r!   )r{   r   r   r   r   NreturnNone)zero)r?   r   r@   rq   rA   rq   r  rr   r  r   rd   rz   )r  r`   r  rq   )r  r`   r-   r   r  rq   )r  r`   r  r   r  r   )r  r`   r-   r   r  r   r  r   )r1   r2   r3   r4   r   MEM_SEMANTICACQUIREr  RELEASERELAXEDACQUIRE_RELEASEr  	ATOMIC_OPADDRMW_OPFADDMINUMINMAXUMAXANDORXORXCHGr  rE   r   r  r  r
  r  r  r  r  r  r  r"  r&  r)  r,  r/  r2  r6  r:  r=  r@  rE  rH  rO  rR  rV  rY  r]  r`  rc  rf  ri  rl  rp  rs  rx  r}  r  r  r  r  r  r  create_si_to_fpcreate_ui_to_fpcreate_fp_to_sicreate_fp_to_uicreate_fp_extcreate_fp_trunccreate_int_castr   r  r  create_faddcreate_fmulcreate_fdivcreate_fremcreate_fsub
create_mulcreate_precise_divfcreate_sdivcreate_udivcreate_sremcreate_urem
create_add
create_sub
create_shlcreate_lshrcreate_minsicreate_minuicreate_minimumfcreate_minnumfcreate_maxsicreate_maxuicreate_maximumfcreate_maxnumfcreate_icmpSLEcreate_icmpSLTcreate_icmpSGEcreate_icmpSGTcreate_icmpULEcreate_icmpULTcreate_icmpUGEcreate_icmpUGTcreate_icmpEQcreate_icmpNEcreate_fcmpOLTcreate_fcmpOGTcreate_fcmpOLEcreate_fcmpOGEcreate_fcmpOEQcreate_fcmpONEcreate_fcmpULTcreate_fcmpUGTcreate_fcmpULEcreate_fcmpUGEcreate_fcmpUEQcreate_fcmpUNE
create_and
create_xor	create_orcreate_int_to_ptrcreate_ptr_to_intr  r  r'  r*  create_clampfcreate_selectr6  r9  r>  
create_cos
create_expcreate_exp2create_iabscreate_floorcreate_ceil
create_logcreate_log2create_precise_sqrtcreate_sqrt
create_sinr\  r`  create_reshaperg  rq  rw  r~  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r:   r0   r!   r   r   r     s     ,";";"C"C  ,";";"C"C  ,";";"C"C((,*C*C*S*S	! 	<..22L//44<..22L//44<..22L//44<..22,--00<..22L//44'#N"%
3+HIGKIKIKIMMMLW
WN
D+D[ POOOOOOOMMOOZO3UF JKNKLKJKNKMJTCKCKJKJKHJMJOJQKNLNLQOPNNLNLQOPNSNMNVNPNSNMNVNPNMMQMMNPNSNVNNNRNMNPNSNVNNNRNPJPJNI&&d8	JT dMZMF<3 >J=J?K>KAL?K=J?KG?K=J3E uNLdN1e
ei]NPTWtlkjrZR0'*Y
 `f4=JNY\[ G	+?,8?8r!   r   c                H   ^ [        U5      mUS.U4S jjn[        XU5        g )N)memberc           
     ~   > U " U0 UR                  5        VVs0 s H  u  p4US:w  d  M  X4_M     snnDST0D6$ s  snnf )Nr   items)ru  argskwargskvsemantics        r   r   _patch_attr.<locals>.<lambda>  s\     :kMS\\^AVM[TQDEDT BFM[AV:k bj:kAVs   99)r	   setattr)objr  ru  r   
new_memberr}  s        @r   _patch_attrr    s$    g&H&, lJ Cz"r!   c                    [         R                  " U 5       H7  u  p#[        R                  R	                  U5      (       d  M+  [        XX15        M9     g r   )inspect
getmembersr   core
is_builtinr  )pkgr   r  ru  s       r   _patch_builtinr    s8    **3/77f%%63 0r!   c                x   ^ S mS nS U l         U4S jU l        S U l        S U l        [	        U5      U l        g )Nc                h    U R                   R                  nUR                  S:X  a  [        U5      $ S$ )Nr   T)r   r   sizer   )r   r   s     r   	_get_bool%_patch_lang_tensor.<locals>._get_bool$  s,    {{ "YY!^tDz55r!   c                   [        [        R                  " U R                  R                  5      U R                  R
                  5      nU R                  R                  5       (       d   e[        U R                  R                  5      nUS   US   sUS'   US'   [        R                  R                  U R
                  U5      n[        R                  R                  X5      $ )Nrh   )r   rK   re  r   r   r   rw  is_blocklistr@   r   r  r   r   )r   r   rC   res_tys       r   _get_transpose*_patch_lang_tensor.<locals>._get_transpose*  s    bll4;;+;+;<dkk>O>OPyy!!####499??++6r?KO(BR##DJJ<ww~~f--r!   c                @    [        U R                  R                  5      $ r   )r   r   r   r   s    r   r   $_patch_lang_tensor.<locals>.<lambda>2  s    C(8(8$9r!   c                   > T" U 5      $ r   r0   )r   r  s    r   r   r  3  s	    9T?r!   c                @    [        U R                  R                  5      $ r   )reprr   r   r   s    r   r   r  4  s    4(8(8#9r!   c                @    [        U R                  R                  5      $ r   )rz   r   r   r   s    r   r   r  5  s    #dkk&6&6"7r!   )	__index__r   __repr____str__propertyT)r   r  r  s     @r   _patch_lang_tensorr  "  s8    6. :F2FO9FO7FN'FHr!   c                  2    \ rS rSrS rS rS rS rS rSr	g)	ReduceScanOpInterfacei9  c                    Xl         X l        g r   r|  
combine_fn)r   r|  r  s      r   rE   ReduceScanOpInterface.__init__;  s    	$r!   c                L    Ub!  U[        U5      :  a  [        SU SU 35      eg g )Nzaxis z out of bounds for shape )rO   r   )r   r@   r|  s      r   
check_axis ReduceScanOpInterface.check_axis?  s4    E
 2uTF*CE7KLL !3r!   c                    U Hi  n[        U[        R                  R                  5      (       d  [	        S[        U5       35      eU R                  UR                  U R                  5        Mk     g )Nzinput must be a tensor, got )	r   r   r  r   r   rw  r  r@   r|  )r   r   r.  s      r   check_tensor"ReduceScanOpInterface.check_tensorC  sN    Cc277>>22 #?S	{!KLLOOCIItyy1 r!   c                j   [        U5      n[        US5      (       aM  UR                  (       a<  UR                  U5      n[        R
                  " U[        UR                  5      5      nO[        R                  " U/US9nUn[        R                  R                  [        XR                  5      U5      $ )Nr@   rI   )r   r$   r@   rR   r   r   r  rK   rL  r  r   r   rT   )r   r  r   r	  ret_types        r   	to_tensorReduceScanOpInterface.to_tensorI  sz     '3  SYY**X&C}}UDO<H((C51CHww~~l3=xHHr!   c                    [        U[        5      (       d  U R                  U45      S   $ U R                  U5        U R	                  U5      n[        U[
        [        45      (       a  [        U5      $ U4$ Nr   )r   tupleapplyr  
apply_implr  )r   r   r  s      r   r  ReduceScanOpInterface.applyS  sb    %''::ui(++% ooe$'dE];;uSzH#Hr!   r  N)
r1   r2   r3   r4   rE   r  r  r  r  r:   r0   r!   r   r  r  9  s    %M2IIr!   r  c                  J   ^  \ rS rSrU 4S jrS rS rS	S jrS rS r	Sr
U =r$ )
	ReduceOpsi[  c                0   > [         TU ]  X5        X0l        g r   )superrE   	keep_dims)r   r|  r  r  	__class__s       r   rE   ReduceOps.__init__]  s    *"r!   c                    / nU Hh  nUb  UR                  U5        M  SnUR                  U R                  UR                  R                  R	                  5       UR
                  5      5        Mj     [        U5      U4$ r  )appendr  r   r   flattenr   r  )r   r   r|  r  r   s        r   unravelReduceOps.unravela  sg    D

4 

4>>$++*:*:*B*B*DdjjQR  Sz4r!   c                  ^ ^^^ T R                   nT R                  TT R                   5      u  mn/ n/ nTS   R                  R                  R                  nUSU XcS-   S  -   nT Hi  nUR                  UR                  R                  5        UR                  [        R                  " XxR                  R                  R                  S95        Mk     [        US   R                  5       GHf  n	[        R                  " X5      mTSU TUS-   S  -   m[        UUU 4S j[        U5       5       5      n
TU   S:X  aH  [        [        U5      5       H.  nX   R                  R                  R                  5       X[   T'   M0     M  [        UUU 4S j[        U5       5       5      nT R                   R"                  " / UQU
Q76 n[%        U[        5      (       d  U4OUn[        [        U5      5       H]  n[%        X   [&        R(                  R*                  5      (       a&  X   R                  R                  R                  5       OX   X[   T'   M_     GMi     / n[        U5       H  u  pT R,                  (       aM  Ub  [        R.                  " X5      nOF[        [        U5      5       H  n[        R.                  " US5      nM     OUc  UR                  5       nUR                  T R1                  UTU	   R                  5      5        M     U$ )Nr   r   rI   c              3  l   >#    U  H)  u  pTR                  UT   TU   R                  5      v   M+     g 7fr   r  r   ).0iirl  r   input_indexr   s      r   	<genexpr>+ReduceOps.generic_reduce.<locals>.<genexpr>z  s1     s]rTYTVq~uRy O O]r   14c              3  l   >#    U  H)  u  pTR                  UT   TU   R                  5      v   M+     g 7fr   r  )r  oior   output_indexr   s      r   r  r    s1     !w`vW\WY$..<%)//"R"R`vr  )r|  r  r   r   r@   r  rK   r  r   rN   r  unravel_indexr  r  rO   ri   r  fnr   r   r  r   r  r  r  )r   r   original_axisr|  
input_dataoutput_datainput_shapeoutput_shaper.  r   input_tuplej	acc_tuplecombine_fn_retr  r   _r  r  s   ``               @@r   generic_reduceReduceOps.generic_reducek  s   		ll5$))4t
Ahoo**00"1T*[-CCCcjjoo.rxxJJOO<Q<QRS  z!}))*A**1:K&q.TAXY1GGLs]fgq]rssK4 A%s;/0A3>>3H3H3M3M3R3R3TKN<0 1 "!w`iju`v!ww	!%!3!3!MY!M!M6@QV6W6W^.]k	s;/0AV`!bggnnW6 W69<3F3F3K3K3P3P3R;D<  N<0 1 +"  -GA~~ ,>>$5D"3{#34!~~dA6 5 &yy{JJt~~dE!HNN;< . 
r!   c                   [        U[        5      (       a  US   OUnS nS nU(       aJ  U R                  U" UR                  R                  U R
                  U R                  S9UR                  5      nU(       aN  U R                  U" UR                  R                  U R
                  U R                  S9[        R                  5      nUb  Ub  XE4$ Ub  U$ Ub  U$ [        S5      e)Nr   r|  keepdimsz-val_reduce_op and idx_reduce_op are both None)r   r  r  r   r   r|  r  r   r   r   r   )r   r   val_reduce_opidx_reduce_opr  idxs         r   min_maxReduceOps.min_max  s    &ue44a%..u||/@/@tyy[_[i[i!jlqlwlwxC..u||/@/@tyy[_[i[i!jlnltltuC?s8O_J_JLMMr!   c                    U R                  [        R                  " UR                  R                  U R
                  U R                  S9UR                  5      $ )Nr  )r  rK   r}  r   r   r|  r  r   r   r   s     r   r}  ReduceOps.sum  s<    ~~bffU\\%6%6TYYQUQ_Q_`bgbmbmnnr!   c                $   U R                   [        R                  R                  :X  a0  U R	                  US   [
        R                  [
        R                  S9$ U R                   [        R                  R                  :X  a0  U R	                  US   [
        R                  [
        R                  S9$ U R                   [        R                  R                  :X  a"  U R	                  US   [
        R                  S S9$ U R                   [        R                  R                  :X  a"  U R	                  US   [
        R                  S S9$ U R                   [        R                  R                  :X  a  U R!                  US   5      $ U R#                  U5      $ )Nr   )r  r  )r  r   standard_argmin_combine_tie_break_leftr  rK   minargmin_argmax_combine_tie_break_leftr  argmax_elementwise_maxnanmax_elementwise_minnanmin_sum_combiner}  r  r  s     r   r  ReduceOps.apply_impl  s   ??bkkHHH<<abii<XX__ J JJ<<abii<XX__ < <<<<a		QU<VV__ < <<<<a		QU<VV__ 8 8888E!H%% &&u--r!   )r  r   )r1   r2   r3   r4   rE   r  r  r  r}  r  r:   __classcell__r  s   @r   r  r  [  s)    # )VN$o. .r!   r  c                  @   ^  \ rS rSrU 4S jrS rS rS rS rSr	U =r
$ )ScanOpsi  c                0   > [         TU ]  X5        X0l        g r   )r  rE   reverse)r   r|  r  r  r  s       r   rE   ScanOps.__init__  s    *r!   c                    U R                  [        R                  " UR                  R                  U R
                  S9UR                  S9/$ Nr  rI   )r  rK   cumsumr   r   r|  r   r  s     r   r   ScanOps.cumsum  s8    ryy):):KSXS^S^_``r!   c                    U R                  [        R                  " UR                  R                  U R
                  S9UR                  S9/$ r  )r  rK   cumprodr   r   r|  r   r  s     r   r  ScanOps.cumprod  s8    rzz%,,*;*;$))LTYT_T_`aar!   c           	       ^ ^^^ / n/ nTS   R                   R                  R                  nT Hi  nUR                  UR                   R                  5        UR                  [        R
                  " XER                   R                  R                  S95        Mk     [        US   R                  5       GH  n[        R                  " Xd5      m[        UUU 4S j[        U5       5       5      nTT R                     S:X  aH  [        [        U5      5       H.  nXx   R                   R                  R                  5       X8   T'   M0     M  [        UU 4S j[        [        T5      5       5       5      m[        UUU 4S j[        U5       5       5      n	T R                  R                   " / U	QUQ76 n
[#        U
[        5      (       d  U
4OU
n	[        [        U5      5       H]  n[#        X   [$        R&                  R(                  5      (       a&  X   R                   R                  R                  5       OX   X8   T'   M_     GM     / n[        U5       H3  u  pgUR                  T R+                  UTU   R                  5      5        M5     U$ )Nr   rI   c              3  l   >#    U  H)  u  pTR                  UT   TU   R                  5      v   M+     g 7fr   r  )r  r  rl  indexr   r   s      r   r  'ScanOps.generic_scan.<locals>.<genexpr>  s/     fPeur%%)//BBPer  c              3  \   >#    U  H!  oTR                   :X  a  TU   S -
  OTU   v   M#     g7f)r   Nr  )r  r   r  r   s     r   r  r    s-     "kYjTU		>58a<uQx#OYjs   ),c              3  l   >#    U  H)  u  pTR                  UT   TU   R                  5      v   M+     g 7fr   r  )r  r  r  r   
prev_indexr   s      r   r  r    s1     !u^tUZUW$..:b	"P"P^tr  )r   r   r@   r  rK   r  r   rN   r  r  r  r  r|  rO   ri   r  r  r   r   r  r   r  )r   r   r  r  r@   r.  r   r   r  r  r  r  r  r  s   ``          @@r   generic_scanScanOps.generic_scan  s   
a$$**Ccjjoo.rxxZZ__5J5JKL  z!}))*A$$Q.EfPYZdPeffDTYY1$s;/0A,0GNN,?,?,D,D,FKN5) 1 #"kY^_bch_iYj"kk
!!u^ghs^t!uu	!%!3!3!FY!F!F6@QV6W6W^.]k	s;/0AOY!bggnnP6 P6IL,?,?,D,D,I,I,K;D<  N5) 1 +"  -GAJJt~~dE!HNN;< .
r!   c           
        / nU R                   (       af  U H_  nUR                  U R                  [        R                  " UR
                  R                  U R                  S9UR                  5      5        Ma     OUnU R                  [        R                  R                  :X  a  U R                  US   5      nONU R                  [        R                  R                  :X  a  U R                  US   5      nOU R!                  U5      nU R                   (       aK  U HE  n[        R                  " UR
                  R                  U R                  S9UR
                  l        MG     U$ )Nr  r   )r  r  r  rK   flipr   r   r|  r   r  r   r  r  r   _prod_combiner  r  )r   r   	new_inputr.  r  s        r   r  ScanOps.apply_impl  s    	<<  

dii0XZ]ZcZc!de  I??bkk666++il+C__ 9 99,,y|,C ##I.C<<"$''#**//		"J

 
r!   )r  )r1   r2   r3   r4   rE   r   r  r  r  r:   r  r  s   @r   r  r    s#    ab< r!   r  c                     SS jn SS jnU [         l        U[         l        U [         R                  l        U[         R                  l        g )Nc                8    [        XU5      R                  U 5      $ r   )r  r  )r   r|  r  r  rz  s        r   _new_reduce'_patch_reduce_scan.<locals>._new_reduce  s    95;;EBBr!   c                8    [        XU5      R                  U 5      $ r   )r  r  )r   r|  r  r  rz  s        r   	_new_scan%_patch_reduce_scan.<locals>._new_scan   s    t177>>r!   )F)r   reduceassociative_scanr  )r  r  s     r   _patch_reduce_scanr    s5    C? BI#B BGGN(BGGr!   c                    S nS	S jnS
S jnS nX l         X l        X0l        [        U l        XR
                  l        [        USS9U l        [        USS9U l	        [        USS9U l
        [        5         g )Nc                `   U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR	                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S	:X  a  UR                  5       $ U R                   S
:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR                  5       $ U R                   S:X  a  UR!                  5       $ U R                   S:X  a  UR#                  5       $ [%        SU  S35      e)Nvoidr   r   r   r   r   r   r   r   rS   r|   r~   r   fp16bf16fp32fp64zfail to convert z to ir type)r  get_void_tyr  r  r  r  r"  r&  r)  r,  r/  r=  r2  r6  r  r
  r  r  r   )r   r   s     r   
_new_to_ir$_patch_lang_core.<locals>._new_to_ir  s   99&&((YY& &&((YY& &&((YY'!''))YY'!''))YY("((**YY'!''))YY("((**YY'!''))YY("((**YY'!''))YY)#))++YY*$**,,YY& &&((YY& &&((YY& ''))YY& ((**+D6=>>r!   c                6    Uc  SnUc  SU pTOXpT[        XEU5      $ )Nr   r   )rN   )arg1arg2steprz  ru  ends         r   
_new_range$_patch_lang_core.<locals>._new_range3  s*    <D<D33U&&r!   c                     U (       d   U5       eg r   r0   )r4  r  s     r   _new_static_assert,_patch_lang_core.<locals>._new_static_assert<  s    Str!   c                   [        U [        R                  5      (       d  U $ [        U[        [        45      (       d  U/OUnU Vs/ s H0  n[        U[        R
                  5      (       a  UR                  OUPM2     nn[        U5      [        S[        U R                  5      5      :w  a  [        SU 35      eU R                  R                  X!5        U $ s  snf )Nr   z$len(values) != len(input.shape) for )r   r   r   r  r  	constexprr-   rO   r  r@   r   r   r.   )r   r  r  r|  s       r   	_set_attr#_patch_lang_core.<locals>._set_attr?  s    %++L!+FT5M!B!B&IOPAZ2<<88!''a?Pv;#aU[[!122CD6JKKd+	 Qs   7Cztt.divisibility)r  ztt.contiguityztt.constancy)NN) )rN   static_rangestatic_assertr  static_printr   to_irr   multiple_ofmax_contiguousmax_constancyr  )langr%  r,  r/  r3  s        r   _patch_lang_corer>  	  sp    $?P'
 J"+D!JJy/@AD!)/BD @Dr!   c                H   U R                   R                  5        VVs/ s H@  u  p[        R                  " U5      (       d  M"  U[        [        R
                  4;   d  M>  UPMB     nnn[        U5      S:  d   S5       eU Hq  n[        U[        5        [        UR                  [        5        U[        :X  a  [        UR                  [        5        [        UR                  5        [        U5        Ms     [        [        R
                  R                  [        5        g s  snnf )Nr   z:triton.language must be visible from within jit'd function)__globals__rx  r  ismoduler   r  rO   r  interpreter_builderr   r   r  r>  tensor_descriptor_base)r  r  r-   langsr=  s        r   _patch_langrE  W  s    #%>>#7#7#9p#9xqW=M=Me=TUY^cegigngnboYoU#9Epu:?XXX?t01t{{$782:499&9:4;;'  277113FG qs   !DDDc                b    [        U S5      (       a  [        U 5      " U6 $ [        U 5      " U5      $ )N_fields)r$   rw  )r.  contentss     r   _tuple_createrI  d  s-     $+3	#:#:49hSS	(@SSr!   c                   [        U [        5      (       Ga  [        R                  " [        R
                  R                  R                  U 5      S 5      n[        R                  nSU s=::  a  S:  a  O  O[        R                  nOqSU s=::  a  S:  a  O  O[        R                  nOPSU s=::  a  S:  a  O  O[        R                  nO/SU s=::  a  S:  a  O  O[        R                  nO[        SU  35      e[        [        R                  " U /US9U5      n[        R                   " X15      $ [#        U S	5      (       a  [        R                  " [        R
                  R                  R                  U 5      S 5      n[        [        R                  " U R%                  5       /[        R                  S9U5      n[        R                   " X15      $ [        U [&        5      (       a  [)        U [+        [,        U 5      5      $ [        U [.        5      (       a  U R0                   Vs/ s H  n[-        U5      PM     nnU R0                  S
   S:X  d   e[        R2                  " S5      US
'   [5        [7        5       5      nUR9                  [-        U R:                  5      U R<                   Vs/ s H  n[-        U5      PM     snUU R>                   Vs/ s H  n[        R2                  " U5      PM     snU R@                  S9$ U $ s  snf s  snf s  snf )Ni   l        l        l         l            l            zUnsupported integer value rI   data_ptrrh   r   )r?   r@   rA   rC   r  )!r   r   r   	str_to_tytritonruntimejitmangle_typerK   r   r   r   rS   r   r   rL  r   r$   rK  r  rI  map_implicit_cvtr
   rA   r2  r	   r   make_tensor_descriptorr?   r@   rC   rd   )r.  tyr   r   srA   r}  r   s           r   rR  rR  n  sG   #s\\&..,,88=tDS 5 HHEc!E!IIEs"U"HHEc!E!IIE9#?@@bhhuE:B?yy$$sJ\\&..,,88=tDbhh'7ryyI2Nyy$$	C		S#mS"9::	C)	*	*-0[[9[=#[9{{2!###ll1o!"4"67..M#((4KPSPYPY5ZPY1mA6FPY5ZdkEH__<VET =?LLOET<Vfifqfq / s 	s J :
 6[<Vs    K=%L
 Lc                    [        U [        R                  R                  R                  5      (       a  U R
                  $ U $ r   )r   rM  rN  rO  TensorWrapperr?   )ts    r   _unwrap_tensorrY    s-    !V^^''5566vvHr!   c                    [        U[        R                  R                  R                  5      (       a3  [        R                  R                  R	                  XR
                  5      $ U $ r   )r   rM  rN  rO  rW  r   )rX  original_tensors     r   _rewrap_tensorr\    sE    /6>>#5#5#C#CDD~~!!//3H3HIIHr!   c                  ,    \ rS rSrS rS rS rS rSrg)GridExecutori  c                   SSK Jn  Xl        X l        X0l        UR
                  R                  5        VVs0 s H  u  pVXT" U5      _M     nnnU Vs/ s H  oWR                  U5      S:X  d  M  UPM     snU l        g s  snnf s  snf )Nr   )_normalize_tyr2  )	rO  r`  r  	arg_namesgridr6   rx  get
constexprs)r   r  ra  rb  r`  r  rT  r6   s           r   rE   GridExecutor.__init__  su    &"	CECUCUC[C[C]^C]xt4r!22C]^,5bID9L9LT9RVa9a4Ib _bs   A<B-Bc                   ^^	 0 m	UU	4S jmU Vs/ s H  nT" U5      PM     nn0 nUR                  5        H  u  pgT" U5      XV'   M     XE4$ s  snf )Nc                  > [        U [        5      (       a  [        U [        TU 5      5      $ [        U [        5      (       aG  [	        T" U R
                  5      U R                  U R                  U R                  U R                  5      $ [        U S5      (       d  U $ [        U 5      nUR                  5       R                  5       T;  a1  UR                  5       nUR                  5       TUR                  5       '   TUR                  5       R                  5          nUR                  SSS9nUR!                  X!R#                  5       UR%                  5       UR'                  5       5        [)        X0S9nU$ )NrK  r   cpu)device)r[  )r   r  rI  rQ  r
   r?   r@   rA   rC   rd   r$   rY  untyped_storagerK  rh  	new_emptyset_storage_offsetr  rj   r\  )r.  unwrapped_argstoragecpu_arg_to_cpustoragess       r   rq  ,GridExecutor._init_args_hst.<locals>._to_cpu  s2   #u%%$S#gs*;<<C!122'CHH%IIKKOOKK  S*--
*3/M,,.779I'779/6{{}))+,}<<>GGIJG#--a->GLL">">"@-BTBTBVXeXlXlXno$WBGNr!   rw  )
r   args_devrz  r.  args_hst
kwargs_hstr,   r-   rq  rr  s
           @@r   _init_args_hstGridExecutor._init_args_hst  sY    	2 -55HSGCLH5 
 ,,.JC%enJO )## 6s   Ac                   ^
^ 0 mU
U4S jm
[        X5       H  u  pVT
" XV5        M     UR                  5        H  u  pxXG   n	T
" X5        M     TR                  5        H  u  pVUR                  U5        M     g )Nc                  > [        U S5      (       aU  [        U 5      [        U5      pU R                  5       UR                  5       4TU R                  5       R                  5       '   g [	        U [
        5      (       a  [        X5       H  u  pT" X5        M     g [	        U [        5      (       a  T" U R                  UR                  5        g g )NrK  )	r$   rY  rj  rK  r   r  zipr
   r?   )arg_devarg_hst	_from_cpurr  s     r   r~  1GridExecutor._restore_args_dev.<locals>._from_cpu  s    w
++#1'#:N7<SBIBYBYB[]d]t]t]vAw002;;=>GU++*-g*?&Wg/ +@G%566',,5 7r!   )r{  rx  r  copy_)r   rt  ru  rz  rv  r|  r}  r,   	kwarg_dev	kwarg_hstr~  rr  s             @@r   _restore_args_devGridExecutor._restore_args_dev  sn    		6 !$H 7Gg' !8 %llnNC"Ii+ - #+//"3WMM'" #4r!   c                   UR                  SS5      (       a  g [        R                  " U R                  5      nUR	                  5        VVs0 s H  u  pEXCR
                  ;   d  M  XE_M     nnnU R                  X5      u  pg[        U R                  5        [        R                  " U R                  /UQ70 UD6nUR	                  5        V	V
s0 s H"  u  pXU R                  ;   a  U
O
[        U
5      _M$     nn	n
[        U R                  5      (       a  U R                  U5      OU R                  n[        U5      S::  d   S5       eUSS[        U5      -
  -  -   n[        R                  " U6    [!        US   5       HU  n[!        US   5       H@  n[!        US   5       H+  n[        R#                  XU5        U R                  " S	0 UD6  M-     MB     MW     U R3                  XX'5        g s  snnf s  sn
n	f ! [$         aD  n[&        R(                  R*                  R,                  (       a  e [/        [1        U5      5      UeS nAff = f)
NwarmupF   z#grid must have at most 3 dimensionsr  r   r   r   r0   )popr  getfullargspecr  rx  ry  rw  rE  getcallargsrd  rR  callablerb  rO   rB  r  rN   r   	ExceptionrM  knobscompilationfront_end_debuggingr   r  r  )r   rt  rz  argspecr{  r|  ru  rv  ry  r  r.  rb  r   r   r   es                   r   __call__GridExecutor.__call__  s   ::h&& ((1#)<<>G>41Q,,5F$!$>G#228DDGG ""477DXDD^b^h^h^jk^jQZQUT__4c-:LL^jk"*499"5"5tyy4994yA~DDD~eq3t9}--(($/		347^tAwA"47^+88qA$ , ( $ 	x6F3 H l  	3||'';;"47+2	3s+   G(%G()G..A'G4 4
I>?H==I)ra  rd  r  rb  N)	r1   r2   r3   r4   rE   rw  r  r  r:   r0   r!   r   r^  r^    s    c"$H#2 Gr!   r^  c                      \ rS rSrS rSrg)ASTTransformeri  c           	        / nUR                    H  nX R                  U5      /-  nM     [        U5      S:  a  [        S5      e[        R
                  " [        R                  " [        R                  " S[        R                  " 5       S9S[        R                  " 5       S9UR                  [        R                  " SS9// S	9Ul	        U$ )
Nr   z&Multiple assignments are not supportedinterpreter_semantic)idctxr  )r-   r   r  F)r-   )funcry  keywords)targetsvisitrO   r   astCall	AttributeNameLoadr-   Constant)r   nodenamestargets       r   visit_AssignASTTransformer.visit_Assign  s    llFjj())E #u:>EFF XXSXX1GSXXZ%X_j#&88:/6:jj#,,UZB[5\gik
 r!   r0   N)r1   r2   r3   r4   r  r:   r0   r!   r   r  r    s    r!   r  c                  L    \ rS rSr\" 5       rS rS rS rS r	S r
S rS rS	rg
)FunctionRewriteri  c                8    Xl         X l        SU l        SU l        g )Nr5  r   )r  rz  filenamedef_file_lineno)r   r  rz  s      r   rE   FunctionRewriter.__init__  s    $%r!   c                L    [         R                  " U R                  5      u  pU R	                  5       u  U l        U l        U R                  U5      U l        U R                  U5      nU R                  U5      nU R                  U5      $ ! [         a    U R                  s $ f = fr   )r  getsourcelinesr  r  _get_jit_fn_file_liner  r  	_find_def
def_lineno_prepare_source_transform_ast_compile_and_exec)r   linesr  r  transformed_asts        r   rewrite_astFunctionRewriter.rewrite_ast%  s    	--dgg6HE /3.H.H.J+t+../""5)--c2%%o66  	77N	s   "B
 
B#"B#c                B    SSK JnJn  U" U" U R                  5      5      $ )Nr   )get_jit_fn_file_lineJITFunction)rO  r  r  r  )r   r  r  s      r   r  &FunctionRewriter._get_jit_fn_file_line9  s    :#K$899r!   c                    Sn[        U5       H0  u  p4UR                  5       R                  S5      (       d  M+  US-   nM2     U$ )Nr   zdef r   )r  strip
startswith)r   r  r  r   lines        r   r  FunctionRewriter._find_def=  s@    
 'GAzz|&&v..U
 ( r!   c                r    XR                   S-
  S  nSR                  U5      n[        R                  " U5      $ )Nr   r5  )r  jointextwrapdedent)r   r  r  s      r   r   FunctionRewriter._prepare_sourceE  s2    oo)*+ggens##r!   c                    [         R                  " U5      nU R                  R                  U5      n[         R                  " U5        U R
                  S-
  n[         R                  " X45        U$ r_  )r  parseast_transformerr  fix_missing_locationsr  increment_lineno)r   r  
parsed_astr  
inc_linenos        r   r  FunctionRewriter._transform_astJ  sY     YYs^
..44Z@!!/2))A-
_9r!   c                   [        XR                  SS9n0 U R                  EnU R                  R                  n[        5       R                  5        H  u  pVXT;  d  M  XdU'   M     [        X$U5        X0R                  R                     $ )Nexec)r  mode)	compiler  rz  r  r@  globalsrx  r  r1   )r   r  compiled_codelocal_namespace
fn_globalsr,   r-   s          r   r  "FunctionRewriter._compile_and_execU  so    --fU)T[[/WW((
!)//+JC$"'3 , 	]8ww//00r!   )r  r  r  r  rz  N)r1   r2   r3   r4   r  r  rE   r  r  r  r  r  r  r:   r0   r!   r   r  r    s-    $&O&7(:$
	1r!   r  c                  P    \ rS rSr% 0 rS\S'   SS jrS r\S 5       r S r	S r
S	rg
)InterpretedFunctioni`  zDict[Callable, Callable]rewritten_fnc                  ^  UT l         [        U40 UD6T l        UT l        U 4S jnUT l        [
        R                  " U5      nUR                  R                  5        Vs/ s H  oUR                  PM     snT l
        g s  snf )Nc                 h   > US   nTR                  5       n[        UTR                  U5      " U 0 UD6$ )Nrb  rewriter^  ra  )ry  rz  rb  r  r   s       r   run)InterpretedFunction.__init__.<locals>.runi  s4    &>DBDNND94J6JJr!   )r  r  rewriterrz  r  r  	signature
parametersr  r  ra  )r   r  rz  r  r  r|  s   `     r   rE   InterpretedFunction.__init__d  sm    (6v6	K
 %%b)	*3*>*>*E*E*GH*GQ&&*GHHs    A>c                    U R                   U R                  ;  a1  U R                  R                  5       U R                  U R                   '   U R                  U R                      $ r   )r  r  r  r  r   s    r   r  InterpretedFunction.rewriter  sJ    77$+++)-)B)B)DDdgg&  ))r!   c                .    U R                   R                  $ r   )r  r1   r   s    r   r1   InterpretedFunction.__name__w  s    wwr!   c                N    U R                  5       n[        X R                  U5      $ r   r  )r   rb  r  s      r   __getitem__InterpretedFunction.__getitem__{  s    \\^B55r!   c                    [        U R                  5        U R                  5       n U" U0 UD6$ ! [         a  n[	        [        U5      5      UeS nAff = fr   )rE  r  r  r  r   r  )r   ry  rz  r  r  s        r   r  InterpretedFunction.__call__  sO    DGG\\^	3t&v&& 	3"47+2	3s   / 
AAA)ra  r  rz  r  r  Nr  )r1   r2   r3   r4   r  r6   rE   r  r  r  r  r:   r0   r!   r   r  r  `  s6    -/L*/I*
    63r!   r  )F
__future__r   r  r  r  typingr   r   r   r   r   numpyrK   rM  triton.languagelanguager   r7   r   triton.language.semanticr	   triton.tools.tensor_descriptorr
   errorsr   	functoolsr   _C.libtritonr   r  r   r   r   r<   r`   ru   r   r   r   r   r   	vectorizer   rZ  r   r[  rS   r"  r   r   r  r  r  r  r  r  r  r>  rE  rI  rR  rB  r  rY  r\  r^  NodeTransformerr  r  r  r0   r!   r   <module>r     s   " 
   . .      ! 3 ; $  6 $   6 4% %P $
& 
& 
&	@='@
# ll45ll45Z<p pJ8 J8Z#4(.I ID].% ].@;# ;|) K\
HTD )* %&9: hG hGVS((  B1 B1J&3 &3r!   