
    h                       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  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'\RP                  " \&\RR                  /S9r*\RP                  " \&\RV                  /S9r,\RP                  " \'\RZ                  /S9r. " S S5      r/ " S S 5      r0S! r1S" r2S# r3 " S$ S%5      r4 " S& S'\45      r5 " S( S)\45      r6S* r7S+ r8S, r9S- r:S. r;\0" 5       r<\" \<5      r=S/ r>S0 r? " S1 S25      r@ " S3 S4\R                  5      rB " S5 S65      rC " S7 S85      rDg)9    )annotationsN)TupleListDict)	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    T/var/www/fran/franai/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__r/   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?   r@   rA   rB   rC   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   rB   onesr   rangelenrA   arangereshaper@   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>   rB   rA   rC   r?   r@   N)r0   r1   r2   r3   rD   r\   r9   r/   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                T    Xl         [        U5      U l        X l        X0l        X@l        g r   )r>   rN   ndimr?   r@   rB   )r   r>   r?   r@   rB   s        r   rD   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 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 alignedzstride must be 16-byte alignedr
   zlast dim must be contiguous)r>   r   itemrN   r@   rb   rB   )r   strides     r   validateTensorDescHandle.validateW   s    yy~~""$r)Q.N0NN.4<< DII---4##$		111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$ )NrG   rf   re   r   z*block offset start must be 16-byte alignedrH   r
   )rN   rb   r>   r   r"   rI   r   rJ   rK   rB   rL   r   rM   rO   rP   r@   rQ   rR   r?   r   rS   )	r   rA   	scalar_tyitemsizerW   rX   rY   rZ   r[   s	            r   r\   %TensorDescHandle.materialize_pointers`   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>   rB   rb   r?   r@   N)r>   r   r?   List[TensorHandle]r@   ro   rB   	List[int])rA   ro   )r0   r1   r2   r3   rD   ri   r\   r9   r/   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)InterpreterOptionss   Nr8   extern_libsFr   debugTsanitize_overflowstrarch)fp8e5fp8e5b16fp8e4nvfp8e4b8fp8e4b15z
Tuple[str]supported_fp8_dtypesr/   !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)r0   r1   r2   r3   ru   r5   rv   rw   ry   r   r   r   r   r   r   r9   r/   r    r   rs   rs   s   sl    KE4"t"D#'^*^46%z6'--/I *I)*!3*%L#%r    rs   c                &   U [         R                  :X  a  [         R                  $ U [         R                  :X  a  [         R                  $ U [         R
                  :X  a  [         R                  $ U [         R                  :X  a  [         R                  $ U $ r   )	rJ   uint8int8uint16int16uint32int32rR   int64rH   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_typerJ   r   rR   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uintrH   r
   r   )getattrrJ   rI   
frombuffertobytesfp_mantissa_widthexponent_biasrQ   r   any
zeros_likerM   maximumminimum_irROUNDING_MODERTNErP   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    r/   N)r0   r1   r2   r3   staticmethodr   r9   r/   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 r          SS 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
   r/   )lhsTyperhsTypes     r   <lambda>-InterpreterBuilder.__init__.<locals>.<lambda>  s    Ir    min_dot_size)ry   rs   optionscodegen_fnsr   r   r   s    r   rD   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_ty.      zzr    c                "    [         R                  $ r   )r   r   r   s    r   get_bf16_tyInterpreterBuilder.get_bf16_ty1      {{r    c                "    [         R                  $ r   )r   r   r   s    r   get_float_tyInterpreterBuilder.get_float_ty4  r  r    c                "    [         R                  $ r   )r   r   r   s    r   get_double_ty InterpreterBuilder.get_double_ty7  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_ty@      xxr    c                "    [         R                  $ r   )r   r   r   s    r   get_int16_tyInterpreterBuilder.get_int16_tyC  r  r    c                "    [         R                  $ r   )r   r   r   s    r   get_uint16_ty InterpreterBuilder.get_uint16_tyF      yyr    c                "    [         R                  $ r   )r   r   r   s    r   get_int32_tyInterpreterBuilder.get_int32_tyI  r  r    c                "    [         R                  $ r   )r   r   r   s    r   get_uint32_ty InterpreterBuilder.get_uint32_tyL  r"  r    c                "    [         R                  $ r   )r   r   r   s    r   get_int64_tyInterpreterBuilder.get_int64_tyO  r  r    c                "    [         R                  $ r   )r   rR   r   s    r   get_uint64_ty InterpreterBuilder.get_uint64_tyR  r"  r    c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e4nv_ty!InterpreterBuilder.get_fp8e4nv_tyU      }}r    c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e4b15_ty"InterpreterBuilder.get_fp8e4b15_tyX      ~~r    c                "    [         R                  $ r   )r   r   r   s    r   get_fp8e4b8_ty!InterpreterBuilder.get_fp8e4b8_ty[  r2  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_tya  r6  r    c                .    [         R                  " X5      $ r   )r   r   )r   elt_ty
addr_spaces      r   
get_ptr_tyInterpreterBuilder.get_ptr_tyd  s    v22r    c                .    [         R                  " X5      $ r   )r   r   )r   r   r?   s      r   get_block_tyInterpreterBuilder.get_block_tyg  s    }}U**r    c                z    [        [        R                  " U/[        R                  S9[        R
                  5      $ NrH   )r   rJ   arraybool_r   r   r   r,   s     r   get_int1InterpreterBuilder.get_int1j  s$    BHHeWBHH=rwwGGr    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   	get_uint8InterpreterBuilder.get_uint8m  $    BHHeWBHH=rxxHHr    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   get_int8InterpreterBuilder.get_int8p  s$    BHHeWBGG<bggFFr    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   
get_uint16InterpreterBuilder.get_uint16s  $    BHHeWBII>		JJr    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   	get_int16InterpreterBuilder.get_int16v  rR  r    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   
get_uint32InterpreterBuilder.get_uint32y  rY  r    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   	get_int32InterpreterBuilder.get_int32|  rR  r    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  rR   r   rL  s     r   
get_uint64InterpreterBuilder.get_uint64  rY  r    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   	get_int64InterpreterBuilder.get_int64  rR  r    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   get_fp16InterpreterBuilder.get_fp16  $    BHHeWBJJ?LLr    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   get_fp32InterpreterBuilder.get_fp32  rl  r    c                z    [        [        R                  " U/[        R                  S9[        R                  5      $ rI  )r   rJ   rJ  r   r   rL  s     r   get_fp64InterpreterBuilder.get_fp64  rl  r    c                T    [        [        R                  " S/[        U5      S9U5      $ Nr   rH   )r   rJ   rJ  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 NonerH   )r   r   r   rJ   rJ  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      $ rI  )r   rJ   rJ  r   r   r   ry  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      $ rI  )r   rJ   	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      $ rI  )r   rJ   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      $ rI  )r$   r   r   rJ   r   r   _interpreterload)
r   rW   r  r  cache_modifiereviction_policyr  rU   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   rW   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   rS   r   r   r   r   r   viewr   r   rQ   )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   rS   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   rS   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   rS   )r   lhsrhsops       r   	binary_opInterpreterBuilder.binary_op  s(    Bsxx2CII4D4DEEr    c                B    U R                  X[        R                  5      $ r   r  rJ   addr   r  r  s      r   r   r    s    "&&)Ir    c                B    U R                  X[        R                  5      $ r   r  rJ   multiplyr  s      r   r   r        "++)Nr    c                B    U R                  X[        R                  5      $ r   r  rJ   divider  s      r   r   r    s    ")))Lr    c                B    U R                  X[        R                  5      $ r   r  rJ   fmodr  s      r   r   r        "'')Jr    c                B    U R                  X[        R                  5      $ r   r  rJ   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  rJ   
left_shiftr  s      r   r   r    s    s(Or    c                B    U R                  X[        R                  5      $ r   )r  rJ   right_shiftr  s      r   r   r    s    "..)Qr    c                B    U R                  X[        R                  5      $ r   r  rJ   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  rJ   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  rJ   
less_equalr  s      r   r   r        DNN3R]],Sr    c                B    U R                  X[        R                  5      $ r   r  rJ   lessr  s      r   r   r        DNN3RWW,Mr    c                B    U R                  X[        R                  5      $ r   r  rJ   greater_equalr  s      r   r   r        DNN3REUEU,Vr    c                B    U R                  X[        R                  5      $ r   r  rJ   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  rJ   equalr  s      r   r   r    s    4>>#BHH+Mr    c                B    U R                  X[        R                  5      $ r   r  rJ   	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  rJ   bitwise_andr  s      r   r   r        s(Pr    c                B    U R                  X[        R                  5      $ r   )r  rJ   bitwise_xorr  s      r   r   r    r  r    c                B    U R                  X[        R                  5      $ r   )r  rJ   
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   rJ   r  r   rS   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   rQ   r  rJ   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   rG   r   )r   r   rJ   r   rR   r   np_umulhi_u64rS   r   rm   rQ   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(  rJ   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(  rJ   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   rI   r   rJ   r   r  r   r   rS   )r   r,  rU   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   )r7  rJ   cosr   r,  s     r   r   r  4      4==bff#=r    c                B    U R                  U[        R                  5      $ r   )r7  rJ   expr@  s     r   r   r  5  rA  r    c                B    U R                  U[        R                  5      $ r   )r7  rJ   exp2r@  s     r   r   r  6      DMM#rww$?r    c                B    U R                  U[        R                  5      $ r   )r7  rJ   absr@  s     r   r   r  7  s    DMM#rvv$>r    c                B    U R                  U[        R                  5      $ r   )r7  rJ   floorr@  s     r   r   r  8  s    T]]3%Ar    c                B    U R                  U[        R                  5      $ r   )r7  rJ   ceilr@  s     r   r   r  9  rF  r    c                B    U R                  U[        R                  5      $ r   )r7  rJ   logr@  s     r   r   r  :  rA  r    c                B    U R                  U[        R                  5      $ r   )r7  rJ   log2r@  s     r   r   r  ;  rF  r    c                B    U R                  U[        R                  5      $ r   r7  rJ   sqrtr@  s     r   r   r  <  s    DMM#rww,Gr    c                B    U R                  U[        R                  5      $ r   rR  r@  s     r   r   r  =  rF  r    c                B    U R                  U[        R                  5      $ r   )r7  rJ   sinr@  s     r   r   r  >  rA  r    c                    UR                   R                  [        R                  :X  a  [	        UR                   5      O[        UR                   5      n[        X!R                  R                  5      $ r   )r   r   rJ   r   np_erf_fp32np_erf_fp64r   rS   )r   r,  r  s      r   
create_erfInterpreterBuilder.create_erf@  sF    '*xx~~'Ck#((#UXU]U]I^C!1!122r    c                    [        S[        R                  " UR                  5      -  UR                  R
                  5      $ Nr
   )r   rJ   rS  r   r   rS   r@  s     r   create_rsqrtInterpreterBuilder.create_rsqrtD  s+    A 113993C3CDDr    c                t    [        UR                  R                  U5      UR                  R                  5      $ r   )r   r   rP   r   rS   )r   r,  r?   allow_reorders       r   r   r  H  s(    \#((JZJZ[`Jacfclclcscs=tr    c                    [        [        R                  " UR                  U5      UR                  R
                  5      $ r   )r   rJ   	transposer   r   rS   )r   r,  perms      r   create_transInterpreterBuilder.create_transJ  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      $ )NrG   rH   )r   r   rI   is_floatingr   r   r   r  rJ   r   matmulrS   )r   r   r   dinput_precisionmax_num_imprecise_acca_datab_datas           r   
create_dotInterpreterBuilder.create_dotM  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      $ rI  )r   rJ   rO   r   r   )r   ret_tystartstops       r   create_make_range$InterpreterBuilder.create_make_rangeV  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      $ )NrH   r   )binsrM   )r   rJ   r  r   r   r   r   r1  r   	histogramlogical_notsumr   )r   r   rx  r  ry  s        r   create_histogram#InterpreterBuilder.create_histogramY  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rz  )r   rJ   take_along_axisr   r   rS   )r   r  indicesrz  s       r   create_gather InterpreterBuilder.create_gatherc  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
   rG   )	r$   rI   maxr   r   rQ   rJ   rR   r   )r   r  offsetrU   element_bitwidthelement_bytewidths         r   create_addptr InterpreterBuilder.create_addptrh  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      $ )NrH   nanzunsupported padding option )r\   r$   r   r   PADDING_OPTIONPAD_ZEROr   rJ   r   r   PAD_NAN	full_likefloatr   r  )r   r  rT   padding_optionr  r  r  rW   rX   rU   r  r  s               r   create_tensor_pointer_load-InterpreterBuilder.create_tensor_pointer_loado  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,   rT   r  r  rW   rX   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   rJ   expand_dimsr   r   rS   )r   r,  rz  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   rJ   rK   r   r   rS   )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   rJ   concatenater   r   rS   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      $ )Nrf   r  )r   rJ   stackr   r   rS   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      $ rt  )r?   r   r   r   r   r   rJ   fullr   r   rS   )r   rr  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                   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   rS   )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   rS   )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valuesru  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02xr/   r   s    r   r   1InterpreterBuilder.create_print.<locals>.<lambda>  s    b3Lr    )	formatter)r   rJ   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   r/   )r   	conditionmessages      r   create_assert InterpreterBuilder.create_assert  s    &WI&yr    c                     U(       d   S5       eg )NzAssume failedr/   )r   r  s     r   create_assume InterpreterBuilder.create_assume  s    )/)yr    c                    g r   r/   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?   r@   rA   rB   rC   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))rN   rA   r   r(   r;   r>   r?   r@   rB   rC   rM   r   )r   r  rA   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X45      nUR                  5         U$ r   )r_   ri   )r   r>   r?   r@   tensor_shaper  descs          r   create_make_tensor_descriptor0InterpreterBuilder.create_make_tensor_descriptor  s      WCr    c           	     |    [        U[        5      (       d   eUR                  U5      u  pVU R                  XVS UUSS9$ )NF)r  r  r  r  )r   r_   r\   r  )r   r  r  r  r  rW   r  s          r   create_descriptor_load)InterpreterBuilder.create_descriptor_load  sN    $ 01111..w7
&&tn7FTY ' [ 	[r    c                P    UR                  U5      u  pEU R                  XBUS S 5      $ r   r  )r   r  r,   r  rW   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   rf   rH   )r>   r   r"   r   rJ   zerosr   r?   rB   	enumerater   r   r   r  )r   r  	x_offsetsy_offsetru  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
   rf   rH   Tzunsupported type )r   namer   rJ   r  rS   rK  	TypeError)r   ru  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    )ry   r   r   r   r   NreturnNone)
r>   r   r?   ro   r@   ro   r  rp   r  r   )r  r_   r  ro   )r  r_   r,   r   r  ro   )r  r_   r  r   r  r   )r  r_   r,   r   r  r   r  r   )r0   r1   r2   r3   r   MEM_SEMANTICACQUIREr  RELEASERELAXEDACQUIRE_RELEASEr  	ATOMIC_OPADDRMW_OPFADDMINUMINMAXUMAXANDORXORXCHGr  rD   r   r  r  r  r  r  r  r  r  r  r   r$  r'  r*  r-  r0  r4  r8  r;  r>  rC  rF  rM  rP  rT  rW  r[  r^  ra  rd  rg  rj  rn  rq  rv  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_selectr4  r7  r<  
create_cos
create_expcreate_exp2create_iabscreate_floorcreate_ceil
create_logcreate_log2create_precise_sqrtcreate_sqrt
create_sinrZ  r^  create_reshapere  ro  ru  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  r9   r/   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jrZR0'*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)rn  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  rn  r   
new_memberrv  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  rn  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      $ )Nrf   )r   rJ   rc  r   r   r   ru  is_blocklistr?   r   r  r   r   )r   r   rB   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>*  s    C(8(8$9r    c                   > T" U 5      $ r   r/   )r   r  s    r   r   r  +  s	    9T?r    c                @    [        U R                  R                  5      $ r   )reprr   r   r   s    r   r   r  ,  s    4(8(8#9r    c                @    [        U R                  R                  5      $ r   )rx   r   r   r   s    r   r   r  -  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)	ReduceScanOpInterfacei1  c                    Xl         X l        g r   rz  
combine_fn)r   rz  r  s      r   rD   ReduceScanOpInterface.__init__3  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 )rN   r   )r   r?   rz  s      r   
check_axis ReduceScanOpInterface.check_axis7  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   ru  r  r?   rz  )r   r   r,  s      r   check_tensor"ReduceScanOpInterface.check_tensor;  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?   rH   )r   r#   r?   rQ   r   r   r  rJ   rJ  r  r   r   rS   )r   r  r   r  ret_types        r   	to_tensorReduceScanOpInterface.to_tensorA  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.applyK  sb    %''::ui(++% ooe$'dE];;uSzH#Hr    r  N)
r0   r1   r2   r3   rD   r  r  r  r  r9   r/   r    r   r  r  1  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$ )
	ReduceOpsiS  c                0   > [         TU ]  X5        X0l        g r   )superrD   	keep_dims)r   rz  r  r  	__class__s       r   rD   ReduceOps.__init__U  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   rz  r  r   s        r   unravelReduceOps.unravelY  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
   rH   c              3  l   >#    U  H)  u  pTR                  UT   TU   R                  5      v   M+     g 7fr   r  r   ).0iirj  r   input_indexr   s      r   	<genexpr>+ReduceOps.generic_reduce.<locals>.<genexpr>r  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  x  s1     !w`vW\WY$..<%)//"R"R`vr  )rz  r  r   r   r?   r  rJ   r  r   rM   r  unravel_indexr  r   rN   rg   r  fnr   r   r  r   r  r  r  )r   r   original_axisrz  
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_reducec  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   rz  keepdimsz-val_reduce_op and idx_reduce_op are both None)r   r  r  r   r   rz  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  rJ   r{  r   r   rz  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  rJ   minargmin_argmax_combine_tie_break_leftr  argmax_elementwise_max_elementwise_min_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d<SS__ < <<<<ad<SS__ 8 8888E!H%% &&u--r    )r  r   )r0   r1   r2   r3   rD   r  r  r  r{  r  r9   __classcell__r  s   @r   r  r  S  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  rD   reverse)r   rz  r  r  r  s       r   rD   ScanOps.__init__  s    *r    c                    U R                  [        R                  " UR                  R                  U R
                  S9UR                  S9/$ Nr  rH   )r  rJ   cumsumr   r   rz  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  rJ   cumprodr   r   rz  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   rH   c              3  l   >#    U  H)  u  pTR                  UT   TU   R                  5      v   M+     g 7fr   r  )r  r  rj  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  rJ   r  r   rM   r  r  r  r   rz  rN   rg   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  rJ   flipr   r   rz  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  )r0   r1   r2   r3   rD   r  r  r  r  r9   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   rz  r  r  rs  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   rz  r  r  rs  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   rR   rz   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;  r0  r4  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   )rM   )arg1arg2steprs  rs  ends         r   
_new_range$_patch_lang_core.<locals>._new_range+  s*    <D<D33U&&r    c                     U (       d   U5       eg r   r/   )r2  r  s     r   _new_static_assert,_patch_lang_core.<locals>._new_static_assert4  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,   rN   r  r?   r   r   r-   )r   r  r  ru  s       r   	_set_attr#_patch_lang_core.<locals>._set_attr7  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) )rM   static_rangestatic_assertr  static_printr   to_irr   multiple_ofmax_contiguousmax_constancyr  )langr  r#  r&  r*  s        r   _patch_lang_corer5    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__rq  r}  ismoduler   r  rN   r  interpreter_builderr   r   r  r5  tensor_descriptor_base)r  r  r,   langsr4  s        r   _patch_langr<  O  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#   ru  )r,  contentss     r   _tuple_creater@  \  s-     $+3	#:#:49hSS	(@SSr    c                   [        U [        5      (       Ga  [        R                  " [        R
                  R                  R                  U 5      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      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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 rH   data_ptrrf   r
   )r>   r?   r@   rB   ) r   r   r   	str_to_tytritonruntimejitmangle_typerJ   r   r   r   rR   r   r   rJ  r   r#   rB  r  r@  map_implicit_cvtr	   r@   r)  r   r   make_tensor_descriptorr>   r?   rB   )r,  tyr   r   sr@   rv  r   s           r   rI  rI  f  s&   #s\\&..,,88=>S 5 HHEc!E!IIEs"U"HHEc!E!IIE9#?@@bhhuE:B?yy$$sJ\\&..,,88=>bhh'7ryyI2Nyy$$	C		S#mS"9::	C)	*	*-0[[9[=#[9{{2!###ll1o!"4"67..sxx(-0YY7Y=#Y725//B/Qa/B	 / 
 	
 J : 8Bs   K0#K5
	 K:c                    [        U [        R                  R                  R                  5      (       a  U R
                  $ U $ r   )r   rD  rE  rF  TensorWrapperr>   )ts    r   _unwrap_tensorrP    s-    !V^^''5566vvHr    c                    [        U[        R                  R                  R                  5      (       a3  [        R                  R                  R	                  XR
                  5      $ U $ r   )r   rD  rE  rF  rN  r   )rO  original_tensors     r   _rewrap_tensorrS    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_tyr)  )	rF  rW  r  	arg_namesgridr5   rq  get
constexprs)r   r  rX  rY  rW  r  rK  r5   s           r   rD   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      (       a<  [	        T" U R
                  5      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$ )NrB  r   cpu)device)rR  )r   r  r@  rH  r	   r>   r?   r@   rB   r#   rP  untyped_storagerB  r_  	new_emptyset_storage_offsetr  rh   rS  )r,  unwrapped_argstoragecpu_arg_to_cpustoragess       r   rh  ,GridExecutor._init_args_hst.<locals>._to_cpu  s+   #u%%$S#gs*;<<C!122'CHH%IIKKOO	  S*--
*3/M,,.779I'779/6{{}))+,}<<>GGIJG#--a->GLL">">"@-BTBTBVXeXlXlXno$WBGNr    rp  )
r   args_devrs  r,  args_hst
kwargs_hstr+   r,   rh  ri  s
           @@r   _init_args_hstGridExecutor._init_args_hst  sY    	0 -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 )NrB  )	r#   rP  ra  rB  r   r  zipr	   r>   )arg_devarg_hst	_from_cpuri  s     r   ru  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    )rr  rq  r  copy_)r   rk  rl  rs  rm  rs  rt  r+   	kwarg_dev	kwarg_hstru  ri  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 dimensions)r
   r   r
   r   r/   )popr}  getfullargspecr  rq  rr  rn  r<  getcallargsr[  rI  callablerY  rN   r9  r  rM   r   	ExceptionrD  knobscompilationfront_end_debuggingr   r  rz  )r   rk  rs  argspecrt  ru  rl  rm  rr  r  r,  rY  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)rX  r[  r  rY  N)	r0   r1   r2   r3   rD   rn  rz  r  r9   r/   r    r   rU  rU    s    c!$F#2 Gr    rU  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,   )funcrr  keywords)targetsvisitrN   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    r/   N)r0   r1   r2   r3   r  r9   r/   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 )Nr,  r   )r  rs  filenamedef_file_lineno)r   r  rs  s      r   rD   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)rF  r  r  r  )r   r  r  s      r   r  &FunctionRewriter._get_jit_fn_file_line2  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_def6  s@    
 'GAzz|&&v..U
 ( r    c                r    XR                   S-
  S  nSR                  U5      n[        R                  " U5      $ )Nr
   r,  )r  jointextwrapdedent)r   r  r  s      r   r   FunctionRewriter._prepare_source>  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_astC  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  rs  r  r7  globalsrq  r  r0   )r   r  compiled_codelocal_namespace
fn_globalsr+   r,   s          r   r  "FunctionRewriter._compile_and_execN  so    --fU)T[[/WW((
!)//+JC$"'3 , 	]8ww//00r    )r  r  r  r  rs  N)r0   r1   r2   r3   r  r  rD   r  r  r  r  r  r  r9   r/   r    r   r  r    s-    $&O&7(:$
	1r    r  c                  D    \ rS rSr0 rS	S jrS r\S 5       r S rS r	Sr
g)
InterpretedFunctioniY  c                   ^  UT l         [        U40 UD6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$ )NrY  rewriterU  rX  )rr  rs  rY  r  r   s       r   run)InterpretedFunction.__init__.<locals>.runa  s4    &>DBDNND94J6JJr    )
r  r  rewriterr  r}  	signature
parametersr  r  rX  )r   r  rs  r  r  ru  s   `     r   rD   InterpretedFunction.__init__]  sf    (6v6	K
 %%b)	*3*>*>*E*E*GH*GQ&&*GHHs   A7c                    U R                   U R                  ;  a1  U R                  R                  5       U R                  U R                   '   U R                  U R                      $ r   )r  rewritten_fnr  r  r   s    r   r  InterpretedFunction.rewritej  sJ    77$+++)-)B)B)DDdgg&  ))r    c                .    U R                   R                  $ r   )r  r0   r   s    r   r0   InterpretedFunction.__name__o  s    wwr    c                N    U R                  5       n[        X R                  U5      $ r   r  )r   rY  r  s      r   __getitem__InterpretedFunction.__getitem__s  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   )r<  r  r  r  r   r  )r   rr  rs  r  r  s        r   r  InterpretedFunction.__call__w  sO    DGG\\^	3t&v&& 	3"47+2	3s   / 
AAA)rX  r  r  r  Nr  )r0   r1   r2   r3   r  rD   r  r  r  r  r9   r/   r    r   r  r  Y  s0    LI*
    63r    r  )E
__future__r   r  r  r}  typingr   r   r   r   numpyrJ   rD  triton.languagelanguager   r6   r   triton.language.semanticr   triton.tools.tensor_descriptorr	   errorsr   	functoolsr   _C.libtritonr   r  r   r   r   r;   r_   rs   r   r   r   r   r   	vectorizer   rX  r   rY  rR   r   r   r   r{  r  r  r  r  r  r  r5  r<  r@  rI  r9  r  rP  rS  rU  NodeTransformerr  r  r  r/   r    r   <module>r     s   " 
   $ $      ! 3 ; $  6 $   6 4# #L $
& 
& 
&	@='@
# ll45ll45Z<p pD8 D8N#4(.I ID].% ].@;# ;|) K\
HT!H )* %&9: gG gGTS((  B1 B1J%3 %3r    