
    "hID                       d dl mZ d dlZd dlZddlmZmZ ddlmZ ddlm	Z	m
Z
 ddlmZ ddlmZ dd	lmZmZmZ dd
lmZ ddlmZ ddlmZ d dlmZ d dlZd dlZd dlZdZdZeeedZ dZ!dZ"e!e!e"dZ#d Z$d%dZ% G d d      Z& G d d      Z' ejP                         d        Z)d Z*d&dZ+d'dZ,d Z- G d d       Z. G d! d"e/      Z0 G d# d$      Z1y)(    )annotationsN   )get_cache_invalidating_env_varsir)backends)	GPUTargetAttrsDescriptor)__version__)OutOfResources)get_cache_managerget_dump_managerget_override_manager)driver)get_sass   )ast_to_ttir)Pathz^\s*tt\.func\s+(?:public\s+)?(@\w+)(\((?:%\w+: [\S\s]+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*(attributes \{[\S\s]+\})?\s+\{\s*$z=\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\))ttirttgirptxz-%\w+: ((?:[^,\s<)]+|<[^>]+>)+(?: {[^}]+})?),?z\.param\s+\.(\w+)c                    t        j                  d|       }t        j                  d|       }|yt        j                  dd|       } |dt        |j	                  d            z   S | S )Nz!tt\.ptr<([^,]+)ztt.nv_tma_desc = 1	nvTmaDescz {[^}]+} *r   )researchsubconvert_type_reprgroup)xmatchtmas      U/var/www/html/sandstorm/venv/lib/python3.12/site-packages/triton/compiler/compiler.pyr   r   ,   sd     II)1-E
)))1
-C

{B"A&u{{1~666H    c                |    d}t        j                  ||       }t        |      dk(  sJ d       t        |d         }|S )Nz&"triton_gpu.num-warps"\s?=\s?(\d+)\s?:r   z(Expected exactly one match for num_warpsr   )r   findalllenint)srcttgir_num_warps_patternnum_warps_matches	num_warpss       r#   _get_num_warps_from_ir_strr-   9   sL    G 

#:C@ !Q&R(RR&%a()Ir$   c                  (    e Zd ZdddZd Zd Zd Zy)	ASTSourceNc                   || _         d| _        |j                  | _        || _        || _        || _        t        | j                  t              rLt        | j                  j                  d            D ci c]  \  }}||j                          c}}| _        n:| j                  j                         D ]  }t        |t              rt        d       | j
                  i | _        n:| j
                  j                         D ]  }t        |t              rt        d       | j                  t               | _        y y c c}}w )Nr   ,zSignature keys must be stringzConstants keys must be string)fnext__name__name	signature	constantsattrs
isinstancestr	enumeratesplitstripkeys	TypeErrorr	   )selfr2   r6   r7   r8   kvs          r#   __init__zASTSource.__init__E   s   KK	""
dnnc*7@AUAUVYAZ7[\tq!al\DN^^((* E!!S)#$CDDE >>!DN^^((* E!!S)#$CDDE ::(*DJ  ]s   6E c                   t        | j                  j                               D cg c]  \  }}|	 }}}t        d | j                  j                         D              }| j                  j
                   d| j                  j                          d| d| }t        j                  |j                  d            j                         S c c}}w )Nc              3  <   K   | ]  \  }}t        |      |f  y wN)r:   ).0rA   rB   s      r#   	<genexpr>z!ASTSource.hash.<locals>.<genexpr>_   s     !Q$!Q3q61+!Qs   -utf-8)sortedr6   itemsr7   r2   	cache_keyr8   hashhashlibsha256encode	hexdigest)r@   rA   rB   
sorted_sigsorted_constantskeys         r#   rN   zASTSource.hash[   s    $*4>>+?+?+A$BCDAqaC
C "!Q$..:N:N:P!QQ""#1TZZ__%6$7qAFVEWX~~cjj12<<>> Ds   Cc                8    t        | j                  | ||||      S )N)contextoptionscodegen_fns
module_map)r   r2   )r@   rX   rY   rZ   rW   s        r#   make_irzASTSource.make_irc   s!    477D'7Xc&02 	2r$   c                    t               S rF   )dictr@   s    r#   parse_optionszASTSource.parse_optionsg   s	    vr$   NNreturnNoner4   
__module____qualname__rC   rN   r[   r_    r$   r#   r/   r/   C   s    +,?2r$   r/   c                  $    e Zd Zd Zd Zd Zd Zy)IRSourcec                   || _         t        |      }|j                  dd  | _        |j	                         | _        t        j                  t        | j                     | j
                  t        j                        }|j                  d      | _        |j                  d      }t        j                  t        | j                     |      }t        |      D ci c]  \  }}|t        |       c}}| _        y c c}}w )Nr   r   )pathr   suffixr3   	read_textr)   r   r   prototype_pattern	MULTILINEr   r5   r&   arg_type_patternr;   r   r6   )r@   rk   r!   r6   typesrA   tys          r#   rC   zIRSource.__init__m   s    	Dz;;qr?>>#		+DHH5txxNKKN	KKN	

+DHH5yA@I%@PQuq"!.r22QQs   C8c                z    t        j                  | j                  j                  d            j	                         S )NrJ   )rO   rP   r)   rQ   rR   r^   s    r#   rN   zIRSource.hashx   s'    ~~dhhoog67AACCr$   c                T    t        j                  | j                  |      }||_        |S rF   )r   parse_mlir_modulerk   rW   )r@   rX   rY   rZ   rW   modules         r#   r[   zIRSource.make_ir{   s$    %%dii9 r$   c                b    | j                   dk(  rdt        | j                        iS t               S )Nr   r,   )r3   r-   r)   r]   r^   s    r#   r_   zIRSource.parse_options   s*    88w!;DHH!EFFvr$   Nrd   rg   r$   r#   ri   ri   k   s    	RD
r$   ri   c                 4   dd l } t        j                  j                  t        j                  j                  t        j                  j	                  t
                          }g }t        t
        d      5 }|t        j                  |j                               j                         gz  }d d d        t        j                  j                  |d      dft        j                  j                  |d      dfg}|D ]  \  }}| j                  |g|      D ]y  }t        |j                  j                  |j                        j                   d      5 }|t        j                  |j                               j                         gz  }d d d        {  t        j                         }t        t        j                  j                  |d      d      5 }	 |j                  d	      }	|	sn|j#                  |	       &	 d d d        |j%                  |j                                t        j                  j                  |d
      }
| j                  |
gd      D ]y  }t        |j                  j                  |j                        j                   d      5 }|t        j                  |j                               j                         gz  }d d d        { t&         dj                  |      z   S # 1 sw Y   ?xY w# 1 sw Y   xY w# 1 sw Y   xY w# 1 sw Y   xY w)Nr   rbcompilerztriton.compiler.r   ztriton.backends.)prefixz_C/libtriton.soi   languageztriton.language.rI   )pkgutilosrk   dirnameabspath__file__openrO   rP   readrR   joinwalk_packagesmodule_finder	find_specr5   originupdateappendr
   )r}   TRITON_PATHcontentsfpath_prefixesrk   r{   liblibtriton_hashchunklanguage_paths              r#   
triton_keyr      s   ''//"''//"''//(2K"LMKH	h	 ;W^^AFFH-779::; 
k:	.0BC	k:	.0BCM & Cf(($(? 	CCc''11#((;BBDI CQW^^AFFH5??ABBC C	CC ^^%N	bggll;(9:D	A )QFF7OE!!%(	  	) OON,,./GGLLj9M$$m_=O$P ?###--chh7>>E 	?1;;=>>H	? 	?? ]chhx0005; ;C C
) )	? 	?s0   46K'6K4(L6L'K14K>LL	c                    |dk(  s|dk(  rt        j                  | |      }||_        |S |dk(  s|dk(  rt        |       j	                         S |dk(  rt        |       j                         S y )Nr   r   llirr   cubin)r   ru   rW   r   rm   
read_bytes)	full_namer3   rW   rv   s       r#   parser      sl    
f}w%%i9 
f}uI((**
g~I))++ r$   c                   t        j                  dd      dk(  ry| j                  t        | j                         | j                  t        | j                         ddg}| j
                  g }4t        fd|D              s|j                         j                  4t        ||dd       D ]  \  }}||_         |sd| _        yd|d	   _        |d
   | _        y)z
    Removes code_generator.py and related files from tracebacks.

    These are uninteresting to the user -- "just show me *my* code!"
    TRITON_FRONT_END_DEBUGGING01Nz"/triton/compiler/code_generator.pyz/ast.pyc              3     K   | ]6  }j                   j                  j                  j                  |      s3| 8 y wrF   )tb_framef_codeco_filenameendswith)rG   r   tbs     r#   rH   z#filter_traceback.<locals>.<genexpr>   s/     V2;;+=+=+I+I+R+RST+U1Vs   4??r   r   )
r~   getenv	__cause__filter_traceback__context____traceback__anyr   tb_nextzip)e	BAD_FILESframes	cur_frame
next_framer   s        @r#   r   r      s     
yy-s3s:{{%}} ' 	-I
 
BF
.ViVVMM"ZZ .
 $'vvabz#: 'J&	' !r
 )r$   c                R	   |t         j                  j                         }t        |t              sJ d       t        |      }t        | t               }|r"t        | t              sJ d       t        |       } | j                         }|j                  t        |xs
 t               fi |      }t               }t                d| j                          d|j                          d|j                          dt        t        |j                                      	}t!        j"                  |j%                  d            j'                         }t)        |      }	t*        j,                  j/                  dd      dk(  }
t*        j,                  j/                  dd      dk(  }|
rt1        | j                               nd }|rt3        | j                               nd }| j4                  d d	 }| d
}|	j7                  |      xs i }|j/                  |      }t*        j,                  j/                  dd      dk(  }|s;|9t9        j:                  t=        |      j?                               }tA        | ||      S ||d|jB                  |}t               }|jE                  ||       tG        |jI                               jK                  | jL                        }|r|dz  }tO        jP                         }tO        jR                  |       |jS                  |       |jU                         }|jW                         }	 | jY                  ||||      }t*        j,                  j/                  dd       }tG        |j                               |d  D ]  \  }} |||      }| d| }|.|j_                  |      x} ta        d|         tc        | ||      }|	je                  ||      ||<   ||je                  ||       ||k(  r0|	j_                  |      }!|jg                  |!       ta        d|!        |} |	je                  t9        jh                  |tj              |d      ||<   |	jm                  ||       |jo                          tA        | ||      S # tZ        $ r}t]        |        d }~ww xY w)Nz target must be of GPUTarget typez'source must be either AST or a filepathrI   rJ   TRITON_KERNEL_OVERRIDEr   r   TRITON_KERNEL_DUMP   .jsonTRITON_ALWAYS_COMPILE)rN   targetr   
USE_IR_LOC.z
Overriding kernel with file zCreating new locations for )defaultF)binary)8r   activeget_current_targetr9   r   make_backendr/   r:   ri   r_   r]   r   r   rN   rK   rL   rO   rP   rQ   rR   r   r~   environgetr   r   r5   	get_groupjsonloadsr   rm   CompiledKernel__dict__
add_stageslistr>   indexr3   r   rW   load_dialectsget_codegen_implementationget_module_mapr[   	Exceptionr   get_fileprintr   putcreate_location_snapshotdumpsvars	put_groupdisable_multithreading)"r)   r   rX   backend	ir_sourceextra_optionsenv_varsrU   rN   fn_cache_managerenable_overrideenable_ir_dumpfn_override_managerfn_dump_manager	file_namemetadata_filenamemetadata_groupmetadata_pathalways_compilemetadatastagesfirst_stagerW   rY   rZ   rv   r   
use_ir_locr3   
compile_irnext_moduleir_filenamer   ir_full_names"                                     r#   compiler      sN   ~113fi(L*LL(6"GsI..I#s#N%NN#sm%%'M##D):DF$Lm$LMG.0H\N!CHHJ<q(8',,.9I3vV^VdVdVfOgKhJi
jC>>#**W-.88:D(. jjnn%=sCsJOZZ^^$8#>#EN>M.sxxz:SW6D&sxxz2$O
 #I$+U+%//0ABHbN"&&'89MZZ^^$;SASHNm7::d=1;;=>c>488  

 	H VFvw'v{{}%++CGG4KqjjlGW'"446K'')JWk:wG d3J/= Z 2"1SE*+>Q>Z>Z[f>g1g0t29+>?	38K&6&:&:;&T{#&[9+44[AL00>/~>?  )9(<(<TZZZ^=_arDI )= )KN$%0.A
 ""$#~t44;  s   5R 	R&R!!R&c                   t        j                         D cg c]*  }|j                  j                  |       s|j                  , }}t	        |      dk7  r't        t	        |       d| j                   d| d       |d   |       S c c}w )Nr   z! compatible backends for target (z) (z). There should only be one.r   )r   valuesrz   supports_targetr'   RuntimeErrorr   )r   r    activess      r#   r   r   2  s    #+??#4[a

8R8RSY8Zqzz[G[
7|q7|n=fnn=MSQXPYYuvx 	x71:f	 \s
    BBc                       e Zd Zd ZddZd Zy)LazyDictc                     || _         g | _        y rF   )dataextras)r@   r   s     r#   rC   zLazyDict.__init__<  s    	r$   c                    | j                   D ]  \  }}| j                   || z  | _         | j                   j                          | j                  S rF   )r   r   clearr@   funcargss      r#   r   zLazyDict.get@  sG    ++ 	0JD$		D$K/DI	0yyr$   c                >    | j                   j                  ||f       y rF   )r   r   r   s      r#   addzLazyDict.addF  s    D$<(r$   Nra   )r4   re   rf   rC   r   r   rg   r$   r#   r   r   :  s    )r$   r   c                      e Zd Zd Zy)AsmDictc                T    |dk(  rt        | d         }nt        d|z        || |<   |S )Nsassr   zUnknown key: '%s')r   KeyError)r@   rU   values      r#   __missing__zAsmDict.__missing__L  s6    &=T']+E.455S	r$   N)r4   re   rf   r  rg   r$   r#   r   r   J  s    r$   r   c                  >     e Zd ZdZdZd Zd Z fdZd Zd Z	 xZ
S )r   Nc           	        ddl m} t        d |j                         D              }t	        j
                  |j                               }t        |d         |d<   |d   }t        |d   |d   |d         |d<    |d	t        t        |j                                           } |di || _        t        | j                  j                        }	|	j                  | j                        | _        || _        || _        | j                  j&                  | _        |j                         D 
cg c]"  \  }
}|
j)                  d
      rt+        |      $ }}
}|	j,                  }t/        |D ci c]B  }|j0                  dd  |j0                  dd  |k(  r|j3                         n|j                         D c}      | _        | j4                  |   | _        d | _        d | _        y c c}}
w c c}w )Nr   )
namedtuplec              3  \   K   | ]$  \  }}|j                  d       st        |       & yw)r   N)r   r   )rG   cps      r#   rH   z*CompiledKernel.__init__.<locals>.<genexpr>`  s$     `$!QAJJW^L_d1g`s   ,,cluster_dimsr   r   arch	warp_sizeKernelMetadatar   r   rg   )collectionsr  nextrL   r   r   rm   tupler   rK   r   r>   r   r   r   pack_metadatapacked_metadatar)   rN   r5   r   r   
binary_extr   rl   r   asmkernelrv   function)r@   r)   r   rN   r  r   r   r   r  r   r	  r
  	asm_filesr  files                  r#   rC   zCompiledKernel.__init__^  s   *`.2F2F2H`a::m5578#(.)A#B (#&vi'8&.&Q\J]^#$4fT(--/=R6ST&22t}}334&44T]]C	MM&&	)7)=)=)?[AqzzRYGZT!W[	[''
!
 KKO$++ab/Z2OT__.UYUcUcUee
  hhz*  \
s   *GG)AG%c                r   | j                   y t        j                  j                         }t        j                  j	                  | j
                  | j                        | _        t        j                  j                  j                  |      d   }| j                  j                  |kD  r!t        | j                  j                  |d      t        j                  j                  j                  | j                  | j                  | j                  j                  |      \  | _         | _        | _        | _        y )Nmax_shared_memzshared memory)rv   r   r   get_current_devicelauncher_clsr)   r   runutilsget_device_propertiessharedr   load_binaryr5   r  r  n_regsn_spills)r@   device
max_shareds      r#   _init_handleszCompiledKernel._init_handles{  s    ;;"113==--dhhF]]((>>vFGWX
==*, !5!5z?SSAGATATA`A`IIt{{DMM$8$8&BB>T]DKr$   c                L    |dk(  r| j                          t        | 	  |      S )Nr  )r'  super__getattribute__)r@   r5   	__class__s     r#   r*  zCompiledKernel.__getattribute__  s&    5= w'--r$   c                b   t         j                  y t        | j                  | j                  |d      }t        | j                  t              r | j                  j                  j                  |S i }d}t        | j                  j                  j                        D ]Q  \  }}|| j                  j                  j                  v r| j                  j                  |   ||<   E||   ||<   |dz  }S |j                  | j                  j                  j                  || j                  |f       |S )N)r5   r  streamr   r   )r   launch_enter_hookr   r5   r  r9   r)   r/   r2   launch_metadatar;   	arg_names
constexprsr7   r   r   )	r@   gridr-  r   retarg_dictarg_idxiarg_names	            r#   r/  zCompiledKernel.launch_metadata  s    ++3		t}}PVWX$((I.$((++2M2M2UJ$TXX[[%:%:; 	KAxDHHKK***%)XX%7%7%A"%)']"1	 	++dDMM8-LM
r$   c                <      j                          d d fd
}|S )N)r-  c                T   | =t         j                  j                         }t         j                  j                  |      }  j                  | g| } j
                  d   d   d   | j                  j                  |t        j                  t        j                  g	|  y )Nr   r   r   )r   r   r  get_current_streamr/  r  r  r  r   r.  launch_exit_hook)r-  r   r%  r/  r2  r@   s       r#   runnerz*CompiledKernel.__getitem__.<locals>.runner  s    ~99;99&A2d224G$GODHHT!Wd1gtAwtG[G[]l#55~7V7V_Y]_r$   )r'  )r@   r2  r<  s   `` r#   __getitem__zCompiledKernel.__getitem__  s    !% 	_ r$   )r4   re   rf   r.  r;  rC   r'  r*  r/  r=  __classcell__)r+  s   @r#   r   r   W  s+     :B.
"r$   r   )r)   r:   )r   BaseExceptionr`   )2
__future__r   rO   r   _C.libtritonr   r   r   backends.compilerr   r	   r   r
   runtime.autotunerr   runtime.cacher   r   r   runtime.driverr   tools.disasmr   code_generatorr   pathlibr   r   	functoolsr~   mlir_prototype_patternptx_prototype_patternrn   mlir_arg_type_patternptx_arg_type_patternrp   r   r-   r/   ri   	lru_cacher   r   r   r   r   r   r]   r   r   rg   r$   r#   <module>rO     s    "   >  :  . U U # # '  	  	 ] X "#   I + !" 
% %P 6 1 1D,"$JV5r) ) 
d 
S Sr$   