+
    ړLj                      a  0 t $ ^ RIHt ^ RIt^ RIt^ RIt^ RIHtHtH	t	H
t
HtHt ^ RIt^ RIt^ RIt^ RIHt ^ RIt^ RIHt ^ RIHt ^ RIHt ^ RIHt ^RIHt ^ R	IHt ^R
I H!t" ^RI H#t$ ^RI%H&t& ]! R4      t'] ! R R4      4       t( ! R R4      t) ! R R4      t*]! RR7       ! R R4      4       t+R t,R t-R t.R t/R t0R t1]Pd                  ! ]0]Pf                  .R7      t4]Pd                  ! ]0]Pj                  .R7      t6]Pd                  ! ]1]Pn                  .R7      t8 ! R R 4      t9 ! R! R"4      t:];! 4       t<]:! 4       t=]! ]=4      t>R#]?R$&    ! R% R&4      t@R' R( ltAR) R* ltBR+ R, ltC ! R- R.4      tD ! R/ R0]D4      tE ! R1 R2]D4      tFR3 R4 ltGR5 R6 ltHR7 tIR8 tJR9 tKR: tL ! R; R<4      tM ! R= R>]P                  4      tO ! R? R@4      tP ! RA RB]]',          4      tQR# )C    )annotationsN)TupleListDictCallableTypeVarOptional)	dataclass)TritonSemantic)KernelInterface)TensorDescriptor)InterpreterError)partial)interpreter)ir)_tuple_createTc                  ~    ] tR t^t$ RtR]R&   R]R&   ]P                  ! ]R7      t	R]R&   R	 t
R
 tR tR tR tRtR# )TensorHandlez
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.ndarraydataztl.dtypedtype)default_factoryr   attrc                	    \        V P                  V P                  4      '       gN   \        R V P                  P                  ^,           RV P                  P
                   RV P                   24      hR# )znumpy data itemsize (z) bits) exceeds dtype primitive_bitwidth (z bits) for triton type N)_validate_np_data_sizer   r   
ValueErroritemsizeprimitive_bitwidthselfs   &F/app/.local/lib/python3.14/site-packages/triton/runtime/interpreter.py__post_init__TensorHandle.__post_init__'   sm    %dii<<4TYY5G5G!5K4L M!!%!>!> ??VW[WaWaVbd e e =    c                	H    \        V P                  P                  4       4      # N)boolr   allr   s   &r!   __bool__TensorHandle.__bool__,   s    DIIMMO$$r$   c                	^    V P                   p\        VR 4      '       d   VP                  pK   V# )
element_ty)r   hasattrr,   )r    r   s   & r!   get_element_tyTensorHandle.get_element_ty/   s(    

e\**$$Er$   c                	^    \        V P                  P                  4       V P                  4      # r&   )r   r   copyr   r   s   &r!   cloneTensorHandle.clone5   s    DIINN,djj99r$   c                	"    W P                   V&   R # r&   )r   )r    keyvalues   &&&r!   set_attrTensorHandle.set_attr8   s    		#r$    N)__name__
__module____qualname____firstlineno____doc____annotations__dataclassesfielddictr   r"   r)   r.   r2   r7   __static_attributes__r9   r$   r!   r   r      sC     O""48D$8e
%:r$   r   c                  "    ] tR t^<tR tR tRtR# )BlockPointerHandlec                	N    Wn         W n        W0n        W@n        WPn        W`n        R # r&   )baseshapestridesoffsetsblock_shapeorder)r    rG   rH   rI   rJ   rK   rL   s   &&&&&&&r!   __init__BlockPointerHandle.__init__>   s!    	
&
r$   c                	   V P                   P                  4       pVP                  ^,          p\        P                  ! V P                   P
                  V P                  4      p\        P                  ! V P                  \        R7      p\        \        V P                  4      4       EF  p^.\        V P                  4      ,          pV P                  V,          Wv&   V P                  V,          P
                  \        P                  ! V P                  V,          4      ,           P                  V4      pWCV,          V P                  V,          P
                  ,          P                  \        P                   4      ,           pWa9   g   K  WXV P"                  V,          P
                  8  ,          V^ 8  ,          pEK  	  \%        W@P                   P&                  P(                  4      p	W3# )   r   )rG   r.   r   npbroadcast_tor   rK   onesr'   rangelenrJ   arangereshaperI   astypeuint64rH   r   r   scalar)
r    boundary_checkdtype_ttn_bytes	ptrs_datamasksdim
bcast_dimsoffptrs_handles
   &&        r!   materialize_pointers'BlockPointerHandle.materialize_pointersF   sS   99++---2OODIINND4D4DE	((5T--./Cs4#3#344J"..s3JO<<$))BIId6F6Fs6K,LLUUV`aC!s]T\\#5F5K5K%K$S$STVT]T]$^^I$tzz#';';!;<qI 0 #9iioo.D.DE!!r$   )rG   rK   rJ   rL   rH   rI   N)r:   r;   r<   r=   rM   re   rC   r9   r$   r!   rE   rE   <   s    "r$   rE   c                  4    ] tR t^VtR R ltR tR R ltRtR# )TensorDescHandlec               (    V ^8  d   QhRRRRRRRR/# )   rG   r   rH   List[TensorHandle]rI   rK   	List[int]r9   )formats   "r!   __annotate__TensorDescHandle.__annotate__X   s+      \ 2D Oa 'r$   c                	b    Wn         \        V4      V n        W n        W0n        W@n        WPn        R # r&   )rG   rV   ndimrH   rI   rK   padding)r    rG   rH   rI   rK   rr   s   &&&&&&r!   rM   TensorDescHandle.__init__X   s'    	J	
&r$   c                	   V P                   P                  P                  4       ^,          ^ 8X  g   Q R4       h\        V P                  4      V P
                  8X  g   Q h\        V P                  4      V P
                  8X  g   Q hV P
                  ^8  g   Q R4       hV P                   P                  P                  pVP                  ^,          pV P                  RR  F9  pVP                  P                  4       V,          pV^,          ^ 8X  d   K4  Q R4       h	  V P                  R,          P                  P                  4       ^8X  g   Q R4       hR# )   zbase must be 16-byte alignedz"descriptor cannot be 0 dimensionalNzstride must be 16-byte alignedzlast dim must be contiguous)
rG   r   itemrV   rI   rq   rK   r   r,   r   )r    	scalar_tyr   stridebyte_strides   &    r!   validateTensorDescHandle.validatea   s   yy~~""$r)Q.N0NN.4<< DII---4##$		111yyA~CCC~IIOO..	//14ll3B'F ++**,x7K#q(J*JJ( ( ||B$$))+q0O2OO0r$   c                   V ^8  d   QhRR/# )rj   rJ   rk   r9   )rm   s   "r!   rn   ro   n   s     " ",> "r$   c                	   \        V4      V P                  8X  g   Q hV P                  P                  P                  pVP
                  ^,          pVR,          P                  V,          ^,          ^ 8X  g   Q R4       h\        P                  ! V P                  P                  V P                  4      p\        P                  ! V P                  \        R7      p\        \        V P                  4      4       EF  p^.\        V P                  4      ,          pV P                  V,          Wv&   W,          P                  \        P                  ! V P                  V,          4      ,           P                  V4      pWCV,          V P                  V,          P                  ,          P!                  \        P"                  4      ,           pV^ V8*  ,          WP$                  V,          P                  8  ,          pEK  	  VP                  \        P"                  8X  g   Q h\'        W@P                  P                  P(                  4      p	W3# )rP   z*block offset start must be 16-byte alignedrQ   rv   )rV   rq   rG   r   r,   r   r   rR   rS   rK   rT   r'   rU   rW   rX   rI   rY   rZ   rH   r   r[   )
r    rJ   rx   r   r_   r`   ra   rb   rc   rd   s
   &&        r!   re   %TensorDescHandle.materialize_pointersn   s   7|tyy(((IIOO..	//14  8+r1Q6d8dd6OODIINND4D4DE	((5T--./Cs4#3#344J"..s3JO<$$ryy1A1A#1F'GGPPQ[\C!^dll36G6L6L%L$T$TUWU^U^$__IQ#X&#

30D0D*DEE 0 "))+++"9iioo.D.DE!!r$   )rG   rK   rq   rr   rH   rI   N)r:   r;   r<   r=   rM   r{   re   rC   r9   r$   r!   rh   rh   V   s    P" "r$   rh   T)frozenc                      ] tR t^t$ RtR]R&   RtR]R&   RtR]R&   RtR	]R
&   Rt	R]R&   Rt
R]R&   RtR]R&   RtR]R&   ^ tR]R&   RtR]R&   RtR# )InterpreterOptionsNzOptional[dict]extern_libsFr'   debugTsanitize_overflowzOptional[str]archzTuple[str, ...]supported_fp8_dtypes!deprecated_fp8_dot_operand_dtypestf32strdefault_dot_input_precisionallowed_dot_input_precisionsintmax_num_imprecise_acc_defaultr   backend_namer9   )fp8e5fp8e5b16fp8e4nvfp8e4b8fp8e4b15)r   tf32x3ieee)r:   r;   r<   r=   r   r?   r   r   r   r   r   r   r   r   r   rC   r9   r$   r!   r   r      sl    "&K&E4"t"D-,c/c9;%;'--4N /N)*!3*%L#%r$   r   c                    \        V\        P                  4      '       d   R # V P                  ^,          pVP                  pV^8  d   ^pW#8  d   R# R # )TF)
isinstancetlpointer_typer   r   )np_arraytl_dtypenp_dtype_bitwidthtl_dtype_bitwidths   &&  r!   r   r      sO    (BOO,, ))A- 33 1,r$   c                .   V \         P                  8X  d   \         P                  # V \         P                  8X  d   \         P                  # V \         P
                  8X  d   \         P                  # V \         P                  8X  d   \         P                  # V # r&   )	rR   uint8int8uint16int16uint32int32rZ   int64rQ   s   &r!   _get_signed_np_dtyper      s[    ww		xx		xx		xxLr$   c                   \        V \        P                  4      '       d%   \        P                  ! \        P
                  4      # / \        P                  \        P                  ! \        4      b\        P                  \        P                  ! \        P                  4      b\        P                  \        P                  ! \        P                  4      b\        P                  \        P                  ! \        P                  4      b\        P                  \        P                  ! \        P                  4      b\        P                  \        P                  ! \        P                  4      b\        P                  \        P                  ! \        P                  4      b\        P                  \        P                  ! \        P                  4      b\        P                  \        P                  ! \        P                  4      b\        P                   \        P                  ! \        P                   4      b\        P"                  \        P                  ! \        P"                  4      b\        P
                  \        P                  ! \        P
                  4      b\        P$                  \        P                  ! \        P                  4      b\        P&                  \        P                  ! \        P                  4      b\        P(                  \        P                  ! \        P                  4      b\        P*                  \        P                  ! \        P                  4      b\        P,                  \        P                  ! \        P                  4      b\        P.                  \        P                  ! \        P                  4      /Cp\        V \        P0                  4      '       da   \        V P2                  \        P                  4      '       d%   \        P                  ! \        P
                  4      # WP2                  ,          # W,          # r&   )r   r   r   rR   r   rZ   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                
   \        \        R VP                   24      p\        \        R VP                   24      p\        P                  ! V P	                  4       VR7      pWaP                  ^,
          ,	          ^,          pVP                  VP
                  ,
          ^,
          pVP                  VP
                  ,
          ^,
          p	V^VP
                  ,          ^,
          ,          p
VP                  pVP                  pWaP
                  ,	          ^V,          ^,
          ,          P                  \        P                  4      pV^ 8H  p\        P                  ! V4      '       d   \        P                  ! V\        P                  R7      p\        VP
                  4       F,  pV
V,	          ^,          pVP
                  V,
          VV^8H  &   K.  	  V
^ 8H  p^W,          ,
          W&   W,
          VVV,          &   W,          W,          ,          ^VP
                  ,          ^,
          ,          W&   \        P                  ! ^ \        P                  ! W,
          V,           ^V	,          ^,
          4      4      pVP                  V4      pVP                  V4      pVP                  VP                  8  d   WP
                  VP
                  ,
          ,	          ^VP
                  ,          ^,
          ,          pV\        P                  P                   8X  d?   V
^VP
                  VP
                  ,
          ^,
          ,          ,          pVV^ 8  ,           pVP                  V4      pMRV
P                  V4      VP
                  VP
                  ,
          ,          ^VP
                  ,          ^,
          ,          pV^ 8H  p\        P                  ! V4      '       d   WaP
                  ,	          ^V,          ^,
          ,          P                  \        P                  4      pV^ 8g  pVV,          p\        P                  ! V\        P                  R7      p^V,
          W,          V,
          ,
          VV&   VV,          VV,          ,	          ^VP
                  VV,          ,
          ,          ,          VV&   VVP                  ^,
          ,          VVP
                  ,          ,          V,          pVP#                  V P$                  4      # )uintrQ   )getattrrR   r   
frombuffertobytesfp_mantissa_widthexponent_biasrY   r   any
zeros_likerU   maximumminimum_irROUNDING_MODERTNErX   rH   )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                .    \         P                  ! V 4      # r&   )matherfxs   &r!   _erfr     s    88A;r$   c                F    \        V 4      \        V4      ,          ^@,	          # )@   )r   )abs   &&r!   
_umulhi_64r     s     FSVO""r$   )otypesc                  &    ] tR tRt]R 4       tRtR# )ExtraFunctionsi  c                	x    \         P                  ! VP                  P                  V P                  W4      V4      # 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$   r9   N)r:   r;   r<   r=   staticmethodr   rC   r9   r$   r!   r   r     s    p pr$   r   c                     ] tR tRt]P
                  P                  ]P
                  P                  ]P
                  P                  ]P
                  P                  ]P
                  P                  ]P
                  P                  ]P
                  P                  ]P
                  P                  /t]P                  P                  ]P                  P                  ]P                  P                  ]P                  P                  ]P                  P                   ]P                  P                   ]P                  P"                  ]P                  P"                  ]P                  P$                  ]P                  P$                  ]P                  P&                  ]P                  P&                  ]P                  P(                  ]P                  P(                  ]P                  P*                  ]P                  P*                  ]P                  P,                  ]P                  P,                  ]P                  P.                  ]P                  P.                  /
tR R ltR tR tR tR tR tR	 tR
 t R t!R t"R t#R t$R t%R t&R t'R t(R t)R t*R t+R t,R t-R t.R t/R t0R t1R t2R t3R t4R t5R  t6R! t7R" t8R# t9R$ t:R% t;R& t<R' t=R( t>R) t?R* t@R+ tAR, tBR- tCR. tDR/ tER0 tFR1 tGR2 tHR3 tIR4 tJR5 tKR6 tLR7 tMR8 tNR9 tOR: tPR; tQR< tRR= tSR> tTR? tUR@ tVRA tWRB tXRC tYRD tZRE t[RF t\RG t]RH t^RI t_RJ t`RK taRL tbRM tcRN tdRO teRP tfRQ tgRR thRS tiRT tjRU tkRV tlRW tmRX tnRY toRZ tpR[ tqR\ trR] tsR^ ttR_ tuR` tvRa twRb txRc tyRd tzRe t{Rf t|Rg t}]Lt~]LtRh tRi tRj tRk tRl tRm tRn tRo tRp tRq tRr tRs tRt tRu tRv tRw tRx tRy tRz tR{ tR| tR} tR~ tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tRR R lltR R ltR R ltR R ltR R ltR tRtR# )InterpreterBuilderi"  c                   V ^8  d   QhRR/# rj   returnNoner9   )rm   s   "r!   rn   InterpreterBuilder.__annotate__7  s     N N$ Nr$   c                	    R V n         \        4       V n        / V n        \        P
                  V P                  R&   R V P                  R&   R # )Nconvert_custom_typesc                    R# )   )r  r  r  r9   )lhsTyperhsTypes   &&r!   <lambda>-InterpreterBuilder.__init__.<locals>.<lambda><  s    Ir$   min_dot_size)r   r   optionscodegen_fnsr   r   r   s   &r!   rM   InterpreterBuilder.__init__7  sB    	)+3A3W3W/0+M(r$   c                	    WP                   ^ ,          8  g   \        R4      hW P                   ^,          8  g   \        R4      hW0P                   ^,          8  g   \        R4      hWV3V n        R# )r   zx >= grid_dim[0]zy >= grid_dim[1]zz >= grid_dim[2]N)grid_dimr   grid_idxr    r   yzs   &&&&r!   set_grid_idxInterpreterBuilder.set_grid_idx>  s^    ==##/00==##/00==##/00q	r$   c                	    WV3V n         R # r&   )r  )r    nxnynzs   &&&&r!   set_grid_dimInterpreterBuilder.set_grid_dimG  s    r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_half_tyInterpreterBuilder.get_half_tyL      zzr$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_bf16_tyInterpreterBuilder.get_bf16_tyO      {{r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_float_tyInterpreterBuilder.get_float_tyR  r  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_double_ty InterpreterBuilder.get_double_tyU  r  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_int1_tyInterpreterBuilder.get_int1_tyX      wwr$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_int8_tyInterpreterBuilder.get_int8_ty[  r+  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_uint8_tyInterpreterBuilder.get_uint8_ty^      xxr$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_int16_tyInterpreterBuilder.get_int16_tya  r2  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_uint16_ty InterpreterBuilder.get_uint16_tyd      yyr$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_int32_tyInterpreterBuilder.get_int32_tyg  r2  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_uint32_ty InterpreterBuilder.get_uint32_tyj  r9  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_int64_tyInterpreterBuilder.get_int64_tym  r2  r$   c                	"    \         P                  # r&   )r   rZ   r   s   &r!   get_uint64_ty InterpreterBuilder.get_uint64_typ  r9  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_fp8e4nv_ty!InterpreterBuilder.get_fp8e4nv_tys      }}r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_fp8e4b15_ty"InterpreterBuilder.get_fp8e4b15_tyv      ~~r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_fp8e4b8_ty!InterpreterBuilder.get_fp8e4b8_tyy  rI  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_fp8e5_tyInterpreterBuilder.get_fp8e5_ty|  r!  r$   c                	"    \         P                  # r&   )r   r   r   s   &r!   get_fp8e5b16_ty"InterpreterBuilder.get_fp8e5b16_ty  rM  r$   c                	.    \         P                  ! W4      # r&   )r   r   )r    elt_ty
addr_spaces   &&&r!   
get_ptr_tyInterpreterBuilder.get_ptr_ty  s    v22r$   c                	.    \         P                  ! W4      # r&   )r   r   )r    r   rH   s   &&&r!   get_block_tyInterpreterBuilder.get_block_ty  s    }}U**r$   c                	    \        \        P                  ! V.\        P                  R 7      \        P
                  4      # rQ   )r   rR   arraybool_r   r   r    r6   s   &&r!   get_int1InterpreterBuilder.get_int1  s$    BHHeWBHH=rwwGGr$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   	get_uint8InterpreterBuilder.get_uint8  $    BHHeWBHH=rxxHHr$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   get_int8InterpreterBuilder.get_int8  s$    BHHeWBGG<bggFFr$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   
get_uint16InterpreterBuilder.get_uint16  $    BHHeWBII>		JJr$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   	get_int16InterpreterBuilder.get_int16  ri  r$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   
get_uint32InterpreterBuilder.get_uint32  rp  r$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   	get_int32InterpreterBuilder.get_int32  ri  r$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  rZ   r   rc  s   &&r!   
get_uint64InterpreterBuilder.get_uint64  rp  r$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   	get_int64InterpreterBuilder.get_int64  ri  r$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   get_fp16InterpreterBuilder.get_fp16  $    BHHeWBJJ?LLr$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   get_fp32InterpreterBuilder.get_fp32  r  r$   c                	    \        \        P                  ! V.\        P                  R 7      \        P                  4      # r`  )r   rR   ra  r   r   rc  s   &&r!   get_fp64InterpreterBuilder.get_fp64  r  r$   c                	Z    \        \        P                  ! ^ .\        V4      R7      V4      # r   rQ   )r   rR   ra  r   )r    types   &&r!   get_null_value!InterpreterBuilder.get_null_value  s!    BHHaSd0CDdKKr$   c                	    V P                   f   \        R4      h\        \        P                  ! V P                   V,          .\        P
                  R7      \        P
                  4      # )Nzgrid_idx is NonerQ   )r  r   r   rR   ra  r   r   r    axiss   &&r!   create_get_program_id(InterpreterBuilder.create_get_program_id  sD    == /00BHHdmmD&9%:"((KRXXVVr$   c                	    \        \        P                  ! V P                  V,          .\        P                  R 7      \
        P                  4      # r`  )r   rR   ra  r  r   r   r  s   &&r!   create_get_num_programs*InterpreterBuilder.create_get_num_programs  s.    BHHdmmD&9%:"((KRXXVVr$   c                	    \        \        P                  ! VP                  \        R 7      \
        P                  4      pRpV P                  WWbW44      # rQ   N)r   rR   	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                	    \        \        P                  ! VP                  \        R 7      \
        P                  4      pV P                  WVRR4      # r  )r   rR   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                	   VP                  4       p\        V4      pVf-   \        \        P                  ! VP
                  VR7      V4      p\        P                  ! VP
                  VP
                  VP
                  V4      p	\        W4      # )NrQ   )r.   r   r   rR   r   r   _interpreterload)
r    ptrsr  r  cache_modifiereviction_policyr  r]   dtype_nprets
   &&&&&&&   r!   r  %InterpreterBuilder.create_masked_load  sg    &&( *= tyy!I8TE		499ejj(KC**r$   c                	n    \         P                  ! VP                  VP                  VP                  4      # r&   )r  storer   )r    r  r6   r  r  r  s   &&&&&&r!   r  &InterpreterBuilder.create_masked_store  s#    !!$))UZZCCr$   c                	   VP                   P                  pVP                  pV\        P                  8X  d   V\        P                  8X  g+   V\        P                  8X  dZ   V\        P                  8X  dE   \        VP                  W4R 4      P                  \        V4      4      p\        WRP                  4      # \        VP                  P                  \        V4      4      VP                  4      # r&   )r   r[   r   r   r   r   r   viewr   r   rY   )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                	$    V P                  W4      # r&   r  r    r  r  s   &&&r!   r  InterpreterBuilder.<lambda>      $..2Or$   c                	$    V P                  W4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	$    V P                  W4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	$    V P                  W4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	$    V P                  W4      # r&   r  r  s   &&&r!   r  r    s    s0Mr$   c                	$    V P                  W4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	$    V P                  W4      # r&   r  )r    r  r  	is_signeds   &&&&r!   r  r    s    T^^C=Zr$   c                	    VP                   P                  pVP                  p\        VP                  WEV4      P	                  \        V4      4      p\        WbP                  4      # r&   )r   r[   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    \        VP                  P                  \        V4      4      VP                  4      # r&   )r   r   r  r   r[   r  s   &&&r!   create_bitcast!InterpreterBuilder.create_bitcast  s%    CHHMM-*ABHOOTTr$   c                	    V! VP                   VP                   4      pVP                  P                  p\        WE4      '       g   VP	                  \        V4      4      p\        WE4      # r&   r   r   r[   r   rY   r   r   )r    lhsrhsopr   r   s   &&&&  r!   	binary_opInterpreterBuilder.binary_op  sN    CHHchh'99##%f77]]=#:;FF--r$   c                	B    V P                  W\        P                  4      # r&   r  rR   addr    r  r  s   &&&r!   r  r    s    "&&)Ir$   c                	B    V P                  W\        P                  4      # r&   r  rR   multiplyr  s   &&&r!   r  r        "++)Nr$   c                	B    V P                  W\        P                  4      # r&   r  rR   divider  s   &&&r!   r  r    s    ")))Lr$   c                	B    V P                  W\        P                  4      # r&   r  rR   fmodr  s   &&&r!   r  r        "'')Jr$   c                	B    V P                  W\        P                  4      # r&   r  rR   subtractr  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r        s(Mr$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    s    "))1Tr$   c                	$    V P                  W4      # r&   create_idivr  s   &&&r!   r  r        )9)9#)Cr$   c                	$    V P                  W4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    s    s(Hr$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r     r  r$   c                	B    V P                  W\        P                  4      # r&   )r  rR   
left_shiftr  s   &&&r!   r  r    s    s(Or$   c                	B    V P                  W\        P                  4      # r&   )r  rR   right_shiftr  s   &&&r!   r  r    s    "..)Qr$   c                	B    V P                  W\        P                  4      # r&   r  rR   r   r  s   &&&r!   r  r        $..2::*Nr$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r        T^^Cbjj-Qr$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r        DNN3RZZ,Pr$   c                	B    V P                  W\        P                  4      # r&   r  rR   r   r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r  	  r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r  
  r  r$   c                	B    V P                  W\        P                  4      # r&   r  rR   
less_equalr  s   &&&r!   r  r        DNN3R]],Sr$   c                	B    V P                  W\        P                  4      # r&   r  rR   lessr  s   &&&r!   r  r        DNN3RWW,Mr$   c                	B    V P                  W\        P                  4      # r&   r  rR   greater_equalr  s   &&&r!   r  r        DNN3REUEU,Vr$   c                	B    V P                  W\        P                  4      # r&   r  rR   greaterr  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r
  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  rR   equalr  s   &&&r!   r  r    s    4>>#BHH+Mr$   c                	B    V P                  W\        P                  4      # r&   r  rR   	not_equalr  s   &&&r!   r  r    s    4>>#BLL+Qr$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r
  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r        DNN3RXX,Nr$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r        DNN3R\\,Rr$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r
  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r    r!  r$   c                	B    V P                  W\        P                  4      # r&   r  r  s   &&&r!   r  r     r#  r$   c                	B    V P                  W\        P                  4      # r&   )r  rR   bitwise_andr  s   &&&r!   r  r  !      s(Pr$   c                	B    V P                  W\        P                  4      # r&   )r  rR   bitwise_xorr  s   &&&r!   r  r  "  r,  r$   c                	B    V P                  W\        P                  4      # r&   )r  rR   
bitwise_orr  s   &&&r!   r  r  #  s    t~~c'Nr$   c                	    \        VP                  \        P                  ! VP                  VP                  4      ,
          VP                  ,          VP                  P
                  4      # r&   )r   r   rR   r  r   r[   r  s   &&&r!   r  InterpreterBuilder.create_idiv'  sC     SXX#(((CCPRUR[R[RbRbccr$   c                	>   \        VP                  P                  4      p\        VP                  P                  4      pVP                  P                  V4      Vn        VP                  P                  V4      Vn        V P	                  W\
        P                  4      # r&   )r   r   r   rY   r  rR   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                	z   VP                   P                  pV\        P                  8X  g   V\        P                  8X  d?   \        \        VP                   VP                   4      VP                  P                  4      # \        \        R VP                  ^,          ^,           24      pVP                   P                  V4      pVP                   P                  V4      p\        P                  ! WV4      VP                  ^,          ,	          p\        VP                  V4      VP                  P                  4      # )r   )r   r   rR   r   rZ   r   np_umulhi_u64r[   r   r   rY   r  )r    r  r  r   compute_dtypelhs_datarhs_dataret_datas   &&&     r!   create_umulhi InterpreterBuilder.create_umulhi5  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                	    V! VP                   VP                   VP                   4      pVP                  P                  p\        WV4      '       g   VP	                  \        V4      4      p\        WV4      # r&   r  )r    r  r  r  r  r   r   s   &&&&&  r!   
ternary_opInterpreterBuilder.ternary_opA  sT    CHHchh

3;;%%%f77]]=#:;FF--r$   c                	D    V P                  WV\        P                  4      # r&   )rA  rR   clip)r    arglohipropagate_nanss   &&&&&r!   r  r  J  s    doocWY[][b[b>cr$   c                	D    V P                  WV\        P                  4      # r&   )rA  rR   where)r    condr  r  s   &&&&r!   r  r  K  s    CQSQYQY1Zr$   c                	    \        VP                  VP                  ,          VP                  ,           VP                  P                  4      # r&   r   r   r   r[   r  s   &&&&r!   
create_fmaInterpreterBuilder.create_fmaM  s,    AFFQVVOaff4aggnnEEr$   c                	b    \        V! VP                  4      VP                  P                  4      # r&   rM  )r    rE  r  s   &&&r!   unary_opInterpreterBuilder.unary_opQ  s!    BsxxL#))*:*:;;r$   c                	N   VP                   pVP                  ^,
          p\        \        RVP                   24      pVP                  P                  V4      p^V,          ^,
          pWV,          P                  \        V4      4      p\        WqP                   P                  4      # )r  r   )	r   r   r   rR   r   r  r   r   r[   )r    rE  r]   mask_bitwidthnp_uint_dtyper   r  r  s   &&      r!   create_fabsInterpreterBuilder.create_fabsT  s    99 33a7d8+F+F*G$HIxx}}]+]"a'{  x!89C!1!122r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   cosr    rE  s   &&r!   r  r  ^      4==bff#=r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   exprZ  s   &&r!   r  r  _  r[  r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   exp2rZ  s   &&r!   r  r  `      DMM#rww$?r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   absrZ  s   &&r!   r  r  a  s    DMM#rvv$>r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   floorrZ  s   &&r!   r  r  b  s    T]]3%Ar$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   ceilrZ  s   &&r!   r  r  c  r`  r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   logrZ  s   &&r!   r  r  d  r[  r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   log2rZ  s   &&r!   r  r  e  r`  r$   c                	B    V P                  V\        P                  4      # r&   rQ  rR   sqrtrZ  s   &&r!   r  r  f  s    DMM#rww,Gr$   c                	B    V P                  V\        P                  4      # r&   rl  rZ  s   &&r!   r  r  g  r`  r$   c                	B    V P                  V\        P                  4      # r&   )rQ  rR   sinrZ  s   &&r!   r  r  h  r[  r$   c                	    VP                   P                  \        P                  8X  d   \	        VP                   4      M\        VP                   4      p\        W!P                  P                  4      # r&   )r   r   rR   r   np_erf_fp32np_erf_fp64r   r[   )r    rE  r  s   && r!   
create_erfInterpreterBuilder.create_erfj  sF    '*xx~~'Ck#((#UXU]U]I^C!1!122r$   c                	    \        ^\        P                  ! VP                  4      ,          VP                  P
                  4      # r  )r   rR   rm  r   r   r[   rZ  s   &&r!   create_rsqrtInterpreterBuilder.create_rsqrtn  s+    A 113993C3CDDr$   c                	t    \        VP                  P                  V4      VP                  P                  4      # r&   )r   r   rX   r   r[   )r    rE  rH   allow_reorders   &&&&r!   r  r  r  s(    \#((JZJZ[`Jacfclclcscs=tr$   c                	    \        \        P                  ! VP                  V4      VP                  P
                  4      # r&   )r   rR   	transposer   r   r[   )r    rE  perms   &&&r!   create_transInterpreterBuilder.create_transt  s(    BLL48#)):J:JKKr$   c                	   VP                   pVP                   pVP                  P                  ^8X  d!   VP                  P                  4       '       g<   VP                  P                  ^8X  d   VP                  P                  4       '       d   \	        WaP                  \
        P                  R4      P                  \        P                  4      p\	        WrP                  \
        P                  R4      P                  \        P                  4      p\        \        P                  ! WgVP                   P                  R7      VP                   ,           VP                  P                  4      # )rP   NrQ   )r   r   r   is_floatingr   r   r   r  rR   r   matmulr[   )r    r   r   dinput_precisionmax_num_imprecise_acca_datab_datas   &&&&&&  r!   
create_dotInterpreterBuilder.create_dotw  s    GG&&!+0C0C0E0EGG&&!+0C0C0E0E#FGGRZZFKKBJJWF#FGGRZZFKKBJJWFBIIfAFFLLIAFFRTUT[T[TbTbccr$   c                	~    \        \        P                  ! W#\        P                  R 7      \        P                  4      # r`  )r   rR   rW   r   r   )r    ret_tystartstops   &&&&r!   create_make_range$InterpreterBuilder.create_make_range  s"    BIIeBBHHMMr$   c                	   Vf?   \        \        P                  ! VP                  \        R7      \
        P                  4      p\        P                  ! VP                  VP                  P                  R7      p\        P                  ! VP                  VP                  \        P                  ! VP                  4      4      p\        P                  ! W^ V3VR7      ^ ,          pV^ ;;,          \        P                  ! VP                  4      P                  4       ,          uu&   \        V\
        P                  4      # )NrQ   )binsrU   weights)r   rR   r  r   r'   r   r   r   rJ  r   	histogramlogical_notsumr   )r    r   r  r  dummy_weightsr  s   &&&&  r!   create_histogram#InterpreterBuilder.create_histogram  s    <TYYd CRWWMD TYYdiiooF xx		499bmmDII.FGLLD	=YZ[\	!tyy15577Irxx00r$   c                	    \        \        P                  ! VP                  VP                  VR 7      VP                  P
                  4      # r  )r   rR   take_along_axisr   r   r[   )r    r  indicesr  s   &&&&r!   create_gather InterpreterBuilder.create_gather  s3    B..sxxDQSVS\S\ScScddr$   c                	   VP                  4       pVP                  p\        ^V^,          4      p\        VP                  WRP                  P                  \        P                  4      ,          ,           VP                  4      # rw  )	r.   r   maxr   r   rY   rR   rZ   r   )r    r  offsetr]   element_bitwidthelement_bytewidths   &&&   r!   create_addptr InterpreterBuilder.create_addptr  sc    %%'#66#3q#89CHH'8;;;M;Mbii;X'XXZ]ZcZcddr$   c                	   VP                  V4      w  rxVP                  4       p	\        V	4      p
Vf   R pMV\        P                  P
                  8X  d.   \        \        P                  ! VP                  V
R7      V	4      pMdV\        P                  P                  8X  d8   \        \        P                  ! VP                  \        R4      V
R7      V	4      pM\        RV 24      hV P                  WxWWV4      # )NrQ   nanzunsupported padding option )re   r.   r   r   PADDING_OPTIONPAD_ZEROr   rR   r   r   PAD_NAN	full_likefloatr   r  )r    r  r\   padding_optionr  r  r  r  r`   r]   r  r  s   &&&&&&&     r!   create_tensor_pointer_load-InterpreterBuilder.create_tensor_pointer_load  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    VP                  V4      w  rgV P                  WbWtV4      # r&   re   r  )r    r  r6   r\   r  r  r  r`   s   &&&&&&  r!   create_tensor_pointer_store.InterpreterBuilder.create_tensor_pointer_store  s)    ..~>''UO\\r$   c                	    \        \        P                  ! VP                  V4      VP                  P
                  4      # r&   )r   rR   expand_dimsr   r   r[   )r    rE  r  s   &&&r!   create_expand_dims%InterpreterBuilder.create_expand_dims  s(    BNN388T:CII<L<LMMr$   c                	    \        \        P                  ! VP                  V4      VP                  P
                  4      # r&   )r   rR   rS   r   r   r[   )r    rE  rH   s   &&&r!   create_broadcast#InterpreterBuilder.create_broadcast  s(    BOOCHHe<cii>N>NOOr$   c                	    \        \        P                  ! VP                  VP                  .4      VP                  P
                  4      # r&   )r   rR   concatenater   r   r[   r  s   &&&r!   
create_catInterpreterBuilder.create_cat  s/    BNNCHHchh+?@#))BRBRSSr$   c                	    \        \        P                  ! VP                  VP                  .RR7      VP                  P
                  4      # )r  r  rv   )r   rR   stackr   r   r[   r  s   &&&r!   create_joinInterpreterBuilder.create_join  s1    BHHchh%9CSYYEUEUVVr$   c                	    \        VP                  R,          VP                  P                  4      \        VP                  R,          VP                  P                  4      3# ).).r   ).r  rM  )r    r  s   &&r!   create_splitInterpreterBuilder.create_split  sE    SXXf-syy/?/?@,sxxX^O_adajajaqaqBrssr$   c           	     	   VP                   p\        VP                  \        P                  4      '       d[   \        \        P                  ! W2P                  ^ ,          \        VP                  4      R7      VP                  P                  4      # \        \        P                  ! W2P                  \        VP                  4      R7      VP                  P                  4      # r  )rH   r   r   r   r   r   rR   fullr   r   r[   )r    r  rE  rH   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           	     	    \        \        P                  ! RVP                  ^ ,          \	        VP
                  4      R7      VP
                  P                  4      # )r  rQ   rw  )r   rR   r  r   r   r   r[   rZ  s   &&r!   create_unsplat!InterpreterBuilder.create_unsplat  s:    BGGE388A;mCII>VWY\YbYbYiYijjr$   c                	   W@P                   9  d   \        R V 24      hV P                   V,          p\        \        P                  ! VP
                  VP
                  VP
                  V4      VP                  P                  4      # )unsupported semantic )ir_sem_to_interpreter_semr   r   r  
atomic_casr   r   r[   )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           	     	p   WP                   9  d   \        R V 24      hWPP                  9  d   \        RV 24      hV P                   V,          pV P                  V,          p\        \        P
                  ! WP                  VP                  VP                  V4      VP                  P                  4      # )zunsupported rmwOp r  )	ir_rmw_op_to_interpreter_rmw_opr   r  r   r  
atomic_rmwr   r   r[   )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                	    \        R 4      h)z4extern_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                	    \        R 4      h)z,inline_asm not supported in interpreter moder  )r    	inlineAsmconstraintsvaluesr  r  packs   &&&&&&&r!   create_inline_asm$InterpreterBuilder.create_inline_asm      !"PQQr$   c                	   R V P                   ^ ,           RV P                   ^,           RV P                   ^,           R2pV'       d   VRV 2,          pV'       d   \        P                  ! RR /R7       V F"  p\        VRVP                   2,           4       K$  	  V'       d   \        P                  ! RR7       R# R# )(z, ) r(   c                    R V R 2# )0x02xr9   r   s   &r!   r  1InterpreterBuilder.create_print.<locals>.<lambda>  s    b3Lr$   )	formatterN)r  rR   set_printoptionsprintr   )r    prefixhexr  isSignedmsgr6   s   &&&&&  r!   create_printInterpreterBuilder.create_print  s    
 $--"#2dmmA&6%7r$--:J9K1MQvh<C52H*IJE#!EJJ<(() $/ r$   c                	&    V'       g	   Q V 4       hR # r&   r9   )r    	conditionmessages   &&&r!   create_assert InterpreterBuilder.create_assert  s    &WI&yr$   c                	$    V'       g   Q R 4       hR# )zAssume failedNr9   )r    r  s   &&r!   create_assume InterpreterBuilder.create_assume  s    )/)yr$   c                	    R # r&   r9   r   s   &r!   create_barrier!InterpreterBuilder.create_barrier  s    r$   c                	b    V Uu. uF  qwP                  4       NK  	  pp\        WW8WV4      # u upi r&   )r2   rE   )	r    rG   rH   rI   rJ   rK   rL   r  new_offsetss	   &&&&&&&  r!   create_make_block_ptr(InterpreterBuilder.create_make_block_ptr  s.    4;<G&||~G<!$w[XX =s   ,c                	   \        VP                  4      \        V4      8w  d   \        R 4      hVP                   Uu. uF  q3P                  4       NK  	  pp\	        VP
                  VP                  VP                  WAP                  VP                  4      p\        \        V4      4       F=  pVP                  V,          ;P                  W&,          P                  ,          un        K?  	  V# u upi )z len(ptr.offsets) != len(offsets))rV   rJ   r   r2   rE   rG   rH   rI   rK   rL   rU   r   )r    r  rJ   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               0    V ^8  d   QhRRRRRRRRRR	R
R/# )rj   rG   r   rH   rk   rI   tensor_shaperl   r  r'   rr   r   r9   )rm   s   "r!   rn   r     s;      , GY dv 4=JNY\r$   c                	@    \        WW4V4      pVP                  4        V# r&   )rh   r{   )r    rG   rH   rI   r  r  rr   descs   &&&&&&& r!   create_make_tensor_descriptor0InterpreterBuilder.create_make_tensor_descriptor  s    WGLr$   c                    V ^8  d   QhRRRR/# )rj   r  rh   r  rk   r9   )rm   s   "r!   rn   r     s      [ [+; [FX [r$   c           	     	0   \        V\        4      '       g   Q hVP                  V4      w  rVVP                  4       p\	        V4      pVP
                  p	V	\        P                  P                  8X  d.   \        \        P                  ! VP                  VR 7      V4      p
MdV	\        P                  P                  8X  d8   \        \        P                  ! VP                  \        R4      VR 7      V4      p
M\!        RV	 24      hV P#                  WVWVRR7      # )rQ   r  zunsupported padding F)r  r  r  )r   rh   re   r.   r   rr   r   r  r  r   rR   r   r   r  r  r  r   r  )r    r  r  r  r  r  r  r]   r  rr   r  s   &&&&&      r!   create_descriptor_load)InterpreterBuilder.create_descriptor_load  s    $ 01111..w7
&&( *,,c((111 tyy!I8TE**222 diiuX!VX`aE3G9=>>&&t57FTY ' [ 	[r$   c               $    V ^8  d   QhRRRRRR/# )rj   r  rh   r6   r   r  rk   r9   )rm   s   "r!   rn   r   #  s(     G G,< G\ G\n Gr$   c                	P    VP                  V4      w  rEV P                  WBVR R 4      # r&   r  )r    r  r6   r  r  r  s   &&&&  r!   create_descriptor_store*InterpreterBuilder.create_descriptor_store#  s+    ..w7
''T4FFr$   c               $    V ^8  d   QhRRRRRR/# )rj   r  rh   	x_offsetsr   y_offsetr9   )rm   s   "r!   rn   r   '  s#     	+ 	+-= 	+, 	+bn 	+r$   c                	   VP                   P                  P                  p\        V4      p\        P
                  ! VP                  P                  ^ ,          VP                  R,          .VR7      pRpRp	\        VP                  4       FA  w  r\        V\        P                  4      V.pV P                  WW4      P                  WzR3&   KC  	  \        Wu4      # )r   rQ   N:NNNrv   )rG   r   r,   r   rR   zerosr   rH   rK   	enumerater   r   r   r  )r    r  r#  r$  r  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               (    V ^8  d   QhRRRRRRRR/# )rj   r  rh   r6   r   r#  r$  r9   )rm   s   "r!   rn   r   2  s+     ? ?.> ?| ?`l ?,8?r$   c                	    \        VP                  4       FZ  w  rV\        VP                  V,          VP                  4      p\        V\        P
                  4      V.pV P                  WV4       K\  	  R # r&   )r'  r   r   r   r   r   r   )	r    r  r6   r#  r$  r   r*  slicer  s	   &&&&&    r!   create_descriptor_scatter,InterpreterBuilder.create_descriptor_scatter2  sT    $Y^^4KA A<E#Hbhh7BG((g> 5r$   c                	4   \        V4      pR VP                  9   d.   \        \        P                  ! ^RVR7      VP
                  4      # V\        P                  8X  d.   \        \        P                  ! ^RVR7      VP
                  4      # \        RV 24      h)r   rQ   Tzunsupported type rv   )r   namer   rR   r  r[   rb  	TypeError)r    r  np_types   && r!   get_all_ones_value%InterpreterBuilder.get_all_ones_value9  ss    %GLL 2W =t{{KK 4w ?MM/v677r$   )r   r
  r  r  r	  N)zero)r:   r;   r<   r=   r   MEM_SEMANTICACQUIREr  RELEASERELAXEDACQUIRE_RELEASEr  	ATOMIC_OPADDRMW_OPFADDMINUMINMAXUMAXANDORXORXCHGr  rM   r  r  r  r  r#  r&  r)  r-  r0  r4  r7  r;  r>  rA  rD  rG  rK  rO  rR  rU  rZ  r]  rd  rg  rk  rn  rr  ru  rx  r{  r~  r  r  r  r  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  r6  r>  rA  create_clampfcreate_selectrN  rQ  rV  
create_cos
create_expcreate_exp2create_iabscreate_floorcreate_ceil
create_logcreate_log2create_precise_sqrtcreate_sqrt
create_sinrt  rx  create_reshaper  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r   r  r  r  r  r  r  r  r   r+  r0  r6  rC   r9   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. 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. dMZMF<3 >J=J?K>KAL?K=J?KG?K=J3E uNLdN1$e
ei]NPTWtlkjrZR0'*Y
[ G	+?8r$   r   r   interpreter_semanticc                  >    ] tR tRtRtR R ltR R ltR R ltR	tR
# )_LangPatchScopeiH  z2Tracks patched attributes so they can be restored.c                   V ^8  d   QhRR/# r   r9   )rm   s   "r!   rn   _LangPatchScope.__annotate__K  s     = =$ =r$   c                	    . V n         R # r&   _changesr   s   &r!   rM   _LangPatchScope.__init__K  s	    :<r$   c               (    V ^8  d   QhRRRRRRRR/# )rj   objobjectr3  r   r6   r   r   r9   )rm   s   "r!   rn   r  N  s(     " "F "# "f " "r$   c                	x    \        W\        4      pV P                  P                  WV34       \	        WV4       R # r&   )r   _MISSINGr  appendsetattr)r    r  r3  r6   originals   &&&& r!   r7   _LangPatchScope.set_attrN  s/    3h/c235!r$   c                   V ^8  d   QhRR/# r   r9   )rm   s   "r!   rn   r  S  s     - - -r$   c                	    V P                   '       dC   V P                   P                  4       w  rpV\        J d   \        W4       KF  \	        WV4       KT  R # r&   )r  popr  delattrr  )r    r  r3  r  s   &   r!   restore_LangPatchScope.restoreS  sA    mmm"&--"3"3"5Cx8#"8, r$   r  N)	r:   r;   r<   r=   r>   rM   r7   r  rC   r9   r$   r!   r  r  H  s    <="
- -r$   r  c                   V ^8  d   QhRR/# rj   r  r  r9   )rm   s   "r!   rn   rn   \  s     * *? *r$   c                8    R V/R lpVP                  WV4       R# )memberc           
         T ! T/ VP                  4        UUu/ uF  w  r4VR 8w  g   K  W4bK  	  uppBR \        /B # u uppi )r   )itemsr  )r  argskwargskvs   $*,  r!   r  _patch_attr.<locals>.<lambda>]  s[     :wMS\\^AVM[TQDEDT BFM[AV:w bv:wAVs   <<N)r7   )r  r3  r  r   r  
new_members   &&&&& r!   _patch_attrr  \  s"    xf xJ 
NN3j)r$   c                   V ^8  d   QhRR/# r  r9   )rm   s   "r!   rn   rn   d  s     ; ; ;r$   c                    \         P                  ! V 4       F9  w  r4\        P                  P	                  V4      '       g   K,  \        WWAV4       K;  	  R # r&   )inspect
getmembersr   core
is_builtinr  )pkgr   r  r3  r  s   &&&  r!   _patch_builtinr  d  s:    **3/77f%%6E: 0r$   c                   V ^8  d   QhRR/# r  r9   )rm   s   "r!   rn   rn   j  s     : :o :r$   c                   a R  oR pVP                  V RR 4       VP                  V RV3R l4       VP                  V RR 4       VP                  V RR	 4       VP                  V R
\        V4      4       R# )c                j    V P                   P                  pVP                  ^8X  d   \        V4      # R# )r  T)r   r   sizer'   )r    r   s   & r!   	_get_bool%_patch_lang_tensor.<locals>._get_booll  s,    {{ "YY!^tDz55r$   c                   \        \        P                  ! V P                  P                  4      V P                  P
                  4      pV P                  P                  4       '       g   Q h\        V P                  P                  4      pVR,          VR,          uVR&   VR&   \        P                  P                  V P
                  V4      p\        P                  P                  W4      # )rj   rv   )r   rR   r}  r   r   r   r  is_blocklistrH   r   r  r   r   )r    r   rK   res_tys   &   r!   _get_transpose*_patch_lang_tensor.<locals>._get_transposer  s    bll4;;+;+;<dkk>O>OPyy!!####499??++6r?KO(BR##DJJ<ww~~f--r$   	__index__c                \    \        V P                  P                  P                  4       4      # r&   )r   r   r   squeezer   s   &r!   r  $_patch_lang_tensor.<locals>.<lambda>z  s    S9I9I9Q9Q9S5Tr$   r)   c                   < S! V 4      # r&   r9   )r    r  s   &r!   r  r  {  s	    IdOr$   __repr__c                @    \        V P                  P                  4      # r&   )reprr   r   r   s   &r!   r  r  |  s    D9I9I4Jr$   __str__c                @    \        V P                  P                  4      # r&   )r   r   r   r   s   &r!   r  r  }  s    3t{{7G7G3Hr$   r   N)r7   property)r   r  r  r  s   && @r!   _patch_lang_tensorr  j  sh    6. 
NN6;(TU	NN6:'CD	NN6:'JK	NN69&HI	NN63 89r$   c                  :    ] tR tRtR tR tR tR tR tR t	Rt
R	# )
ReduceScanOpInterfacei  c                	    Wn         W n        R # r&   r  
combine_fn)r    r  r  s   &&&r!   rM   ReduceScanOpInterface.__init__  s    	$r$   c                	T    Ve$   V\        V4      8  d   \        RV RV 24      hR # R # )Nzaxis z out of bounds for shape )rV   r   )r    rH   r  s   &&&r!   
check_axis ReduceScanOpInterface.check_axis  s4    E
 2uTF*CE7KLL !3r$   c                	    V Fj  p\        V\        P                  P                  4      '       g   \	        R \        V4       24      hV P                  VP                  V P                  4       Kl  	  R# )zinput must be a tensor, got N)	r   r   r  r   r   r  r  rH   r  )r    r   rE  s   && r!   check_tensor"ReduceScanOpInterface.check_tensor  sN    Cc277>>22 #?S	{!KLLOOCIItyy1 r$   c                	t   \        V4      p\        VR 4      '       dO   VP                  '       d=   VP                  V4      p\        P
                  ! V\        VP                  4      4      pM\        P                  ! V.VR7      pTp\        P                  P                  \        WP                  4      V4      # )rH   rQ   )r   r-   rH   rY   r   r   r  rR   ra  r  r   r   r[   )r    r  r   r(  ret_types   &&&  r!   	to_tensorReduceScanOpInterface.to_tensor  s|     '3  SYYY**X&C}}UDO<H((C51CHww~~l3=xHHr$   c                	    \        R 4      h)z,apply_impl must be implemented by subclassesr  r    r   s   &&r!   
apply_impl ReduceScanOpInterface.apply_impl  r  r$   c                	    \        V\        4      '       g   V P                  V34      ^ ,          # V P                  V4       V P	                  V4      p\        V\
        \        34      '       d   \        V4      # V3# )r   )r   tupleapplyr  r  r  )r    r   r  s   && r!   r  ReduceScanOpInterface.apply  sb    %''::ui(++% ooe$'dE];;uSzH#Hr$   r  N)r:   r;   r<   r=   rM   r  r  r  r  r  rC   r9   r$   r!   r  r    s$    %M2IRIr$   r  c                  J   a  ] tR tRtV 3R ltR tR tR	R ltR tR t	Rt
V ;t# )
	ReduceOpsi  c                	2   < \         SV `  W4       W0n        R # r&   )superrM   	keep_dims)r    r  r  r  	__class__s   &&&&r!   rM   ReduceOps.__init__  s    *"r$   c                	    . pV Fi  pVe   VP                  V4       K  ^ pVP                  V P                  VP                  P                  P	                  4       VP
                  4      4       Kk  	  \        V4      V3# r&   )r  r  r   r   flattenr   r  )r    r   r  r  r   s   &&&  r!   unravelReduceOps.unravel  sg    D

4 

4>>$++*:*:*B*B*DdjjQR  Sz4r$   c                	>  a aaa S P                   pS P                  SS P                   4      w  op. p. pS^ ,          P                  P                  P                  pV^ V Wc^,           R ,           pS Fl  pVP                  VP                  P                  4       VP                  \        P                  ! WxP                  P                  P                  R7      4       Kn  	  \        V^ ,          P                  4       EF  p	\        P                  ! W4      oS^ V SV^,           R ,           o\        ;QJ d"    . VVV 3R l\        V4       4       F  NK  	  5M! VVV 3R l\        V4       4       4      p
SV,          ^ 8X  dQ   \        \        V4      4       F6  pW,          P                  P                  P                  4       W[,          S&   K8  	  K  \        ;QJ d"    . VVV 3R l\        V4       4       F  NK  	  5M! VVV 3R l\        V4       4       4      pS P                   P"                  ! . VOV
O5!  p\%        V\        4      '       g   V3MTp\        \        V4      4       Fn  p\%        W,          \&        P(                  P*                  4      '       d+   W,          P                  P                  P                  4       MW,          W[,          S&   Kp  	  EK  	  . p\        V4       F  w  rS P,                  '       dO   Ve   \        P.                  ! W4      pMG\        \        V4      4       F  p\        P.                  ! V^ 4      pK  	  MVf   VP                  4       pVP                  S P1                  VSV	,          P                  4      4       K  	  V# )r   NrQ   c              3  |   <"   T F1  w  rSP                  VS,          SV,          P                  4      x  K3  	  R # 5ir&   r  r   ).0iir  r   input_indexr    s   &  r!   	<genexpr>+ReduceOps.generic_reduce.<locals>.<genexpr>  s1     s]rTYTVq~uRy O O]r   9<c              3  |   <"   T F1  w  rSP                  VS,          SV,          P                  4      x  K3  	  R # 5ir&   r  )r  oior   output_indexr    s   &  r!   r   r    s1     !w`vW\WY$..<%)//"R"R`vr  )r  r  r   r   rH   r  rR   r&  r   rU   r  unravel_indexr  r'  rV   rw   r  fnr   r   r  r   r  r  r  )r    r   original_axisr  
input_dataoutput_datainput_shapeoutput_shaperE  r   input_tuplej	acc_tuplecombine_fn_retr  r   _r  r  s   ff               @@r!   generic_reduceReduceOps.generic_reduce  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%%s]fgq]rssK4 A%s;/0A3>>3H3H3M3M3R3R3TKN<0 1 "E!w`iju`v!wEE!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                	   \        V\        4      '       d
   V^ ,          MTpRpRpV'       dN   V P                  V! VP                  P                  V P
                  V P                  R7      VP                  4      pV'       dR   V P                  V! VP                  P                  V P
                  V P                  R7      \        P                  4      pVe   Ve   WE3# Ve   V# Ve   V# \        R4      h)r   Nr  keepdimsz-val_reduce_op and idx_reduce_op are both None)r   r  r  r   r   r  r  r   r   r   r   )r    r   val_reduce_opidx_reduce_opr  idxs   &&&&  r!   min_maxReduceOps.min_max  s    &ue44a%..u||/@/@tyy[_[i[i!jlqlwlwxC..u||/@/@tyy[_[i[i!jlnltltuC?s8O_J_JLMMr$   c                	    V P                  \        P                  ! VP                  P                  V P
                  V P                  R 7      VP                  4      # )r  )r  rR   r  r   r   r  r  r   r  s   &&r!   r  ReduceOps.sum  s<    ~~bffU\\%6%6TYYQUQ_Q_`bgbmbmnnr$   c                	n   V P                   \        P                  P                  8X  d8   V P	                  V^ ,          \
        P                  \
        P                  R7      # V P                   \        P                  P                  8X  d8   V P	                  V^ ,          \
        P                  \
        P                  R7      # V P                   \        P                  P                  8X  d*   V P	                  V^ ,          \
        P                  RR7      # V P                   \        P                  P                  8X  d*   V P	                  V^ ,          \
        P                  RR7      # V P                   \        P                  P                  8X  d   V P!                  V^ ,          4      # V P#                  V4      # )r   )r  r  N)r  r   standard_argmin_combine_tie_break_leftr  rR   minargmin_argmax_combine_tie_break_leftr  argmax_elementwise_maxnanmax_elementwise_minnanmin_sum_combiner  r  r  s   &&r!   r  ReduceOps.apply_impl  s   ??bkkHHH<<abii<XX__ J JJ<<abii<XX__ < <<<<a		QU<VV__ < <<<<a		QU<VV__ 8 8888E!H%% &&u--r$   )r  r&   )r:   r;   r<   r=   rM   r  r  r  r  r  rC   __classcell__r  s   @r!   r  r    s)    # )VN$o. .r$   r  c                  @   a  ] tR tRtV 3R ltR tR tR tR tRt	V ;t
# )ScanOpsi  c                	2   < \         SV `  W4       W0n        R # r&   )r  rM   reverse)r    r  r  r1  r  s   &&&&r!   rM   ScanOps.__init__  s    *r$   c                	    V P                  \        P                  ! VP                  P                  V P
                  R 7      VP                  R7      .# r  rQ   )r  rR   cumsumr   r   r  r   r  s   &&r!   r5  ScanOps.cumsum  s8    ryy):):KSXS^S^_``r$   c                	    V P                  \        P                  ! VP                  P                  V P
                  R 7      VP                  R7      .# r4  )r  rR   cumprodr   r   r  r   r  s   &&r!   r8  ScanOps.cumprod  s8    rzz%,,*;*;$))LTYT_T_`aar$   c           	     	n  a aaa . p. pS^ ,          P                   P                  P                  pS Fl  pVP                  VP                   P                  4       VP                  \        P
                  ! WEP                   P                  P                  R7      4       Kn  	  \        V^ ,          P                  4       EF*  p\        P                  ! Wd4      o\        ;QJ d"    . VVV 3R l\        V4       4       F  NK  	  5M! VVV 3R l\        V4       4       4      pSS P                  ,          ^ 8X  dQ   \        \        V4      4       F6  pWx,          P                   P                  P                  4       W8,          S&   K8  	  K  \        ;QJ d*    . VV 3R l\        \        S4      4       4       F  NK  	  5M#! VV 3R l\        \        S4      4       4       4      o\        ;QJ d"    . VVV 3R l\        V4       4       F  NK  	  5M! VVV 3R l\        V4       4       4      p	S P                  P                   ! . V	OVO5!  p
\#        V
\        4      '       g   V
3MT
p	\        \        V4      4       Fn  p\#        W,          \$        P&                  P(                  4      '       d+   W,          P                   P                  P                  4       MW,          W8,          S&   Kp  	  EK-  	  . p\        V4       F7  w  rgVP                  S P+                  VSV,          P                  4      4       K9  	  V# )r   rQ   c              3  |   <"   T F1  w  rSP                  VS,          SV,          P                  4      x  K3  	  R # 5ir&   r  )r  r  r  indexr   r    s   &  r!   r   'ScanOps.generic_scan.<locals>.<genexpr>  s/     fPeur%%)//BBPer  c              3  v   <"   T F.  qSP                   8X  d   SV,          ^,
          MSV,          x  K0  	  R# 5i)r  Nr  )r  r   r<  r    s   & r!   r   r=  #  s-     "kYjTU		>58a<uQx#OYjs   69c              3  |   <"   T F1  w  rSP                  VS,          SV,          P                  4      x  K3  	  R # 5ir&   r  )r  r  r  r   
prev_indexr    s   &  r!   r   r=  $  s1     !u^tUZUW$..:b	"P"P^tr  )r   r   rH   r  rR   r&  r   rU   r  r  r  r'  r  rV   rw   r  r  r   r   r  r   r  )r    r   r
  r  rH   rE  r   r   r  r  r  r  r<  r@  s   ff          @@r!   generic_scanScanOps.generic_scan  s:   
a$$**Ccjjoo.rxxZZ__5J5JKL  z!}))*A$$Q.E5fPYZdPef55fPYZdPeffDTYY1$s;/0A,0GNN,?,?,D,D,FKN5) 1 #U"kY^_bch_iYj"kUU"kY^_bch_iYj"kk
!E!u^ghs^t!uEE!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           
     	   . pV P                   '       dj   V Fb  pVP                  V P                  \        P                  ! VP
                  P                  V P                  R 7      VP                  4      4       Kd  	  MTpV P                  \        P                  P                  8X  d   V P                  V^ ,          4      pMSV P                  \        P                  P                  8X  d   V P                  V^ ,          4      pMV P!                  V4      pV P                   '       dO   V FH  p\        P                  ! VP
                  P                  V P                  R 7      VP
                  n        KJ  	  V# r  )r1  r  r  rR   flipr   r   r  r   r  r   r   r*  r5  _prod_combiner8  rA  )r    r   	new_inputrE  r  s   &&   r!   r  ScanOps.apply_impl0  s    	<<<  

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

 
r$   )r1  )r:   r;   r<   r=   rM   r5  r8  rA  r  rC   r,  r-  s   @r!   r/  r/    s#    ab< r$   r/  c                   V ^8  d   QhRR/# r  r9   )rm   s   "r!   rn   rn   D  s     ; ;o ;r$   c                    RR lpRR lpV P                  \        RV4       V P                  \        RV4       V P                  \        P                  RV4       V P                  \        P                  RV4       R# )Fc                8    \        WV4      P                  V 4      # r&   )r  r  )r   r  r  r  r  s   &&&&,r!   _new_reduce'_patch_reduce_scan.<locals>._new_reduceH  s    95;;EBBr$   c                8    \        WV4      P                  V 4      # r&   )r/  r  )r   r  r  r1  r  s   &&&&,r!   	_new_scan%_patch_reduce_scan.<locals>._new_scanK  s    t177>>r$   reduceassociative_scanN)F)r7   r   r  )r  rK  rN  s   &  r!   _patch_reduce_scanrR  D  sY    C? 
NN2x-	NN2)95	NN277Hk2	NN277.	:r$   c                   V ^8  d   QhRR/# r  r9   )rm   s   "r!   rn   rn   T  s     K K/ Kr$   c           	        R  pRR lpRR lpR pVP                  V RV4       VP                  V RV4       VP                  V RV4       VP                  V R\        4       VP                  V P                  R	V4       VP                  V R
\        VRR7      4       VP                  V R\        VRR7      4       VP                  V R\        VRR7      4       \	        V4       R# )c                   V P                   R 8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP	                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R	8X  d   VP                  4       # V P                   R
8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP                  4       # V P                   R8X  d   VP!                  4       # V P                   R8X  d   VP#                  4       # \%        RV  R24      h)voidr   r   r   r   r   r   r   r   rZ   r   r   r   fp16bf16fp32fp64zfail to convert z to ir type)r3  get_void_tyr)  r-  r0  r4  r7  r;  r>  rA  rD  rR  rG  rK  r  r  r#  r&  r   )r    r   s   &&r!   
_new_to_ir$_patch_lang_core.<locals>._new_to_irV  s   99&&((YY& &&((YY& &&((YY'!''))YY'!''))YY("((**YY'!''))YY("((**YY'!''))YY("((**YY'!''))YY)#))++YY*$**,,YY& &&((YY& &&((YY& ''))YY& ((**+D6=>>r$   Nc                :    Vf   ^pVf   ^ T rTMYrT\        WEV4      # r&   )rU   )arg1arg2stepr  r  ends   &&&,  r!   
_new_range$_patch_lang_core.<locals>._new_range~  s*    <D<D33U&&r$   c                $    V '       g   Q V4       hR # r&   r9   )rK  r  s   &&r!   _new_static_assert,_patch_lang_core.<locals>._new_static_assert  s    Str$   c                   \        V \        P                  4      '       g   V # \        V\        \        34      '       g   V.MTpV Uu. uF1  p\        V\        P
                  4      '       d   VP                  MTNK3  	  pp\        V4      \        ^\        V P                  4      4      8w  d   \        RV 24      hV P                  P                  W!4       V # u upi )r  z$len(values) != len(input.shape) for )r   r   r   r  r  	constexprr6   rV   r  rH   r   r   r7   )r   r  r3  r  s   &&& r!   	_set_attr#_patch_lang_core.<locals>._set_attr  s    %++L!+FT5M!B!B&IOPAZ2<<88!''a?Pv;#aU[[!122CD6JKKd+	 Qs   7CrU   static_rangestatic_assertstatic_printto_irmultiple_ofztt.divisibility)r3  max_contiguousztt.contiguitymax_constancyztt.constancy)NN) )r7   r  r   r   rR  )langr  r\  rc  rf  rj  s   &&    r!   _patch_lang_coreru  T  s    $?P'
 
NN4*-	NN44	NN4*<=	NN4/	NN4::w
3	NN4	@Q(RS	NN4)79?+ST	NN4').*QRur$   c                n   \        4       pV P                  P                  4        UUu. uFB  w  r#\        P                  ! V4      '       g   K#  V\
        \
        P                  39   g   K@  VNKD  	  ppp\        V4      ^8  g   Q R4       hV Fv  p\        V\        V4       \        VP                  \        V4       V\
        8X  d   \        VP                  \        V4       \        VP                  V4       \        WQ4       Kx  	  \        \
        P                  P                  \        V4       V# u uppi )r  z:triton.language must be visible from within jit'd function)r  __globals__r  r  ismoduler   r  rV   r  interpreter_builderr   r   r  ru  tensor_descriptor_base)r  r  r  r6   langsrt  s   &     r!   _patch_langr|    s    E#%>>#7#7#9p#9xqW=M=Me=TUY^cegigngnboYoUU#9Epu:?XXX?t0%8t{{$7?2:499&95A4;;.%  277113FNL qs    D1D1*D1c                (   \        V \        4      '       Ed"   \        P                  ! \        P
                  P                  P                  V 4      R 4      p\        P                  pRT u;8:  d   R8  d   M M\        P                  pMwRT u;8:  d   R8  d   M M\        P                  pMTR	T u;8:  d   R8  d   M M\        P                  pM1RT u;8:  d   R
8  d   M M\        P                  pM\        RV  24      h\        \        P                  ! V .VR7      V4      p\        P                   ! W14      # \#        V R4      '       d   \        P                  ! \        P
                  P                  P                  V 4      R 4      p\        \        P                  ! V P%                  4       .\        P                  R7      V4      p\        P                   ! W14      # \        V \&        4      '       d   \)        V \+        \,        V 4      4      # \        V \.        4      '       d   V P0                   Uu. uF  p\-        V4      NK  	  ppV P0                  R,          ^8X  g   Q h\        P2                  ! ^4      VR&   \4        P7                  \-        V P8                  4      V P:                   Uu. uF  p\-        V4      NK  	  upTV P<                   Uu. uF  p\        P2                  ! V4      NK  	  upV P>                  R7      # V # u upi u upi u upi )NzUnsupported integer value rQ   data_ptr)rG   rH   rI   rK   r  l        i   l        l            l         l            rv   ) r   r   r   	str_to_tytritonruntimejitmangle_typerR   r   r   r   rZ   r   r   ra  r   r-   r~  r  r   map_implicit_cvtr   rI   ri  r  make_tensor_descriptorrG   rH   rK   rr   )rE  tyr   r   srI   r   s   &      r!   r  r    s2   #s\\&..,,88=tDS 5 HHEc!E!IIEs"U"HHEc!E!IIE9#?@@bhhuE:B?yy$$sJ\\&..,,88=tDbhh'7ryyI2Nyy$$	C		S#mS"9::	C)	*	*-0[[9[=#[9{{2!###ll1o#::chh@W\_\e\eAf\eWX-PQBR\eAfpwadapapGqap\]UVapGqJM++ ; W 	W J : BgGqs   2L,L

Lc                    \        V \        P                  P                  P                  4      '       d   V P
                  # V # r&   )r   r  r  r  TensorWrapperrG   )ts   &r!   _unwrap_tensorr    s-    !V^^''5566vvHr$   c                    \        V\        P                  P                  P                  4      '       d4   \        P                  P                  P	                  WP
                  4      # V # r&   )r   r  r  r  r  r   )r  original_tensors   &&r!   _rewrap_tensorr    sE    /6>>#5#5#C#CDD~~!!//3H3HIIHr$   c                  4    ] tR tRt. 3R ltR tR tR tRtR# )GridExecutori  c                	   ^RI Hp Wn        W n        W0n        W@n        VP                  P                  4        UUu/ uF  w  rgWe! V4      bK  	  pppV Uu. uF  qhP                  V4      R8X  g   K  VNK  	  upV n	        R# u uppi u upi )r  )_normalize_tyri  N)
r  r  r  	arg_namesgridpre_run_hooksr?   r  get
constexprs)	r    r  r  r  r  r  r3  r  r?   s	   &&&&&    r!   rM   GridExecutor.__init__  s}    &"	*CECUCUC[C[C]^C]xt4r!22C]^,5bID9L9LT9RVa9a44Ib _bs   BB2Bc                	   aa	 / o	VV	3R  loV Uu. uF  pS! V4      NK  	  pp/ pVP                  4        F  w  rgS! V4      WV&   K  	  WE3# u upi )c                  < \        V \        4      '       d   \        V \        SV 4      4      # \        V \        4      '       dS   \	        S! V P
                  4      V P                  V P                  V P                  V P                  V P                  4      # \        V R 4      '       g   V # \        V 4      pVP                  4       P                  4       S9  d2   VP                  4       pVP                  4       SVP                  4       &   SVP                  4       P                  4       ,          pVP!                  ^ RR7      pVP#                  W!P%                  4       VP'                  4       VP)                  4       4       \+        W0R7      pV# )r~  cpu)device)r  )r   r  r   r  r   rG   rH   rI   rK   rr   round_f32_to_tf32r-   r  untyped_storager~  r  	new_emptyset_storage_offsetr  ry   r  )rE  unwrapped_argstoragecpu_arg_to_cpustoragess   &   r!   r  ,GridExecutor._init_args_hst.<locals>._to_cpu  s;   #u%%$S#gs*;<<C!122'CHH%IIKKOOKK))  S*--
*3/M,,.779I'779/6{{}))+,}<<>GGIJG#--a->GLL">">"@-BTBTBVXeXlXlXno$WBGNr$   )r  )
r    args_devr  rE  args_hst
kwargs_hstr5   r6   r  r  s
   &&&     @@r!   _init_args_hstGridExecutor._init_args_hst  sY    	4 -55HSGCLH5 
 ,,.JC%enJO )## 6s   Ac                	   a
a / oV
V3R  lo
\        W4       F  w  rVS
! WV4       K  	  VP                  4        F  w  rxWG,          p	S
! W4       K  	  SP                  4        F  w  rVVP                  V4       K  	  R# )c                  < \        V R 4      '       dW   \        V 4      \        V4      rV P                  4       VP                  4       3SV P                  4       P                  4       &   R# \	        V \
        4      '       d   \        W4       F  w  rS! W4       K  	  R# \	        V \        4      '       d    S! V P                  VP                  4       R# R# )r~  N)	r-   r  r  r~  r   r  zipr   rG   )arg_devarg_hst	_from_cpur  s   &&r!   r  1GridExecutor._restore_args_dev.<locals>._from_cpu  s    w
++#1'#:N7<SBIBYBYB[]d]t]t]vAw002;;=>GU++*-g*?&Wg/ +@G%566',,5 7r$   N)r  r  r  copy_)r    r  r  r  r  r  r  r5   	kwarg_dev	kwarg_hstr  r  s   &&&&&     @@r!   _restore_args_devGridExecutor._restore_args_dev  sn    		6 !$H 7Gg' !8 %llnNC"Ii+ - #+//"3WMM'" #4r$   c                	   \         P                  ! V P                  4      pVP                  4        UUu/ uF  w  rEWCP                  9   g   K  WEbK  	  pppV P                  W4      w  rgV P                   F  pV! V/ VB  K  	  \        V P                  4      p	 \         P                  ! V P                  .VO5/ VB p
V
P                  4        UUu/ uF#  w  rYV P                  9   d   TM
\        V4      bK%  	  p
pp\        V P                  4      '       d   V P                  V
4      MV P                  p\        V4      ^8:  g   Q R4       hVR^\        V4      ,
          ,          ,           p\        P                  ! V!    \!        V^ ,          4       F]  p\!        V^,          4       FD  p\!        V^,          4       F+  p\        P#                  WV4       V P                  ! R/ V
B  K-  	  KF  	  K_  	   T	P3                  4        T P5                  YY'4       R# u uppi u uppi   \$         dF   p\&        P(                  P*                  P,                  '       d   h \/        \1        T4      4      ThRp?ii ; i  T	P3                  4        i ; i)   z#grid must have at most 3 dimensionsNrw  r9   )r  getfullargspecr  r  r  r  r  r|  getcallargsr  r  callabler  rV   ry  r  rU   r  	Exceptionr  knobscompilationfront_end_debuggingr   r  r  r  )r    r  r  argspecr  r  r  r  hookpatch_scoper  r3  rE  r  r   r  r  es   &*,               r!   __call__GridExecutor.__call__)  s    ((1#)<<>G>41Q,,5F$!$>G#228D&&D()j) ' "$''*	" &&twwHHZHDbfblblbnobnU^UYD!8#mC>PPbnDo&.tyy&9&9499T?tyyDt9>H#HH>%1s4y=11D,,d3	7tAwA"47^!&tAwA/<<Q1E GGOdO "0 , ( !x6F? H p  7<<++???&tAw/Q67
 !sJ   HH7I8 )H:BI8 A3H% I8 %I50A I00I55I8 8J
)r  r  r  r  r  N)	r:   r;   r<   r=   rM   r  r  r  rC   r9   r$   r!   r  r    s    :< c#$J#2$Gr$   r  c                      ] tR tRtR tRtR# )ASTTransformeriP  c           	     	   . pVP                    F  pW P                  V4      .,          pK  	  \        V4      ^8  d   \        R4      h\        P
                  ! \        P                  ! \        P                  ! R\        P                  ! 4       R7      R\        P                  ! 4       R7      VP                  \        P                  ! RR7      .. R7      Vn	        V# )	r  z&Multiple assignments are not supportedr  )idctxr  )r6   r   r  F)r6   )funcr  keywords)targetsvisitrV   r   astCall	AttributeNameLoadr6   Constant)r    nodenamestargets   &&  r!   visit_AssignASTTransformer.visit_AssignR  s    llFjj())E #u:>EFF XXSXX1GSXXZ%X_j#&88:/6:jj#,,UZB[5\gik
 r$   r9   N)r:   r;   r<   r=   r  rC   r9   r$   r!   r  r  P  s    r$   r  c                  N    ] tR tRt]! 4       tR tR tR tR t	R t
R tR tR	tR
# )FunctionRewriteri`  c                	:    Wn         W n        R V n        ^ V n        R# )rs  N)r  r  filenamedef_file_lineno)r    r  r  s   &&,r!   rM   FunctionRewriter.__init__c  s    $%r$   c                	N    \         P                  ! V P                  4      w  rT P	                  4       w  T n        T n        T P                  T4      T n        T P                  T4      pT P                  T4      pT P                  T4      #   \         d    T P                  u # i ; i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_astj  s    	--dgg6HE /3.H.H.J+t+../""5)--c2%%o66  	77N	s   "B
 
B$#B$c                	B    ^RI HpHp V! V! V P                  4      4      # )r  )get_jit_fn_file_lineJITFunction)r  r  r  r  )r    r  r  s   &  r!   r  &FunctionRewriter._get_jit_fn_file_line~  s    :#K$899r$   c                	    ^ p\        V4       F5  w  r4VP                  4       P                  R4      '       g   K,  V^,           pK7  	  V# )r   zdef )r'  strip
startswith)r    r  r  r   lines   &&   r!   r  FunctionRewriter._find_def  s@    
 'GAzz|&&v..U
 ( r$   c                	z    WP                   ^,
          R pRP                  V4      p\        P                  ! V4      # )r  Nrs  )r  jointextwrapdedent)r    r  r  s   && r!   r   FunctionRewriter._prepare_source  s2    oo)*+ggens##r$   c                	    \         P                  ! V4      pV P                  P                  V4      p\         P                  ! V4       V P
                  ^,
          p\         P                  ! W44       V# rw  )r  parseast_transformerr  fix_missing_locationsr  increment_lineno)r    r  
parsed_astr  
inc_linenos   &&   r!   r  FunctionRewriter._transform_ast  sY     YYs^
..44Z@!!/2))A-
_9r$   c                	"   \        WP                  R R7      p/ V P                  CpV P                  P                  p\        4       P                  4        F  w  rVWT9  g   K  WdV&   K  	  \        W$V4       W0P                  P                  ,          # )exec)r  mode)	compiler  r  r  rw  globalsr  r  r:   )r    r  compiled_codelocal_namespace
fn_globalsr5   r6   s   &&     r!   r  "FunctionRewriter._compile_and_exec  so    --fU)T[[/WW((
!)//+JC$"'3 , 	]8ww//00r$   )r  r  r  r  r  N)r:   r;   r<   r=   r  r  rM   r  r  r  r  r  r  rC   r9   r$   r!   r  r  `  s-    $&O&7(:$
	1r$   r  c                  Z    ] tR tRt$ / tR]R&   R R ltR tR tR t	]
R	 4       t R
 tRtR# )InterpretedFunctioni  zDict[Callable, Callable]rewritten_fnc                   V ^8  d   QhRR/# r   r9   )rm   s   "r!   rn    InterpretedFunction.__annotate__  s     I I Ir$   c                	    Wn         \        V3/ VB V n        W n        . V n        \
        P                  ! V4      pVP                  P                  4        Uu. uF  qDP                  NK  	  upV n
        R # u upi r&   )r  r  rewriterr  r  r  	signature
parametersr  r3  r  )r    r  r  r  r  s   &&,  r!   rM   InterpretedFunction.__init__  sa    (6v6%%b)	*3*>*>*E*E*GH*GQ&&*GHHs   A5c               	    V'       d   R # V P                  4       p\        WPP                  WP                  4      ! V/ VB # r&   )rewriter  r  r  )r    r  warmupr  r  r  s   &$$*, r!   runInterpretedFunction.run  s5    \\^B6H6HI4ZSYZZr$   c                	b    \        V4      '       g   Q hV P                  P                  V4       R # r&   )r  r  r  )r    r  s   &&r!   add_pre_run_hook$InterpretedFunction.add_pre_run_hook  s$    ~~~!!$'r$   c                	    V P                   V P                  9  d2   V P                  P                  4       V P                  V P                   &   V P                  V P                   ,          # r&   )r  r  r  r  r   s   &r!   r  InterpretedFunction.rewrite  sJ    77$+++)-)B)B)DDdgg&  ))r$   c                	.    V P                   P                  # r&   )r  r:   r   s   &r!   r:   InterpretedFunction.__name__  s    wwr$   c                	    \        V P                  4       V P                  4       p V! V/ VB #   \         d   p\	        \        T4      4      ThR p?ii ; ir&   )r|  r  r  r  r   r  )r    r  r  r  r  s   &*,  r!   r  InterpretedFunction.__call__  sO    DGG\\^	3t&v&& 	3"47+2	3s   / AAA)r  r  r  r  r  N)r:   r;   r<   r=   r  r?   rM   r  r  r  r  r  rC   r9   r$   r!   r  r    s<    -/L*/I[(*
    3r$   r  )R__conditional_annotations__
__future__r   r  r  r  typingr   r   r   r   r   r	   r   numpyrR   r  triton.languagelanguager   r@   r
   triton.language.semanticr   triton.runtime.jitr   triton.tools.tensor_descriptorr   errorsr   	functoolsr   _C.libtritonr   r  r   r   _utilsr   r   r   rE   rh   r   r   r   r   r   r   r   	vectorizer   rr  r   rs  rZ   r9  r   r   r  r  ry  r  r?   r  r  r  r  r  r  r/  rR  ru  r|  r  r  r  r  NodeTransformerr  r  r  )r&  s   @r!   <module>r5     s   " " 
   A A      ! 3 . ; $  6 $ "CL   @" "4(" ("V $
& 
& 
& 	@='@
# ll45ll45Z<p p^8 ^8B 8(* '56I'J n J- -(*;:."I "IJ].% ].@;# ;|; K\ BnG nGbS((  B1 B1J'3/!, '3r$   