+
    ړLj                       ^ RI HtHt ^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RI	t	^ RI
t
^ RIHt ^ RIHt ^ RIHt ^ RIHtHtHtHtHtHtHtHt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$H%t%H&t&H't' ^RI(H)t) ^ RI*H+t+H,t,H-t- Rt.Rt/]! R4      t0 ! R R]Pb                  4      t2R R lt3 ! R R4      t4R>R lt5 ! R R]]0,          4      t6R t7R t8R t9 ! R R4      t:/ t;] ! R  R!4      4       t<R" t=R# t> ! R$ R%]:]6]0,          4      t?]R& R' l4       t@]R(RR)RR*RR+RR,RR-RR.R/R/ R0 ll4       t@R?R(RR)RR*RR+RR,RR-RR.R/R1 R2 lllt@ ! R3 R44      tA ! R5 R64      tBR7 tCR8 tD ! R9 R:]:4      tE ! R; R<]:4      tFR= tGR# )@    )annotationsdivisionN)defaultdict)	dataclass)cached_property)	CallableGenericIterableOptionalTypeVaroverloadDictAnyTuple)BaseBackend)
ModuleType)knobs)driver)_async_compile)find_paths_ifget_iterable_pathtype_canonicalisation_dictis_namedtuple)get_cache_key)get_cache_invalidating_env_varsnative_specialize_implirztriton.languagez"triton.experimental.gluon.languageTc                     a  ] tR t^"tRtR V 3R llt]R 4       tR tR 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V ;t# )DependenciesFindera  
This AST visitor is used to find dependencies of a JITFunction. This can
be used to invalidate a JITFunction's hash when its source code -- or
that of its dependencies -- changes.

This visitor also keeps track of the global variables touched by the
JITFunction.  When we launch the kernel, we check that these have the same
values as they did when we ran this visitor.  If not, we raise an error (or
otherwise we could recompile).
c                   V ^8  d   QhRR/# )   returnNone )formats   ">/app/.local/lib/python3.14/site-packages/triton/runtime/jit.py__annotate__DependenciesFinder.__annotate__.   s     )0 )0 )0    c                	   < \         SV `  4        Wn        \        P                  ! VP                  R 4      4      V n        W n        W0n        0 RmV n	        \        \        RR0V n        / V n        RV n        R# )utf-8copymathFN>
   intlenmaxminlistfloatprintrangegetattr
isinstance)super__init__namehashlibsha256encodehasherglobals	nonlocalssupported_python_builtinsGLUON_MODULETRITON_MODULEsupported_modulesused_global_valsvisiting_arg_default_value)selfr;   r@   rA   src	__class__s   &&&&&r'   r:   DependenciesFinder.__init__.   sn    	nnSZZ%89 "*
& 	"
" TV*/'r*   c                	6    V P                   P                  4       # N)r?   	hexdigestrH   s   &r'   retDependenciesFinder.retY   s    {{$$&&r*   c                	    \         P                  ! VP                  4      '       d   R # \        VRR4      pVP	                  \
        4      # )T
__module__ )inspect	isbuiltinfuncr7   
startswithrD   )rH   noderW   modules   &&& r'   _is_triton_builtin%DependenciesFinder._is_triton_builtin]   s9    TYY''|R0  //r*   c                	l   \        V\        4      '       g   Q hV P                  P                  4       VP                  P                  4       ,           Ff  pVw  r4V P                  V,          w  rTVP                  V,          w  rdWV8w  g   K9  \	        R V RV RV P
                   RVP                   RV R24      h	  V P                  P                  VP                  4       VP                  pV\        \        VRR4      4      ,          pV P                  P                  VP                  R4      4       R	# )
Global variable z has value z when compiling z, but inner kernel z has conflicting value z7 from when it was first compiled.  This is not allowed.noinlineFr,   N)r8   JITCallablerF   keysRuntimeErrorr;   __name__update	cache_keystrr7   r?   r>   )rH   rW   kvar_name_v1v2func_keys   &&      r'   _update_hashDependenciesFinder._update_hashc   s)   $,,,, &&++-0E0E0J0J0LLLAKH))!,EB))!,EBx"&xjB4?OPTPYPY{Zmnrn{n{m|  }T  UW  TX  XO  P  M 	$$T%:%:;>>Cj%8998??734r*   c                	v   ^RI Hp Ve   \        V4      \        J d   R# \	        VRR4      '       d'   VP
                   F  pV P                  V4       K  	  R# \	        VRR4      '       d   R# \	        VRR4      R8X  d   R# \        V\        4      '       d   V P                  V4       R# \        V4      '       d6   \        V\        4      '       g    \        W4      '       g   \        R	V 24      hV P                  '       d   R# Ve1   \        P                  ! V4      V3V P                  V\!        V4      3&   R# )
r"   	constexprN__triton_aggregate__F__triton_builtin__rS   rT   ztriton.language.extra.libdevicez!Unsupported function referenced: )language.corerq   typer   r7   
hash_attrsrecord_referencer8   r`   rm   callablerb   rG   r-   deepcopyrF   id)rH   valvar_dictr;   rq   attrs   &&&&  r'   rw   #DependenciesFinder.record_referenceu   s   - ;$s)z13.66%%d+ '3,e44 3b)-NNc;''c"C==C!6!6z#?Y?Y!B3%HII ***;?==;Mx:XD!!4H"67r*   c                	N  a  \        VP                  4      \        P                  J d   VP                  # VP                  S P
                  9   d   R # V 3R lpV! VP                  4      w  r4VP                  S P                  9   d   V# S P                  W4VP                  4       V# )Nc                   < SP                   P                  V R 4      pVe   VSP                   3# SP                  P                  V R 4      pVe   VSP                  3# R# )NNN)r@   getrA   )r;   r{   rH   s   & r'   name_lookup2DependenciesFinder.visit_Name.<locals>.name_lookup   s]    ,,""4.CDLL((..$$T40CDNN**r*   )ru   ctxastStorerz   local_namesrB   rw   )rH   rY   r   r{   r|   s   f&   r'   
visit_NameDependenciesFinder.visit_Name   s|    >SYY&77N77d&&&	 $DGG,77d444JcTWW5
r*   c                	`    VP                    Uu. uF  q P                  V4      NK  	  up# u upi rM   )eltsvisit)rH   rY   elts   && r'   visit_TupleDependenciesFinder.visit_Tuple   s&     ,09959C

39555s   +c                	H   V P                  VP                  4      p\        V\        P                  4      '       d   V P                  VP                  4      pK=  \        VR R4      pVe   W0P                  9   d   R# \        W!P                  4      pV P                  V4       V# )rc   rT   N)	r   valuer8   r   	Attributer7   rE   r}   rw   )rH   rY   lhslhs_namerP   s   &&   r'   visit_Attribute"DependenciesFinder.visit_Attribute   sy    jj$cmm,,**SYY'C3
B/;(&<&<<c99%c"
r*   c                	    VP                   P                    Uu0 uF  q"P                  kK  	  upV n        V P                  V4       R # u upi rM   )argsargr   generic_visit)rH   rY   r   s   && r'   visit_FunctionDef$DependenciesFinder.visit_FunctionDef   s6    /3yy~~>~GG~>4  ?s   Ac                	  a  V 3R  lp\         P                  ! VP                  VP                  VP                  '       d   VP                  .M. VP
                  4       F  pS P                  V4       K  	  V! VP                  4       VP                  e   S P                  VP                  4       V! VP                  4       R# )c                   <  SP                   '       d   Q hR Sn         V  F  pVf   K	  SP                  V4       K  	  RSn         R#   RSn         i ; i)TNF)rG   r   )defaultsexprrH   s   & r'   visit_defaults:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sQ    8:::::26/$D'

4( % 38/%/s   A A A 	AN)
	itertoolschainposonlyargsr   vararg
kwonlyargsr   kw_defaultskwargr   )rH   rY   r   r   s   f&  r'   visit_arguments"DependenciesFinder.visit_arguments   s    	8 ??4#3#3TYYQUQ\Q\Q\bdfjfufuvCJJsO w 	t''(::!JJtzz"t}}%r*   c                	    V P                  V4      p\        V\        4      '       d%   V ;P                  \	        V4      ,          un        R # V P                  P                  V4       R # rM   )r   r8   r3   r   setadd)rH   rY   targets   && r'   visitAssnTarget"DependenciesFinder.visitAssnTarget   sH     D!fd##F+  (r*   c                	    \        VP                  4      ^8w  d   \        R4      hV P                  VP                  ^ ,          4       V P	                  V4       R# )   z2Simultaneous multiple assignment is not supported.N)r0   targets	TypeErrorr   r   rH   rY   s   &&r'   visit_AssignDependenciesFinder.visit_Assign   sG    t||!
 PQQT\\!_- 	4 r*   c                	^    V P                  VP                  4       V P                  V4       R # rM   r   r   r   r   s   &&r'   visit_AnnAssign"DependenciesFinder.visit_AnnAssign   $    T[[) 	4 r*   c                	^    V P                  VP                  4       V P                  V4       R # rM   r   r   s   &&r'   	visit_ForDependenciesFinder.visit_For  r   r*   )	r@   r?   r   r;   rA   rE   rB   rF   rG   r   )rc   rS   __qualname____firstlineno____doc__r:   propertyrP   r[   rm   rw   r   r   r   r   r   r   r   r   r   __static_attributes____classcell__rJ   s   @r'   r    r    "   sj    	)0 )0V ' '05$%N06
	!
&@)!!! !r*   r    c                   V ^8  d   QhRR/# r"   r#   rf   r%   )r&   s   "r'   r(   r(     s     D D Dr*   c                   ^ RI Hu Hp \        V \        4      '       d   V P                  4       p V P                  R4      '       dF   V P                  R4      p \        V 4      p V P                  R4      '       g   Q hRV R,          ,           # V P                  R4      '       d   R\        V RR
 4      ,           # V P                  R4      '       d   R\        V R,          4      ,           # V P                  R4      '       d   \        V P                  R4      4      # M\        WP                  4      '       d   R\        V P                  4       2# \        WP                  4      '       d   V P                  p M.\        V \        4      '       d   V P                  p M\	        V 4      p \         P"                  ! V P%                  RR	4      V 4      # )r   Nzconst const**kr   NNztl._trT   )triton.language.corelanguagecorer8   rf   striprX   removeprefix_normalize_tyendswithpointer_type
element_tydtyper;   ru   rc   r   r   replace)tyr   s   & r'   r   r     sW   ''"cXXZ=="")Br"B==%%%%"R&= ;;sr#2w///==r"v...== !788  	B))	*	*=/011	B

	#	#WW	B		[[W%))"**T2*>CCr*   c                      ] tR tRtRtR R lt]R 4       t]R R l4       t]R R	 l4       t	]R
 4       t
]R 4       t]R 4       t]R 4       tRtR# )KernelParami+  zBRepresents a parameter (name plus metadata) to a @jit'ed function.c               (    V ^8  d   QhRRRRRRRR/# )r"   numr/   paramzinspect.Parameterdo_not_specializebooldo_not_specialize_on_alignmentr%   )r&   s   "r'   r(   KernelParam.__annotate__.  s0     M MC M(9 Md M15Mr*   c                	6    Wn         W n        W0n        W@n        R # rM   )r   _paramr   r   )rH   r   r   r   r   s   &&&&&r'   r:   KernelParam.__init__.  s    !2.L+r*   c                	.    V P                   P                  # rM   )r   r;   rO   s   &r'   r;   KernelParam.name5  s    {{r*   c                   V ^8  d   QhRR/# r   r%   )r&   s   "r'   r(   r   :  s     5 5C 5r*   c                	    V P                   P                  '       d4   V P                   P                  \        P                  P                  8X  d   R # \        V P                   P                  4      # )rT   )r   
annotationrU   	Parameteremptyr   rO   s   &r'   r   KernelParam.annotation9  sG    {{%%%)?)?7CTCTCZCZ)ZT[[3344r*   c                   V ^8  d   QhRR/# r   r%   )r&   s   "r'   r(   r   @  s       r*   c                	    V P                   pVP                  R 4      '       d   VR,          pM VP                  R4      '       d
   VR,          pV\        \        P                  ! 4       4      9   d   V P                   # R# )r   :r"   NNr   r   rT   )r   rX   r   r   values)rH   as   & r'   annotation_typeKernelParam.annotation_type?  s_    OO<<"A\\#"A.55788??"r*   c                	     R V P                   9   # rp   )r   rO   s   &r'   is_constexprKernelParam.is_constexprJ  s    doo--r*   c                	    V P                   '       d   R # RV P                  9   ;'       g    V P                  P                  R4      # )Fr   r   )r   r   rX   rO   s   &r'   is_constKernelParam.is_constN  s7    $//)MMT__-G-G-MMr*   c                	.    V P                   P                  # rM   )r   defaultrO   s   &r'   r   KernelParam.defaultT  s    {{"""r*   c                	d    V P                   P                  \        P                  P                  8g  # rM   )r   r   rU   r   r   rO   s   &r'   has_defaultKernelParam.has_defaultX  s#    {{""g&7&7&=&===r*   )r   r   r   r   N)rc   rS   r   r   r   r:   r   r;   r   r   r   r   r   r   r  r   r%   r*   r'   r   r   +  s    LM     5 5
   . . N N
 # # > >r*   r   c                :    R pRp\        \        WW4      ^ ,          # )FT)r   r   )r   
specializer   aligns   &&  r'   mangle_typer  ]  s!    HE!+sjPQRSSr*   c                  :    ] tR tRt$ R]R&   R tR tR R ltRtR	# )
KernelInterfaceic  r   runc               	^    V P                   ! \        \        P                  V4      R VRR/VB # )gridwarmupT)r
  map
MockTensor
wrap_dtype)rH   r  r   kwargss   &$*,r'   r  KernelInterface.warmupf  s.    xxZ5J5JD1Q\T\$\U[\\r*   c               	    \        R 4      h)zrun not implemented)NotImplementedError)rH   r  r  r   r  s   &$$*,r'   r
  KernelInterface.runi  s    !"788r*   c                   V ^8  d   QhRR/# )r"   r#   r   r%   )r&   s   "r'   r(   KernelInterface.__annotate__l  s     Z Z1 Zr*   c                   a a VV 3R l# )z
A JIT function is launched with: fn[grid](*args, **kwargs).
Hence JITFunction.__getitem__ returns a callable proxy that
memorizes the grid.
c                 0   < SP                   ! V R SRR/VB # )r  r  F)r
  )r   r  r  rH   s   *,r'   <lambda>-KernelInterface.__getitem__.<locals>.<lambda>r  s    txx$'YT'Y%'YRX'Yr*   r%   )rH   r  s   ffr'   __getitem__KernelInterface.__getitem__l  s     ZYr*   r%   N)	rc   rS   r   r   __annotations__r  r
  r  r   r%   r*   r'   r	  r	  c  s    	
F]9Z Zr*   r	  c                    VP                  4        UUu/ uF  w  rWYWP                  P                  R 8X  d   \        V4      MlVP                  P                  R8X  d   RVP                  /MCVP                  P                  R8X  d(   RVP
                   RVP                  P                   2/MTbK  	  ppp^ RIpRT RTRVP                  4        U	u. uF  p	\        V	4      NK  	  up	R	\        VP                  4       4      R
VP                  4        U	u. uF  p	\        V	4      NK  	  up	R\        VP                  4       4      RVP                  RXRVP                  /	p
VP                  V
4      pV# u uppi u up	i u up	i )r   rq   JITFunctionjit_function:Nr;   	signatureconstant_keysconstant_vals
attrs_keys
attrs_valsoptionskeyr   )itemsrJ   rc   rf   r   rZ   fnr   jsonra   r3   r   __dict__dumps)r;   r#  	constantsattrsr(  r)  r   r   r,  xobjserialized_objs   &&&&&&&     r'   serialize_specialization_datar4  v  sh   
 $//+	 ,JC 	??33w>SZ??##{2 FQRWR]R]D^??##}4 :HELL>YZ[`[c[c[p[pZqIr8s:?	@ ,	   k9oQZQ_Q_Qa?bQaAQQa?bdsY ,%**,0O,Qa,0OQ]_cdidpdpdr_s7##UC6??C
 ZZ_N @c0Os   BE0E6E;c                   \        V P                  4      \        V4      8X  g   Q h. p\        V P                  P                  4       V4       EF>  w  rEVP                  '       d   VP                  RV R24       K/  VP                  '       d   RMRpVP                  '       d   RMRpVP                  '       d   RMRpRV RV RV RV R2	p	VP                  '       d   \        VP                  \        4      '       d,   VP                  R8X  g   VP                  R,          R9   d   R	pV'       d&   VP                  R
VP                   RV	 R24       EK
  VP                  R
VP                   R24       EK,  VP                  V	 4       EKA  	  R p
RRP                  \        \        WP                  P                  4       4      4      R.,           4       RRP                  V P                  P                  4        Uu. uF  pRV RV 2NK  	  up4       RRP                  V4       R2pV P                  P                  4        UUu/ uF?  w  rLVP                   \"        P$                  P&                  Jg   K/  RV 2VP                   bKA  	  ppp\(        pWR&   W-R&   \*        VR&   \-        W4       VR,          # u upi u uppi )a  
Equivalent to sig.bind followed by apply_defaults. This generates a
native Python function (using exec) which can be memoized on a per-kernel
basis to avoid having to run these expensive functions -- which constitute
much of the kernel launch overhead -- every time we run the kernel.
z("constexpr", )TrueFalsezspecialize_impl(backend, , u1:Nr"   NFz("z",) + z[1:]z", None)c                    V ^,          P                   \        P                  P                  J d
   V ^ ,          # V ^ ,           RV ^ ,           2# )r   z	=default_)r   rU   r   r   )r1  s   &r'   r  0create_function_from_signature.<locals>.<lambda>  sC    AaDLLG,=,=,C,CCAaDaAaD6QZ[\]^[_Z`Iaar*   z
def dynamic_func(z	**optionsz):
    params = {'z': z}
    specialization = [,z-]
    return params, specialization, options
default_specialize_implbackendr`   dynamic_func)fpbf)r0   
parameterszipra   r   appendr   r   r   r   r8   rf   joinr3   r  r*  r   rU   r   r   r   r`   exec)sigkparamsrA  specializationr;   kpr   r  r  rP   r   	func_bodyr   func_namespacer@  s   &&&            r'   create_function_from_signaturerP    s    s~~#g,...N++-w7???!!N4&":;!#v'H$&$8$8$8fJ!@@@GfE-dV2hZr*RPUwVWXC!!!b00#66))T1R5G5G5K|5[%*
"))Br/A/A.B&T*RS #))Br/A/A.B(*KL%%/' 8, bC))DS..*>*>*@!ABk]RST U		3>>;N;N;PQ;P4QtfCv.;PQRS Txx/0 1I >>//11KD== 1 1 7 77 	)(4&5==(1   -O(7$% '9$/N=! 	# .))+ Rs   >J;,K 9K c                8    V P                    R V P                   2# ).)rS   r   r+  s   &r'   get_full_namerT    s    mm_Aboo.//r*   c                      ] tR tRtR tR t]R R l4       tR tR t	]R 4       t
R	 R
 ltR tR tR t]! ]]R7      tRtR# )r`   i  c                	   Wn         \        P                  ! V4      V n         \        P                  ! V4      w  V n        V n        \        T4      T n	        \        P                  ! 4       T n        \        P                  ! RP                  T P                  4      4      pT\         P"                  ! RT\         P$                  4      P'                  4       R pY0n        RT n        / T n        TP.                  T n        TP0                  T n        TP2                  T n        TP4                  T n        TP6                  T n        R#   \         d   p\        R 4      ThRp?ii ; i)z1@jit functions should be defined in a Python fileNrT   z^def\s+\w+\s*\()r+  rU   r#  getsourcelinesraw_srcstarting_line_numberOSError
ValueErrorrT  _fn_name	threadingRLock
_hash_locktextwrapdedentrH  research	MULTILINEstart_srchashrF   r   rc   r   __globals__rS   )rH   r+  erI   s   &&  r'   r:   JITCallable.__init__  s    **2.	Y6=6L6LR6P3DL$3 &b)#//+ oobggdll34")).R\\BHHJKL		 TV zzOO>>--7  	YPQWXX	Ys   #E EEEc                	   V P                   pVP                  f   V P                  # \        VP                  P
                  VP                  4       UUu/ uF  w  r#W#P                  bK  	  pppV P                  V,          # u uppi rM   )r+  __closure__rh  rF  __code__co_freevarscell_contents)rH   r+  r;   cellrA   s   &    r'   get_capture_scopeJITCallable.get_capture_scope  sq    WW>>!###@CBKKD[D[]_]k]k@lm@l*$T---@l	m)++ ns   Bc                   V ^8  d   QhRR/# r   r%   )r&   s   "r'   r(   JITCallable.__annotate__  s      3 r*   c                	   V P                   ;_uu_ 4        V P                  e   V P                  uuR R R 4       # RV P                   2V n        \        P                  ! V P
                  4      P                  p\        V P                  V P                  VV P                  R7      pVP                  V P                  4       4       VP                  \        V P                  4      ,           V n        \        \!        VP"                  P%                  4       4      4      V n        ^ RIHp T ;P                  \        V P"                  P%                  4        UUUu. uF   w  w  rEw  re\+        Wc4      '       g   K  WF3NK"  	  uppp4      ,          un        \,        P.                  ! V P                  P1                  R4      4      P3                  4       V n        R R R 4       V P                  # u upppi   + '       g   i     T P                  # ; i)Nz
recursion:)r;   r@   rA   rI   rp   r,   )r_  rg  r\  rU   getclosurevarsr+  rA   r    rh  rI   r   parserP   rf   rY  dictsortedrF   r*  r   rq   r8   r<   r=   r>   rN   )rH   rA   dependencies_finderrq   r;   ri   r{   s   &      r'   re   JITCallable.cache_key  sv    ___yy$yy _
 %T]]O4DI..tww7AAI"4$--QUQaQamv9=#C%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!6II9=9N9N9T9T9V=9V"5)4Xc!+C!; *tk9V= > >I  tyy'7'7'@AKKMDI# $ yy	= _$ yys*   G,DG,G%4G%;AG,%G,,H	c                	,    \        V P                  4      # rM   )rg  re   rO   s   &r'   __hash__JITCallable.__hash__
  s    DNN##r*   c                	(   \         P                  ! V P                  4      p\        V\         P                  4      '       g   Q h\        VP                  4      ^8X  g   Q h\        VP                  ^ ,          \         P                  4      '       g   Q hV# )r   )r   rw  rf  r8   Moduler0   bodyFunctionDef)rH   trees   & r'   rw  JITCallable.parse  se    yy#$

++++499~"""$))A,8888r*   c                	    ^ RI Hp V! V 4      # )r   )constexpr_type)r   r  )rH   r  s   & r'   ru   JITCallable.type  s    7d##r*   c                    V ^8  d   QhRRRR/# )r"   handleszlist[ir.value]r#   r$   r%   )r&   s   "r'   r(   rt    s      > d r*   c                	    R # rM   r%   )rH   r  s   &&r'   _flatten_irJITCallable._flatten_ir  s    r*   c                     RV n         Wn        R# )a  
The only method allowed to modify src.
Bypasses the __setattr__ restriction by calling super().__setattr__ directly.

Note that it is the callers responsibility to make sure any triton functions that call this function have the `.hash` value reset to None.
N)rg  rf  )rH   new_srcs   &&r'   _unsafe_update_srcJITCallable._unsafe_update_src  s     		r*   c                	    \        R 4      h)zqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorrO   s   &r'   _set_srcJITCallable._set_src)  s     ( ) 	)r*   c                	    V P                   # rM   )rf  rO   s   &r'   _get_srcJITCallable._get_src.  s    yyr*   )fgetfset)r   rh  rS   rc   r   r\  r_  rf  r+  rg  rX  r#  rY  rF   N)rc   rS   r   r   r:   rq  r   re   r}  rw  ru   r  r  r  r  rI   r   r%   r*   r'   r`   r`     sa     (D,  ,$ $ $)
 x
0Cr*   r`   c                  6    ] tR tRt$ R]R&   R]R&   R]R&   RtR	# )
JitFunctionInfoi7  r   rZ   rf   r;   r   r!  r%   N)rc   rS   r   r   r  r   r%   r*   r'   r  r  7  s    
Ir*   r  c                   a \        V4      \        V4      3pV P                  VR 4      pVe   V# V3R lo\        S! V4      4      \        V4      ,           pW@V&   V# )Nc                  < \        V \        4      '       d   V  Uu. uF  pS! V4      NK  	  up# \        V 4      '       d(   V  Uu. uF  pS! V4      NK  	  ppV P                  ! V!  # \        V \        4      '       d3   \        ;QJ d    . V3R  lV  4       F  NK  	  5# ! V3R  lV  4       4      # \        V \
        4      '       d   V P                  # V # u upi u upi )c              3  4   <"   T F  pS! V4      x  K  	  R # 5irM   r%   ).0r   replace_callabless   & r'   	<genexpr>?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>L  s     ?3C*3//3s   )r8   r3   r   rJ   tupler`   re   )r2  r   resultsr  s   &  r'   r  ,compute_cache_key.<locals>.replace_callablesE  s    c4  69:cs%c*c::39<=#(-G==='**U##5?3?5?5?3???[))== 
 ;=s   CC)r  rf   r   )kernel_key_cacherL  r(  r)  re   r  s   &&&  @r'   compute_cache_keyr  >  sb     #g,
/C $$S$/I
 %n56WEI%Sr*   c                    \        V \        4      '       g   V # \        V 4       F  w  r\        V4      W&   K  	  \	        V 4      # rM   )r8   r3   	enumerateconvert_to_tuple_if_listr  )iteminested_values   &  r'   r  r  V  s>    dD!! %T?*<8 + ;r*   c                  t   a  ] tR tRtR tR R ltR tR tR tR t	R	 t
RV 3R
 lltR tR tR tR tRtV ;t# )r   ib  c                	    R # Fr%   rO   s   &r'   is_gluonJITFunction.is_gluond  s    r*   c                   V ^8  d   QhRR/# )r"   r#   zbool | Noner%   )r&   s   "r'   r(   JITFunction.__annotate__g  s     .
 .
 
.
r*   c
                	   V'       g   R # V P                   P                  p
V P                   P                  pRP                  \	        V P
                  V^,          4       UUu. uF  w  rVP                   RV 2NK  	  upp4      pV
 RVP                   RVP                   RVP                   RVP                   RVP                   RV R	2p\        V P                   4      p\        VW6V^ ,          WrV4      pR
VRVRVRVP                  RVP                  RVP                  RVP                  RVP                  RVP                  RVRVRV	/pV! VV\        WV 4      RV/VCV	RR7      # u uppi )Nr9  z: z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=z, launch_cooperative_grid=](r6  r#  devicer/  	num_warpsnum_ctas
num_stagesenable_fp_fusionlaunch_cooperative_gridextern_libsconfigsspecialization_data	is_warmupr)  F)r)  reprr+  compileis_manual_warmupalready_compiled)r+  r   rS   rH  rF  paramsr;   r  r  r  r  r  rT  r4  r  r  )rH   hookr)  r#  r   r  r/  r(  r  r  r;   rZ   r   r   	arg_reprsr  	full_namer  r  s   &&&&&&&&&&         r'   
_call_hookJITFunction._call_hookg  s    ww####IIc$++WZ[\W]F^_F^%**Rt4F^_`	{7#4#4"5[AQAQ@RR_`g`r`r_s  tG  HO  H`  H`  Ga  a{  |C  |[  |[  {\  \^  _h  ^i  ij  k!$''*	;Iy]def]gip<BD f**((',, 8 8%w'F'F7..w!#6
 vT2C*6*&"
 	
- `s   "E,
c                b    \        V4      '       g   Q hV P                  P                  V4       R# )zu
Add a hook that will be executed prior to the execution of run
function with args and kwargs passed into the kernel
N)rx   pre_run_hooksrG  )rH   r  s   &&r'   add_pre_run_hookJITFunction.add_pre_run_hook  s&    
 ~~~!!$'r*   c                    ^RI HpHpHpHp \
        P                  P                  4       pV! V4      pWn        W n        W0n        \        V P                  V P                  V4      p/ / WVV3# )z!
Precompute as much as possible.
)CompiledKernelr  	ASTSourcemake_backend)compilerr  r  r  r  r   activeget_current_targetrP  r#  r  )rH   r  r  r  r  r   rA  binders   &       r'   create_binderJITFunction.create_binder  s\     	PO113v&,"/WU2v..r*   c           
     	   VP                  V4      pV P                   Uu. uF  qfP                  NK  	  ppV Uu. uF  qf^ ,          NK  	  pp\        Wx4       U	U
u/ uF  w  rWbK	  	  pp	p
RV9  g   Q R4       hRV9  g   Q R4       hRV9  g   Q R4       hV F-  p	WP                  9  g   K  W9  g   K  \        RV	,          4      h	  \        VR 4      pV Uu/ uF&  q\        \        VP                  4       4      V4      bK(  	  ppV Uu. uF  qf^ ,          R	8X  d   R
MV^,          NK  	  pp\        VR 4      pV U	u/ uF  qVP                  \        W4      4      bK  	  pp	W[W3# u upi u upi u up
p	i u upi u upi u up	i )r   device_typez=device_type option is deprecated; current target will be usedr  z8device option is deprecated; current device will be usedstreamz8stream option is deprecated; current stream will be usedz2Keyword argument %s was specified but unrecognisedc                    VR 8H  # rp   r%   )ri   r{   s   &&r'   r  (JITFunction._pack_args.<locals>.<lambda>  s	    3+;Mr*   rq   rT   c                "    \        V\        4      # rM   )r8   rf   )ri   r1  s   &&r'   r  r    s    Z35Gr*   )parse_optionsr  r;   rF  r-  KeyErrorr   r   r3   r   
parse_attr)rH   rA  r  
bound_argsrL  r(  r1  sigkeyssigvalsrg   vr#  
constexprspathattrvalsr0  s   &&&&&&          r'   
_pack_argsJITFunction._pack_args  s   ''/#';;/;a66;/!/0AQ440(+G(=>(=fqQT(=	>F*k,kk*v%a'aa%v%a'aa%A(((Q-=SVWWXX  #7,MN
[ef[eSW-d:3D3D3F.GNN[e
fAOPAA$+-B1Q47Ph(GHPUVPU1G&&'8'EFFPUV:44% 00> gPVs"   E(E-E2,E8!E=?#Fc               	   VP                  R V P                  4      ;'       g    \        P                  P                  VR &   \        P                  P
                  VR&   \        P                  P                  4       p\        P                  P                  V4      pV P                   F  pV! V/ VB  K  	  V P                  V,          w  rrpV! V/ VB w  rp\        P                  P                  e7   \        P                  P                  4       w  ppVP                  RV R24       \        WV4      pVP                  VR4      pVf6   V P                  WWV4      w  ppppV P!                  VVVVVVV4      pVf   R# \#        4       pV P$                  P'                  4        F9  w  w  ppw  ppVP                  VV4      ;pV8w  g   K'  \)        RV RV RV 24      h	  V'       g   Vf   Q h\+        V4      '       d	   V! V4      p\-        V4      pV^ ,          pV^8  d
   V^,          M^pV^8  d
   V^,          M^p VP.                  ! W.VP1                  4       O5!  p!VP2                  ! VVV VVP4                  VP6                  V!\        P                  P8                  \        P                  P:                  .	VP1                  4       O5!   V# )debuginstrumentation_modeNz("custom_pipeline", r6  r^   z1 has changed since we compiled this kernel, from z to )r   r  r   runtimecompilationr  r   r  get_current_deviceget_current_streamr  device_cachesadd_stages_inspection_hookrG  r  r  _do_compileobjectrF   r*  rb   rx   r0   launch_metadatar   r
  functionpacked_metadatalaunch_enter_hooklaunch_exit_hook)"rH   r  r  r   r  r  r  r  kernel_cacher  r   rA  r  r  rL  r(  inspect_stages_keyinspect_stages_hashr)  kernelr#  r  r0  not_presentr;   ri   r{   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  s"   &$$*,                             r'   r
  JITFunction.run  s    **Wdjj9PPU]]=P=Pw).):):)O)O%& 11311&9 &&D$!&! ' CGBTBTU[B\? /5d.Ef.E+
G ==33?6;mm6^6^6`3 3!!$89L8MQ"OP 0'J!!#t, >48OOGU_DK5M1GY
E %%c9fj'SXZ`aF~ h.2.C.C.I.I.K*IT1*\&**4==#E"&tf,]^a]bbfgmfnoq q /L
 ###~~J'D	I!WF )AT!W1F )AT!W1F$44TXJDUDUDWXOJJvvvvvH^H^`o}}668V8VnYcYjYjYlnr*   c                	X    V P                   f   V P                  # V P                  V4      # rM   )_reprr\  )rH   ri   s   &&r'   r  JITFunction.repr  s"     $

 2t}}E

1Er*   c	           	     	  < V'       d   TM. pV'       d   TM. p\         SV `  V4       VP                  V n        W n        W0n        W@n        Wpn        Wn        V \        V P                   R V P                  P                   2&   . V n        \        V P                  P                  P!                  4       4       F`  w  rW9   ;'       g    V
P"                  V9   pW9   ;'       g    V
P"                  V9   pV P                  P%                  \'        WW4      4       Kb  	  \)        V P*                  4      V n        RV n        WPn        W`n        V P                   Uu. uF  qP"                  NK  	  upV n        V P                   Uu. uF"  qP6                  '       g   K  VP8                  NK$  	  upV n        . V n        R# u upi u upi )r"  N)r9   r:   rS   rZ   versionr   r   r  r  _triton_jit_function_registryr+  r   r  r  r#  rE  r   r;   rG  r   r   r  r  r  r  r_   	arg_namesr   r   r  r  )rH   r+  r  r   r   r  r_   r  r  r  r   dnsdns_oaprJ   s   &&&&&&&&&     r'   r:   JITFunction.__init__   sh   1B-Ki)Goq&mm!2.L+
.QU%Qtww7K7K6L&MN!$..";";"B"B"DEHA(KKEJJ:K,KC8hhEJJJh<hFKK{1SAB F )););< 
  +/++6+Q&&+6*.++H+Q5155+H  	 7Hs   F:F?F?c           
     	  a ^ RI p^ RIHo \        P                  P                  4       pVP                  V4      pVR,          V P                  8w  d#   \        RVR,           RV P                   24      h\        \        VR,          4      pVR,          pV P                  V,          w   rxrVR,          p
VP                  V
8w  d   \        RV
 RV 24      hV3R lp\        WV4       UUu/ uF  w  rW! V4      bK  	  ppp\        \        VR	,          4      pVR
,          p\        \        VV4      4      pVR,          P                  4        UUu/ uF  w  rV\!        V4      bK  	  pppVR,          P                  4        UUu/ uF)  w  rT\#        V\$        4      '       d   \        V4      MTbK+  	  pppVR,          pV	P'                  V4      pV P)                  VVVVVVRR7      # u uppi u uppi u uppi )r   Nr;   zSpecialization data is for z but trying to preload for r$  r%  r   c                D  < SP                   P                  V 4      '       d   SP                  V 4      # \        V \        4      '       dW   R V 9   d   SP	                  V R ,          4      # RV 9   d1   V R,          pV\
        9   d   \
        V,          # \        RV R24      hV # )rq   r!  zUnable to resolve JITFunction z for preload)r   is_dtyper8   rx  rq   r  rb   )r   jf_keytls   & r'   _decode_constant-JITFunction.preload.<locals>._decode_constant7  s    xx  ''xx&%&&%'<<k(:;;!U*">2F!>><VDD&)Gx|'\]]Lr*   r&  r'  r#  r(  r)  T)r  )r,  triton.languager   r   r  r  loadsr\  rb   r  r  r  r-  rF  rx  r*  r  r8   r3   r  r  )rH   r  r,  r  deserialized_objr$  r%  ri   r   rA  deserialized_targetr  r)  r   r  r&  r'  r0  r#  r(  r  s   &&                  @r'   preloadJITFunction.preload&  s   $113::&9:F#t}}4-.>v.F-GGbcgcpcpbqrt tE#3O#DE(9#'#5#5f#= 1g.x8 ??11!<=P<QQlmsltuvv	 FIEfgEfzsc+E22Ef
g 0 >?
%l3
SZ01 M]]hLiLoLoLqrLqjcS2599Lq	r /y9??A
A
 E4!8!8ueCA 	 
 u%''0   
 	
 h s
s   +G(G./G4c                	l  a aaaaaaaaaaa S P                   S,          w  oporS P                  \        P                  P                  SSSSSSS.S4	      '       d   R # S P                  S SSS4      o\        P                  P                  4       p
V
eK   \        4       o\        SV	SS4      pVVV VV3R lpVVVVVVV VVV3
R lpV
P                  WV4      pVSS&   V# S P                  SSSP                  R7      pVSS&   S P                  \        P                  P                  SSSSSSS.S4	       V# )Nc                 B   < SP                  SSSP                  S R 7      # ))r   r(  	_env_vars)r  r-  )env_varsr(  rH   rI   r   s   r'   async_compile.JITFunction._do_compile.<locals>.async_compilei  s!    ||C@P@P\d|eer*   c                v   <
 V SS&   SP                  \        P                  P                  SSS	SSSS.S
4	       R # rM   )r  r   r  jit_post_compile_hook)r  r0  r  r  r  r)  r(  rH   r#  r   r  s   &r'   finalize_compile1JITFunction._do_compile.<locals>.finalize_compilel  s;    $*S! C CS)U[]ceo '%&:r*   )r   r(  )r  r  r   r  jit_cache_hookr  r   active_moder   r   r   submitr  r-  r&  )rH   r)  r#  r  r  r(  r0  r  ri   rA  
async_modere   r#  r'  r  r"  r  rI   r   s   ffffffff       @@@@r'   r  JITFunction._do_compile[  s1   .2.@.@.H+a??5==77iQWYcelotnu!# #nnT9j%@#//335
!68H%c7GXFIf f: :
  &&yAQRF &L 	 \\#fg>N>N\OF &LOOEMM??iQWY_akmt"GV-r*   c                	    \        R 4      h)z:Cannot call @triton.jit'd outside of the scope of a kernel)rb   rH   r   r  s   &*,r'   __call__JITFunction.__call__z  s    WXXr*   c                	P    R V P                    RV P                  P                   R2# )zJITFunction(r"  r6  )rZ   r+  r   rO   s   &r'   __repr__JITFunction.__repr__}  s&    dkk]!DGG,@,@+ACCr*   )r  r  r  r  r  r  r  r  r   r   r  r  rZ   r_   r  r  r  )NNNNNNN)rc   rS   r   r   r  r  r  r  r  r
  r  r:   r  r  r0  r3  r   r   r   s   @r'   r   r   b  sL    .
`(/507rF$ L3
j>YD Dr*   r   c                    V ^8  d   QhRRRR/# r"   r+  r   r#   zJITFunction[T]r%   )r&   s   "r'   r(   r(     s      A . r*   c                    R # rM   r%   rS  s   &r'   jitr8    s    r*   r  r  r  r   r   r  r_   c               4    V ^8  d   QhRRRRRRRRRRR	RR
R/# )r"   r  Optional[Callable]r  r   Optional[Iterable[int | str]]r   r  Optional[bool]r_   r#   zCallable[[T], JITFunction[T]]r%   )r&   s   "r'   r(   r(     sS     
 
 
 (	

 5
 %B
 
 
 #
r*   c                    R # rM   r%   )r  r  r  r   r   r  r_   s   $$$$$$$r'   r8  r8    s     r*   c               8    V ^8  d   QhRRRRRRRRRRR	R
RR
RR/# )r"   r+  zOptional[T]r  r:  r  r   r;  r   r  r<  r_   r#   zKernelInterface[T]r%   )r&   s   "r'   r(   r(     s[     4 44 	4
 (4 54 %B4 4 4 4r*   c               J   aaaaaaa R VVVVVVV3R llpV e	   V! V 4      # V# )a  
Decorator for JIT-compiling a function using the Triton compiler.

:note: When a jit'd function is called, arguments are
    implicitly converted to pointers if they have a :code:`.data_ptr()` method
    and a `.dtype` attribute.

:note: This function will be compiled and run on the GPU. It will only have access to:

       * python primitives,
       * builtins within the triton package,
       * arguments to this function,
       * other jit'd functions

:param fn: the function to be jit-compiled
:type fn: Callable
c                    V ^8  d   QhRRRR/# r6  r%   )r&   s   "r'   r(   jit.<locals>.__annotate__  s      a N r*   c                   < \        V 4      '       g   Q h\        P                  P                  '       d   ^RIHp V! V SSSSSSSR7      # \        V SSSSSSSR7      # )r   )InterpretedFunction)r  r   r   r  r_   r  r  )rx   r   r  	interpretinterpreterrC  r   )	r+  rC  r  r   r   r  r_   r  r  s	   & r'   	decoratorjit.<locals>.decorator  so    |||=="""8&r7N_Fdlq08tUdf f "3/M! /	 	r*   r%   )	r+  r  r  r  r   r   r  r_   rF  s	   &ddddddd r'   r8  r8    s&    : & 
~} r*   c                  Z    ] tR tRtRt]R 4       tR
R ltR t]R 4       t	]R 4       t
R	tR# )r  i  zf
Can be used in place of real tensors when calling:
    kernel.warmup(MockTensor(torch.float32), ...)
c                	t    V P                   P                  R 8X  d   V P                  R8X  d   \        V 4      # V # )r   torch)rJ   rc   rS   r  )r   s   &r'   r  MockTensor.wrap_dtype  s.    ==!!W,71Jc?"
r*   Nc                	,    Vf   ^.pWn         W n        R # rM   r   shape)rH   r   rN  s   &&&r'   r:   MockTensor.__init__  s    =CE

r*   c                	    ^.pV P                   R,           F"  pVP                  VR,          V,          4       K$  	  \        \        V4      4      # )r   r   r   )rN  rG  r  reversed)rH   stridessizes   &  r'   strideMockTensor.stride  sB    #JJrNNDNN72;-. #Xg&''r*   c                 	    ^ # r   r%   r%   r*   r'   data_ptrMockTensor.data_ptr      r*   c                 	    ^ # rW  r%   r%   r*   r'   	ptr_rangeMockTensor.ptr_range  rZ  r*   rM  rM   )rc   rS   r   r   r   staticmethodr  r:   rT  rX  r\  r   r%   r*   r'   r  r    sM    
  
(    r*   r  c                  X    ] tR tRtR tR tR tR R ltR tR t	R	 t
R
 tR tR tRtR# )TensorWrapperi  c                	    W n         Wn        VP                  V n        VP                  V n        V P                  P                  V n        R # rM   )r   basedatar  rN  )rH   rb  r   s   &&&r'   r:   TensorWrapper.__init__  s1    
	II	kkYY__
r*   c                	6    V P                   P                  4       # rM   )rb  rX  rO   s   &r'   rX  TensorWrapper.data_ptr  s    yy!!##r*   c                	6    V P                   P                  ! V!  # rM   )rb  rT  )rH   r   s   &*r'   rT  TensorWrapper.stride  s    yy&&r*   c                   V ^8  d   QhRR/# r   r%   )r&   s   "r'   r(   TensorWrapper.__annotate__  s     ; ; ;r*   c                	<    R V P                    RV P                   R2# )zTensorWrapper[r  r6  )r   rb  rO   s   &r'   __str__TensorWrapper.__str__  s    

|2dii[::r*   c                	6    V P                   P                  4       # rM   )rb  element_sizerO   s   &r'   ro  TensorWrapper.element_size  s    yy%%''r*   c                	^    \        V P                  P                  4       V P                  4      # rM   )r`  rb  cpur   rO   s   &r'   rr  TensorWrapper.cpu  s    TYY]]_djj99r*   c                	P    V P                   P                  VP                   4       R # rM   )rb  copy_)rH   others   &&r'   ru  TensorWrapper.copy_  s    		

#r*   c                	^    \        V P                  P                  4       V P                  4      # rM   )r`  rb  cloner   rO   s   &r'   ry  TensorWrapper.clone  s    TYY__.

;;r*   c                	`    \        V P                  P                  V4      V P                  4      # rM   )r`  rb  tor   )rH   r  s   &&r'   r|  TensorWrapper.to  s     TYY\\&14::>>r*   c                	`    \        V P                  P                  V4      V P                  4      # rM   )r`  rb  	new_emptyr   )rH   sizess   &&r'   r  TensorWrapper.new_empty  s"    TYY007DDr*   )rb  rc  r  r   rN  N)rc   rS   r   r   r:   rX  rT  rl  ro  rr  ru  ry  r|  r  r   r%   r*   r'   r`  r`    s5    %$';(:$<?Er*   r`  c                   \        V \        4      '       d=   WP                  P                  8X  d   V P                  # \        V P                  V4      # \	        V R 4      '       d   \        W4      # \        R\        V 4       R24      h)rX  zCannot reinterpret a rR  )r8   r`  rb  r   hasattrr   ru   )tensorr   s   &&r'   reinterpretr    sm    &-((KK%%%;; !e44		$	$V++/V~Q?@@r*   c                J   T p\        V\        4      '       g   VP                  pK$  VP                  P                  P                  pVP
                  p\        VP                  4       F6  w  rEVP                  4       P                  R 4      '       g   K,  W4,          p W#3# 	  W#3# )zdef )
r8   r`   r+  rm  co_filenamerY  r  rX  r   rX   )r+  base_fn	file_name
begin_lineidxlines   &     r'   get_jit_fn_file_liner  *  s    G+..**

##//I--J w/	::<""6**J  	 0   r*   c                  2    ] tR tRtR t]R 4       tR tRtR# )BoundConstexprFunctioni<  c                	    Wn         W n        R # rM   )__self____func__)rH   instancer+  s   &&&r'   r:   BoundConstexprFunction.__init__>  s     r*   c                	.    V P                   P                  # rM   )r  re   rO   s   &r'   re    BoundConstexprFunction.cache_keyB  s    }}&&&r*   c                	B    V P                   ! V P                  .VO5/ VB # rM   r  r  r/  s   &*,r'   r0  BoundConstexprFunction.__call__F  s    }}T]]<T<V<<r*   r  N)	rc   rS   r   r   r:   r   re   r0  r   r%   r*   r'   r  r  <  s      ' '=r*   r  c                  <   a  ] tR tRtV 3R ltR tRR/R ltRtV ;t# )ConstexprFunctioniJ  c                	&   < \         SV `  V4       R # rM   )r9   r:   )rH   r+  rJ   s   &&r'   r:   ConstexprFunction.__init__L  s    r*   c                	$    Ve   \        W4      # V # rM   )r  )rH   r2  objclasss   &&&r'   __get__ConstexprFunction.__get__O  s    ?)#44r*   	_semanticNc               	*   ^ RI HpHp V Uu. uF
  qd! V4      NK  	  ppVP                  4        UUu/ uF  w  rxWt! V4      bK  	  pppV P                  ! V/ VB p	Vf   V	# \
        P                  P                  '       d   V	# V! V	4      # u upi u uppi )r   )_unwrap_if_constexprrq   )r   r  rq   r*  r+  r   r  rD  )
rH   r  r   r  r  rq   r1  rg   r  ress
   &$*,      r'   r0  ConstexprFunction.__call__U  s    H156A$Q'6;A<<>J>!!)!,,>J ggt&v&J =="""J~ 7Js
   B
Br%   )	rc   rS   r   r   r:   r  r0  r   r   r   s   @r'   r  r  J  s      r*   r  c                    \        V 4      # )z
Wraps an arbitrary Python function so that it can be called at
compile-time on constexpr arguments in a Triton function and
returns a constexpr result.
)r  rS  s   &r'   constexpr_functionr  h  s     R  r*   r  rM   )H
__future__r   r   r   r-   r<   rU   r   r]  rb  r`  collectionsr   dataclassesr   	functoolsr   typingr   r	   r
   r   r   r   r   r   r   triton.backendsr   typesr   rT   r   r   r   _utilsr   r   r   r   cacher   triton._C.libtritonr   r   r   rD   rC   r   NodeVisitorr    r   r   r  r	  r4  rP  rT  r`   r  r  r  r  r   r8  r  r`  r  r  r  r  r  r%   r*   r'   <module>r     s    , 
      	  # ! % ] ] ] '     ` `   [ [!3CLg! g!^D4/> />dTZgaj Z&$9*x0i1 i1X !#    0	\D+q1 \DH	 
 
 


  $
 +/	

 8<
 EI
 !
  $
 

4 4  $	4
 +/4 8<4 EI4 !4  $4 4x B"E "EJA!$=[ = <!r*   