
    wiID                       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j(                    d             Z)d Z*d*d Z+d+d!Z,d" Z- G d# d$          Z. G d% d&e/          Z0 G d' d(          Z1dS ),    )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|           }|dS t          j        dd|           } |%dt          |                    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      h/root/.openclaw/workspace/chatterbox_venv_py311/lib/python3.11/site-packages/triton/compiler/compiler.pyr   r   ,   sn     I)1--E
))1
-
-C
{
{B""A&u{{1~~6666H    srcstrc                    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)r%   ttgir_num_warps_patternnum_warps_matches	num_warpss       r#   _get_num_warps_from_ir_strr.   9   sY    G 
#:C@@ !!Q&&&(R&&&%a())Ir$   c                  *    e Zd Zdd	dZd Zd Zd ZdS )
	ASTSourceNreturnNonec                \   || _         d| _        |j        | _        || _        || _        || _        t          | j        t                    r7d t          | j        
                    d                    D             | _        n@| j                                        D ]&}t          |t                    st          d          '| j        i | _        n@| j                                        D ]&}t          |t                    st          d          '| j        t                      | _        d S d S )Nr   c                >    i | ]\  }}||                                 S  )strip.0kvs      r#   
<dictcomp>z&ASTSource.__init__.<locals>.<dictcomp>M   s&    \\\tq!a\\\r$   ,zSignature keys must be stringzConstants keys must be string)fnext__name__name	signature	constantsattrs
isinstancer&   	enumeratesplitkeys	TypeErrorr	   )selfr=   rA   rB   rC   r9   s         r#   __init__zASTSource.__init__E   s3   K	""
dnc** 	E\\yAUAUVYAZAZ7[7[\\\DNN^((** E E!!S)) E#$CDDDE>!DNN^((** E E!!S)) E#$CDDDE:(**DJJJ r$   c                   d t          | j                                                  D             }t          d | j                                        D                       }| j        j         d| j                                         d| d| }t          j	        |
                    d                                                    S )Nc                    g | ]\  }}|S r5   r5   r7   s      r#   
<listcomp>z"ASTSource.hash.<locals>.<listcomp>\   s    CCCDAqaCCCr$   c              3  >   K   | ]\  }}t          |          |fV  d S N)r&   r7   s      r#   	<genexpr>z!ASTSource.hash.<locals>.<genexpr>_   s0      !Q!Q$!Q3q661+!Q!Q!Q!Q!Q!Qr$   -utf-8)sortedrA   itemsrB   r=   	cache_keyrC   hashhashlibsha256encode	hexdigest)rI   
sorted_sigsorted_constantskeys       r#   rV   zASTSource.hash[   s    CCF4>+?+?+A+A$B$BCCC
 "!Q!Q$.:N:N:P:P!Q!Q!QQQ"XXTZ__%6%6XXXXFVXX~cjj1122<<>>>r$   c                6    t          | j        | ||||          S )N)contextoptionscodegen_fns
module_map)r   r=   )rI   r`   ra   rb   r_   s        r#   make_irzASTSource.make_irc   s)    47D'7Xc&02 2 2 	2r$   c                    t                      S rO   )dictrI   s    r#   parse_optionszASTSource.parse_optionsg   s    vvr$   NNr1   r2   r?   
__module____qualname__rJ   rV   rc   rg   r5   r$   r#   r0   r0   C   sZ        + + + + +,? ? ?2 2 2    r$   r0   c                  &    e Zd Zd Zd Zd Zd ZdS )IRSourcec                   || _         t          |          }|j        dd          | _        |                                | _        t          j        t          | j                 | j        t          j	                  }|
                    d          | _        |
                    d          }t          j        t          | j                 |          }d t          |          D             | _        d S )Nr   r   c                4    i | ]\  }}|t          |          S r5   )r   )r8   r9   tys      r#   r;   z%IRSource.__init__.<locals>.<dictcomp>v   s'    QQQuq"!.r22QQQr$   )pathr   suffixr>   	read_textr%   r   r   prototype_pattern	MULTILINEr   r@   r(   arg_type_patternrE   rA   )rI   rr   r!   rA   typess        r#   rJ   zIRSource.__init__m   s    	Dzz;qrr?>>##	+DH5txNNKKNN	KKNN	
+DH5yAAQQ	%@P@PQQQr$   c                ~    t          j        | j                            d                                                    S )NrR   )rW   rX   r%   rY   rZ   rf   s    r#   rV   zIRSource.hashx   s,    ~dhoog6677AACCCr$   c                H    t          j        | j        |          }||_        |S rO   )r   parse_mlir_modulerr   r_   )rI   r`   ra   rb   r_   modules         r#   rc   zIRSource.make_ir{   s"    %di99 r$   c                `    | j         dk    rdt          | j                  iS t                      S )Nr   r-   )r>   r.   r%   re   rf   s    r#   rg   zIRSource.parse_options   s.    8w!;DH!E!EFFvvr$   Nrj   r5   r$   r#   rn   rn   k   sV        	R 	R 	RD D D  
    r$   rn   c                    dd l } t          j                            t          j                            t          j                            t
                                        }g }t          t
          d          5 }|t          j        |	                                          
                                gz  }d d d            n# 1 swxY w Y   t          j                            |d          dft          j                            |d          dfg}|D ]\  }}|                     |g|          D ]}t          |j                            |j                  j        d          5 }|t          j        |	                                          
                                gz  }d d d            n# 1 swxY w Y   t          j                    }t          t          j                            |d          d          5 }	 |	                    d
          }	|	sn|                    |	           .	 d d d            n# 1 swxY w Y   |                    |
                                           t          j                            |d          }
|                     |
gd          D ]}t          |j                            |j                  j        d          5 }|t          j        |	                                          
                                gz  }d d d            n# 1 swxY w Y   t&           d                    |          z   S )Nr   rbcompilerztriton.compiler.r   ztriton.backends.)prefixz_C/libtriton.soTi   languageztriton.language.rQ   )pkgutilosrr   dirnameabspath__file__openrW   rX   readrZ   joinwalk_packagesmodule_finder	find_specr@   originupdateappendr
   )r   TRITON_PATHcontentsfpath_prefixesrr   r   liblibtriton_hashchunklanguage_paths              r#   
triton_keyr      s   NNN'//"'//"'//(2K2K"L"LMMKH	h		 ;W^AFFHH--7799::; ; ; ; ; ; ; ; ; ; ; ; ; ; ; 
k:	.	.0BC	k:	.	.0BCM & C Cf(($(?? 	C 	CCc'11#(;;BDII CQW^AFFHH55??AABBC C C C C C C C C C C C C C C	C
 ^%%N	bgll;(9::D	A	A )Q	)FF7OOE !!%(((		) 	) ) ) ) ) ) ) ) ) ) ) ) ) ) ) OON,,..///GLLj99M$$m_=O$PP ? ?##--ch77>EE 	?11;;==>>H	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	?chhx0000sH   :=CC
C%=F..F25F2<0H99H= H==L  L$	'L$	c                   |dk    s|dk    rt          j        | |          }||_        |S |dk    s|dk    r!t          |                                           S |dk    r!t          |                                           S d S )Nr   r   llirr   cubin)r   r{   r_   r   rt   
read_bytes)	full_namer>   r_   r|   s       r#   parser      s    
f}}w%i99 
f}}uI((***
g~~I))+++ ~r$   eBaseExceptionc                   t          j        dd          dk    rdS | j        t          | j                   | j        t          | j                   ddg}| j        g }9t          fd|D                       s|                               j        9t          ||dd                   D ]\  }}||_        |s	d| _        dS d|d	         _        |d
         | _        dS )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  d   K   | ]*}j         j        j                            |          &|V  +d S rO   )tb_framef_codeco_filenameendswith)r8   r   tbs     r#   rP   z#filter_traceback.<locals>.<genexpr>   s?      VV2;+=+I+R+RST+U+UV1VVVVVVr$   r   r   )
r   getenv	__cause__filter_traceback__context____traceback__anyr   tb_nextzip)r   	BAD_FILESframes	cur_frame
next_framer   s        @r#   r   r      s    
y-s33s::{%%%} ''' 	-I
 
BF
.VVVViVVVVV 	MM"Z .
 $'vvabbz#:#: ' 'J&	 $!r
 )r$   c                
   |t           j                                        }t          |t                    s
J d            t          |          }t          | t                     }|r.t          | t                    s
J d            t          |           } | 	                                }|	                    t          |pt                      fi |          }t                      }t                       d|                                  d|                                 d|                                 dt          t          |                                                     	}t!          j        |                    d                                                    }t)          |          }	t*          j                            dd          dk    }
t*          j                            dd          dk    }|
r!t1          |                                           nd }|r!t3          |                                           nd }| j        d d	         }| d
}|	                    |          pi }|                    |          }t*          j                            dd          dk    }|sF|Dt9          j        t=          |                                                    }tA          | ||          S ||d|j!        |}t                      }|"                    ||           tG          |$                                          %                    | j&                  }|r|dz  }tO          j(                    }tO          j)        |           |)                    |           |*                                }|+                                }	 | ,                    ||||          }n"# tZ          $ r}t]          |            d }~ww xY wt*          j                            dd           }tG          |                                          |d          D ]\  }} |||          }| d| }|:|/                    |          x} #ta          d|             tc          | ||          }|	2                    ||          ||<   ||2                    ||           ||k    r<|	/                    |          }!|3                    |!           ta          d|!            |}|	2                    t9          j4        |tj                    |d          ||<   |	6                    ||           |7                                 tA          | ||          S )Nz target must be of GPUTarget typez'source must be either AST or a filepathrQ   rR   TRITON_KERNEL_OVERRIDEr   r   TRITON_KERNEL_DUMP   .jsonTRITON_ALWAYS_COMPILE)rV   targetr   
USE_IR_LOC.z
Overriding kernel with file zCreating new locations for )defaultF)binary)8r   activeget_current_targetrD   r   make_backendr0   r&   rn   rg   re   r   r   rV   rS   rT   rW   rX   rY   rZ   r   r   environgetr   r   r@   	get_groupjsonloadsr   rt   CompiledKernel__dict__
add_stageslistrG   indexr>   r   r_   load_dialectsget_codegen_implementationget_module_maprc   	Exceptionr   get_fileprintr   putcreate_location_snapshotdumpsvars	put_groupdisable_multithreading)"r%   r   r`   backend	ir_sourceextra_optionsenv_varsr]   rV   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_stager_   ra   rb   r|   r   
use_ir_locr>   
compile_irnext_moduleir_filenamer   ir_full_names"                                     r#   compiler      s^   ~1133fi((LL*LLL(6""GsI...I #s##NN%NNN#smm%%''M##D):DFF$L$Lm$L$LMMG.00H\\
j
jCHHJJ
j
j
j
j',,..
j
j3vV^VdVdVfVfOgOgKhKh
j
jC>#**W--..88::D(.. jnn%=sCCsJOZ^^$8#>>#EN>MW.sxxzz:::SW6DN&sxxzz222$O
 #I$+++%//0ABBHbN"&&'899MZ^^$;SAASHN 9m7:d=11;;==>>c>4888   
 	H VVFvw'''v{{}}%%++CG44K qjllGW'"""4466K''))JWk:wGG    d33J//=  Z j22"**S**+>Q>Z>Z[f>g>g1g0t>9>>???	388K&6&:&:;&T&T{#&[999+44[AAL00>>>>>>???(8(<(<TZZ^=_=_=_arDI )= )K )KN$%0.AAA
 ""$$$#~t444s    O 
O8#O33O8c                      fdt          j                    D             }t          |          dk    r*t          t          |           d j         d| d           |d                    S )Nc                R    g | ]#}|j                                       |j         $S r5   )r   supports_target)r8   r    r   s     r#   rM   z make_backend.<locals>.<listcomp>3  s1    [[[a
8R8RSY8Z8Z[qz[[[r$   r   z! compatible backends for target (z) (z). There should only be one.r   )r   valuesr)   RuntimeErrorr   )r   activess   ` r#   r   r   2  s    [[[[8?#4#4[[[G
7||q7||vvfnvvQXvvvx x 	x71:fr$   c                  "    e Zd Zd ZddZd ZdS )LazyDictc                "    || _         g | _        d S rO   )dataextras)rI   r   s     r#   rJ   zLazyDict.__init__<  s    	r$   r1   r2   c                    | j         D ]\  }}| j         || z  | _        | j                                          | j        S rO   )r   r   clearrI   funcargss      r#   r   zLazyDict.get@  sI    + 	0 	0JD$	DD$K/DIIyr$   c                >    | j                             ||f           d S rO   )r   r   r  s      r#   addzLazyDict.addF  s"    D$<(((((r$   Nri   )r?   rk   rl   rJ   r   r  r5   r$   r#   r   r   :  sF             ) ) ) ) )r$   r   c                      e Zd Zd ZdS )AsmDictc                l    |dk    rt          | d                   }nt          d|z            || |<   |S )Nsassr   zUnknown key: '%s')r   KeyError)rI   r]   values      r#   __missing__zAsmDict.__missing__L  s@    &==T']++EE.4555S	r$   N)r?   rk   rl   r  r5   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dl m} t          d |                                D                       }t	          j        |                                          }t          |d                   |d<   |d         }t          |d         |d         |d                   |d<    |d	t          t          |                                                              } |di || _        t          | j        j                  }	|	                    | j                  | _        || _        || _        | j        j        | _        d
 |                                D             }
|	j        t+          fd|
D                       | _        | j                 | _        d | _        d | _        d S )Nr   )
namedtuplec              3  d   K   | ]+\  }}|                     d           t          |          V  ,dS )r   Nr   r   r8   cps      r#   rP   z*CompiledKernel.__init__.<locals>.<genexpr>`  s>      ``$!QAJJW^L_L_`d1gg``````r$   cluster_dimsr   r   arch	warp_sizeKernelMetadatac                \    g | ])\  }}|                     d           t          |          *S )r   r  r  s      r#   rM   z+CompiledKernel.__init__.<locals>.<listcomp>n  s4    [[[AqzzRYGZGZ[T!WW[[[r$   c                    i | ]K}|j         d d         |j         d d         k    r|                                n|                                LS )r   N)rs   r   rt   )r8   file
binary_exts     r#   r;   z+CompiledKernel.__init__.<locals>.<dictcomp>p  se     
 
 
 KO$+abb/Z2O2OT__...UYUcUcUeUe
 
 
r$   r5   )collectionsr  nextrT   r   r   rt   tupler   rS   r   rG   r   r   r   pack_metadatapacked_metadatar%   rV   r@   r  r  asmkernelr|   function)rI   r%   r   rV   r  r   r   r   r  r   	asm_filesr  s              @r#   rJ   zCompiledKernel.__init__^  s   ******``.2F2F2H2H```aa:m557788#(.)A#B#B (#&vi'8&.&Q\J]^^#$4fT(--//=R=R6S6STT&2222t}344&44T]CC	M&	[[)=)=)?)?[[[	'
 
 
 
 
!
 
 
   hz* r$   c                   | j         d S t          j                                        }t          j                            | j        | j                  | _        t          j        j        	                    |          d         }| j        j
        |k    rt          | j        j
        |d          t          j        j                            | j        | j        | j        j
        |          \  | _         | _        | _        | _        d S )Nmax_shared_memzshared memory)r|   r   r   get_current_devicelauncher_clsr%   r   runutilsget_device_propertiessharedr   load_binaryr@   r#  r$  n_regsn_spills)rI   device
max_shareds      r#   _init_handleszCompiledKernel._init_handles{  s    ;"F1133=--dhFF](>>vFFGWX
=*,, !5z?SSSAGATA`A`It{DM$8&BB BB>T]DKr$   c                z    |dk    r|                                   t                                          |          S )Nr*  )r3  super__getattribute__)rI   r@   	__class__s     r#   r6  zCompiledKernel.__getattribute__  s6    5==   ww''---r$   c                   t           j        d S t          | j        | j        |d          }t          | j        t                    r| j        j        j	        |S i }d}t          | j        j        j                  D ]>\  }}|| j        j        j        v r| j        j        |         ||<   .||         ||<   |dz  }?|                    | j        j        j	        || j        |f           |S )N)r@   r$  streamr   r   )r   launch_enter_hookr   r@   r$  rD   r%   r0   r=   launch_metadatarE   	arg_names
constexprsrB   r  r   )	rI   gridr9  r  retarg_dictarg_idxiarg_names	            r#   r;  zCompiledKernel.launch_metadata  s    +34	t}PVWWXX$(I.. 	$(+2M2UJ$TX[%:;; 	 	KAxDHK***%)X%7%A""%)']"1+dDM8-LMMM
r$   c                D                                        d d fd
}|S )N)r9  c                0   | =t           j                                        }t           j                            |          }  j        | g|R  } j        d         d         d         | j        j        |t          j	        t          j
        g	|R   d S )Nr   r   r   )r   r   r(  get_current_streamr;  r*  r$  r!  r   r:  launch_exit_hook)r9  r  r1  r;  r>  rI   s       r#   runnerz*CompiledKernel.__getitem__.<locals>.runner  s    ~99;;99&AA2d24G$GGGODHT!Wd1gtAwtG[]l#5~7V_Y]_ _ _ _ _ _r$   )r3  )rI   r>  rH  s   `` r#   __getitem__zCompiledKernel.__getitem__  sP    !% 	_ 	_ 	_ 	_ 	_ 	_ 	_ 	_ r$   )r?   rk   rl   r:  rG  rJ   r3  r6  r;  rI  __classcell__)r7  s   @r#   r   r   W  s           :B B B. . . . .
  "      r$   r   )r%   r&   )r   r   rh   )2
__future__r   rW   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_patternru   mlir_arg_type_patternptx_arg_type_patternrw   r   r.   r0   rn   	lru_cacher   r   r   r   r   r   re   r  r   r5   r$   r#   <module>rZ     s   " " " " " "   > > > > > > > >       : : : : : : : :       . . . . . . U U U U U U U U U U # # # # # # # # # # # # ' ' ' ' ' '       				     				 ] X "#    I + !"  
 
 
   % % % % % % % %P       6 1 1 1D, , ,"$ "$ "$ "$JV5 V5 V5 V5r  ) ) ) ) ) ) ) ) 
 
 
 
 
d 
 
 
S S S S S S S S S Sr$   