
    rhT                       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	 ddlm
Z
mZ ddlmZ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Zd dlZd dlZd dlZd dlZdZ de iZ!dZ"de"iZ#d Z$ G d d      Z% G d d      Z& ejN                         d        Z( ejN                         d        Z)d Z*d$dZ+ G d d      Z,d%dZ-d&dZ. G d d      Z/ G d  d!e0      Z1 G d" d#      Z2y)'    )annotationsN   )get_cache_invalidating_env_varsir)backends)Language)BaseBackend	GPUTarget)__version__knobs)OutOfResources)get_cache_managerget_dump_managerget_override_manager)driver)get_sass)Pathz=\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\)ptxz\.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 {[^}]+} *   )researchsubconvert_type_reprgroup)xmatchtmas      k/var/www/html/ai-insurance-compliance-backend/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                  (    e Zd ZdddZd Zd Zd Zy)	ASTSourceNc                   || _         t        j                  | _        d| _        |j
                  | _        || _        t               | _	        |g|j                         D ]T  \  }}t        |t              r|j                  j                  |      fn|}t        |t              sJ || j                  |<   V |xs
 t               | _        t        | j                  t              rLt#        | j                  j%                  d            D ci c]  \  }}||j'                          c}}| _        y | j                  j)                         D ]  }t        |t              rt+        d       y c c}}w )Nttir,zSignature keys must be string)fnr   TRITONlanguageext__name__name	signaturedict	constantsitems
isinstancestr	arg_namesindextupleattrs	enumeratesplitstripkeys	TypeError)selfr)   r/   
constexprsr8   kvs          r"   __init__zASTSource.__init__6   s$    KK	"!"((* &11;As1CR\\''*-!!U+++$%q!& _df
dnnc*7@AUAUVYAZ7[\tq!al\DN^^((* E!!S)#$CDDE ]s   E(c           	        t        | j                  j                               D cg c]  \  }}|	 }}}d }dj                  t        | j                  j                               D cg c]  \  }} ||       c}}      }| j
                  j                   dt        | j                         d| d| }t        j                  |j                  d            j                         S c c}}w c c}}w )Nc                H    t        | d      r| j                  S t        |       S )N	cache_key)hasattrrE   r4   )r   s    r"   <lambda>z ASTSource.hash.<locals>.<lambda>L   s    71k+BAKK A r#   -utf-8)sortedr/   r2   joinr1   r)   rE   r4   r8   hashlibsha256encode	hexdigest)r>   r@   rA   
sorted_sigget_keyconstants_keykeys          r"   hashzASTSource.hashJ   s    $*4>>+?+?+A$BCDAqaC
CN@T@T@V9W!XA'!*!XY""#1S_$5Qzl!M?S~~cjj12<<>>	 D!Xs   C**C0
c                >    ddl m}  || j                  | ||||      S )Nr   )ast_to_ttir)contextoptionscodegen_fns
module_map)code_generatorrV   r)   )r>   rX   rY   rZ   rW   rV   s         r"   make_irzASTSource.make_irQ   s$    /477D'7Xc&02 	2r#   c                    t               S N)r0   r>   s    r"   parse_optionszASTSource.parse_optionsV   s	    vr#   NNreturnNoner-   
__module____qualname__rB   rT   r\   r`    r#   r"   r%   r%   4   s    E(?2
r#   r%   c                  $    e Zd Zd Zd Zd Zd Zy)IRSourcec                   || _         t        |      }|j                  dd  | _        t        j
                  | _        |j                         | _        t        j                  |       |j                  |       | j                  dk(  rt        j                  t        | j                     | j                  t        j                        }|j                  d      | _        |j                  d      }t        j"                  t$        | j                     |      }t'        |      D ci c]  \  }}|t)        |       c}}| _        y t        j,                  | j                   |      | _        | j.                  j1                         }	d|	z   | _        | j.                  j3                  |	      }
| j.                  j5                  |
      }t'        |      D ci c]  \  }}||
 c}}| _        y c c}}w c c}}w )Nr   r   r   @)pathr   suffixr,   r   r*   r+   	read_textsrcr   load_dialectsr   r   prototype_pattern	MULTILINEr   r.   findallarg_type_patternr9   r   r/   parse_mlir_modulemoduleget_entry_func_nameget_functionget_function_signature)r>   rm   rW   backendr    r/   typesr@   tyfn_namefuncOpfunc_tys               r"   rB   zIRSource.__init__\   sn   	Dz;;qr? >>#
!g& 88uII/9488R\\READIAIJJ/99EEDMeDTU51ba!22!66UDN..tyy'BDKkk557GgDI[[--g6Fkk88@G1:71CD2aeDDN V Es   "G&G,c                z    t        j                  | j                  j                  d            j	                         S )NrI   )rL   rM   rp   rN   rO   r_   s    r"   rT   zIRSource.hashu   s'    ~~dhhoog67AACCr#   c                <    || j                   _        | j                   S r^   )rw   rW   )r>   rX   rY   rZ   rW   s        r"   r\   zIRSource.make_irx   s    %{{r#   c                    | j                   dk(  r(| j                  j                  d      }|J d       d|iS t               S )Nttgirzttg.num-warpsz'Unable to parse ttg.num-warps attribute	num_warps)r,   rw   get_int_attrr0   )r>   r   s     r"   r`   zIRSource.parse_options|   sF    88w00AI(S*SS(++vr#   Nre   rh   r#   r"   rj   rj   Z   s    E2Dr#   rj   c                    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#        j$                  d      j'                  d	      d
   }	t        t        j                  j                  |d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   jx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.)prefix
EXT_SUFFIX._Cz
libtriton.i   r+   ztriton.language.rH   )pkgutilosrm   dirnameabspath__file__openrL   rM   readrO   rK   walk_packagesmodule_finder	find_specr.   origin	sysconfigget_config_varr:   updateappendr   )r   TRITON_PATHcontentsfpath_prefixesrm   r   liblibtriton_hashr,   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

"
"<
0
6
6s
;B
?C	bggll;
3%.@A4	H )AFF7OE!!%(	  	) OON,,./GGLLj9M$$m_=O$P ?###--chh7>>E 	?1;;=>>H	? 	?? ]chhx0007; ;C C) )	? 	?s0   46L6L9(L,96L9LL),L69M	c                Z    t         j                  j                  j                  |       d   S )Nmax_shared_mem)r   activeutilsget_device_properties)devices    r"   r   r      s#    ==44V<=MNNr#   c                    |dk(  s|dk(  rt        j                  | |      }||_        |S |dk(  s
|dk(  s|dk(  rt        |       j	                         S |dk(  s|dk(  rt        |       j                         S y )Nr'   r   llirr   amdgcncubinhsaco)r   rv   rW   r   ro   
read_bytes)	full_namer,   rW   rw   s       r"   parser      sx    
f}w%%i9 
f}uxI((**
g~I))++ (r#   c                ,   t         j                  j                  ry| j                  t	        | j                         | j
                  t	        | j
                         ddg}|D cg c]"  }|j                  dt        j                        $ }}| j                  g }4t        fd|D              s|j                         j                  4t        ||dd       D ]  \  }}||_         |sd| _	        yd|d   _        |d   | _	        yc c}w )	z
    Removes code_generator.py and related files from tracebacks.

    These are uninteresting to the user -- "just show me *my* code!"
    Nz"/triton/compiler/code_generator.pyz/ast.py/c              3     K   | ]6  }j                   j                  j                  j                  |      s3| 8 y wr^   )tb_framef_codeco_filenameendswith).0r   tbs     r"   	<genexpr>z#filter_traceback.<locals>.<genexpr>   s/     V2;;+=+=+I+I+R+RST+U1Vs   4??r   r   r   )r   compilationfront_end_debugging	__cause__filter_traceback__context__replacer   sep__traceback__anyr   tb_nextzip)e	BAD_FILESbad_fileframes	cur_frame
next_framer   s         @r"   r   r      s    ,,{{%}} ' 	-I @II8!!#rvv.III	
BF
.ViVVMM"ZZ .
 $'vvabz#: 'J&	' !r
 )! Js   ''Dc                  ,    e Zd ZddZddZddZddZy)	CompileTimerc                `    t        j                          | _        d | _        g | _        d | _        y r^   )timestartir_initialization_endlowering_stage_endsstore_results_endr_   s    r"   rB   zCompileTimer.__init__   s'     IIK
37"<> /3r#   c                6    t        j                          | _        y r^   )r   r   r_   s    r"   finished_ir_initializationz'CompileTimer.finished_ir_initialization   s    %)YY["r#   c                b    | j                   j                  |t        j                         f       y r^   )r   r   r   )r>   
stage_names     r"   stage_finishedzCompileTimer.stage_finished   s!      ''TYY[(ABr#   c                l   t        j                          }| j                  || _        n|| _        dd}g }| j                  }| j                  D ]!  \  }}|j	                  | |||      f       |}# t        j                   || j                  | j                        | ||| j                              S )Nc                *    |yt        || z
  dz        S )Nr   i@B )int)r   ends     r"   deltazCompileTimer.end.<locals>.delta   s    {ew.//r#   )ir_initializationlowering_stagesstore_results)r   floatr   zfloat | Nonerc   r   )r   r   r   r   r   r   CompileTimesr   )r>   	timestampr   lowering_stage_durationsstage_startr   	stage_ends          r"   r   zCompileTimer.end   s    IIK	%%-)2D&%.D"	0
 $& 00%)%=%= 	$!J	$++Z{I9V,WX#K	$ !!#DJJ0J0JK4T-C-CD
 	
r#   Nrb   )r   r4   rc   rd   )rc   zknobs.CompileTimes)r-   rf   rg   rB   r   r   r   rh   r#   r"   r   r      s    41C
r#   r   c                B   t         j                  j                  }|r
t               }|t        j
                  j                         }t        |t              sJ d       t        |      }t        | t               }|r8t        | t              sJ d       t        j                         }t        | ||      } | j                         }|j                  t!        |xs
 t!               fi |      }t#               }	t%                d| j'                          d|j'                          d|j'                          dt        t)        |	j+                                      	}
t-        j.                  |
j1                  d            j3                         }t5        |      }t         j                  j6                  }t         j                  j8                  }t         j                  j:                  }|rt=        | j'                               nd }|rt?        | j'                               nd }| j@                  d d }| d}|jC                  |      xs i }|jE                  |      }t         j                  jF                  }|sF|DtI        | ||      }|r3 || |jJ                  jM                         |jO                         d       |S ||d	|jP                  |	}tR        |d
<   t!               }|jU                  ||| jV                         tY        |j[                               j]                  | j^                        }|r|dz  }t        | t              s:t        j                         }t        j`                  |       |ja                  |       |jc                  |      }|je                         }	 | jg                  |||      }|r'| d| j^                   }|jm                  ||      ||<   n| d}|jm                  ||      ||<   t         j                  jn                  } |r5| r3|jq                  | jr                         tu        d| jr                          |rjw                          tY        |j+                               |d  D ]  \  }!}" |"||      }#| d|! }|6|jE                  dd       x}$rP|$jy                  d|!       r<t{        |$|!|      }#n.|j}                  |      x}%rtu        d|%        t{        |%|!|      }#|r|!dv r|jm                  |#|      ||<   ||jm                  |#|       | |!k(  r0|j}                  |      }&|#jq                  |&       tu        d|&        |#}|sj                  |!        |jm                  t        j                  |t              |d      ||<   |j                  ||       t         j                  j                  s|j                          |r || ||jO                         d       tI        | ||      S # th        $ r}tk        |        d }~ww xY w)Nz target must be of GPUTarget typez'source must be either AST or a filepathrH   rI      .jsonT)rp   metadatametadata_grouptimes	cache_hit)rT   targettriton_versionr   r   z.sourcezCreating new locations for ir_overridez
Overriding kernel with file )r   r   json)defaultF)binary)Fr   r   listenerr   r   r   get_current_targetr3   r
   make_backendr%   r4   r   rW   rj   r`   r0   r   r   rT   rJ   r2   rL   rM   rN   rO   r   overridedump_irstore_binary_onlyr   r   r.   	get_groupgetalways_compileCompiledKernelr   _asdictr   __dict__r   
add_stagesr+   listr<   r6   r,   rq   get_codegen_implementationget_module_mapr\   	Exceptionr   put
use_ir_loccreate_location_snapshotrm   printr   r   r   get_filer   r   dumpsvars	put_groupenable_asandisable_multithreading)'rp   r   rX   compilation_listenertimerr{   	ir_sourcerW   extra_optionsenv_varsrS   rT   fn_cache_managerenable_overrideenable_ir_dumpstore_only_binaryfn_override_managerfn_dump_manager	file_namemetadata_filenamer   metadata_pathr   resr   stagesfirst_stagerY   rZ   rw   r   ir_filenamer
  r,   
compile_irnext_moduler   r   ir_full_names'                                          r"   compiler(    s    ,,55~113fi(L*LL(6"GsI..I#s#N%NN#**,sGW-%%'M##D):DF$Lm$LMG.0H\N!CHHJ<q(8',,.9I3vV^VdVdVfOgKhJi
jC>>#**W-.88:D(. ''00O&&..N));;>M.sxxz:SW6D&sxxz2$O
 #I$+U+%//0ABHbN"&&'89M&&55Nm7S.$7 --/-iik 
  

 	H "-HVFvw5v{{}%++CGG4Kq c8$**,
!g&44W=K'')JWk:wG
 "1SWWI.&6&:&:6;&O{#"7+&6&:&:6;&O{#""--JZ''1+CHH:67((*/= &Z 2"1SE*&  (||M4@@@kFZFZ]^_b^c[dFe#Kg>-66{CCYC29+>?	38K!s.H'H*:*>*>{K*XN;'&[9+44[AL00>/~>?  %/&2 )9(<(<TZZZ^=_arDI )= )KN$%0.A ((&&( x^c^g^g^i',	. #~t44  s    X 	XXX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   valuesr   supports_targetlenRuntimeErrorr{   )r   r   activess      r"   r   r     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 Zd Zy)LazyDictc                     || _         g | _        y r^   )dataextras)r>   r2  s     r"   rB   zLazyDict.__init__  s    	r#   c                    | j                   D ]  \  }}| j                   || z  | _         | j                   j                          | j                  S r^   )r3  r2  clearr>   funcargss      r"   r   zLazyDict.get  sG    ++ 	0JD$		D$K/DI	0yyr#   c                >    | j                   j                  ||f       y r^   )r3  r   r6  s      r"   addzLazyDict.add  s    D$<(r#   N)r-   rf   rg   rB   r   r:  rh   r#   r"   r0  r0    s    )r#   r0  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>   rS   values      r"   __missing__zAsmDict.__missing__  s6    &=T']+E.455S	r#   N)r-   rf   rg   rA  rh   r#   r"   r<  r<    s    r#   r<  c                  6     e Zd Zd Zd Z fdZd Zd Z xZS )r  c           	        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   )r   cps      r"   r   z*CompiledKernel.__init__.<locals>.<genexpr>  s$     `$!QAJJW^L_d1g`s   ,,cluster_dimsr   r{   arch	warp_sizeKernelMetadatar   r   rh   )collectionsrD  nextr2   r   loadsro   r7   r
   rJ   r  r<   r   r   r   pack_metadatapacked_metadatarp   rT   r.   r   r   
binary_extr<  rn   r   asmkernelrw   function)r>   rp   r   rT   rD  r   r   r   rK  r{   rF  rG  	asm_filesrQ  files                  r"   rB   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                   | j                   y t        j                  j                         }t        j                  j	                  | j
                  | j                        | _        t        |      }| j                  j                  |kD  r!t        | j                  j                  |d      t        | j                  d      rR| j                  j                  <d}| j                  j                  |kD  r!t        | j                  j                  |d      t        j                  j                  j                  | j                  | j                   | j                  j                  |      \  | _         | _        | _        | _        | _        t        j                  j+                         j,                  }| j                  j.                  |z  | j(                  kD  r.t        | j                  j.                  |z  | j(                  d      y )Nzshared memory	tmem_sizei   ztensor memorythreads)rw   r   r   get_current_devicelauncher_clsrp   r   runr   sharedr   rF   rX  r   load_binaryr.   rS  rT  n_regsn_spillsn_max_threadsr   rJ  r   )r>   r   
max_sharedmax_tmem_sizerJ  s        r"   _init_handleszCompiledKernel._init_handles  sr   ;;"113==--dhhF#F+
==*, !5!5z?SS4==+.4==3J3J3VM}}&&6$T]]%<%<m_]]U[UbUbUhUhUtUtIIt{{DMM$8$8&VBRT]DK@RMM446@@	==""Y.1C1CC !8!89!DdFXFXZcdd Dr#   c                L    |dk(  r| j                          t        | 	  |      S )Nr\  )rd  super__getattribute__)r>   r.   	__class__s     r"   rg  zCompiledKernel.__getattribute__  s&    5= w'--r#   c                   t         j                  j                  y t        | j                  | j
                  |d      }t        | j                  t              r | j                  j                  j                  |S i }d}t        | j                  j                  j                        D ]  \  }}||   ||<   |dz  } |j                  | j                  j                  j                  || j                  |f       |S )N)r.   rT  streamr   r   )r   runtimelaunch_enter_hookr0  r.   rT  r3   rp   r%   r)   launch_metadatar9   r5   r:  r   )	r>   gridrj  r8  retarg_dictarg_idxiarg_names	            r"   rm  zCompiledKernel.launch_metadata  s    ==**2		t}}PVWX$((I.$((++2M2M2UJ$TXX[[%:%:; 	KAx!%gHXqLG	 	++dDMM8-LM
r#   c                <      j                          d d fd
}|S )N)rj  c                |   | =t         j                  j                         }t         j                  j                  |      }  j                  | g| } j
                  d   d   d   | j                  j                  |t        j                  j                  t        j                  j                  g	|  y )Nr   r   r   )r   r   rZ  get_current_streamrm  r\  rT  rP  r   rk  rl  launch_exit_hook)rj  r8  r   rm  rn  r>   s       r"   runnerz*CompiledKernel.__getitem__.<locals>.runner  s    ~99;99&A2d224G$GODHHT!Wd1gtAwtG[G[]l]]44emm6T6T]W[]r#   )rd  )r>   rn  rx  s   `` r"   __getitem__zCompiledKernel.__getitem__  s    !% 	] r#   )	r-   rf   rg   rB   rd  rg  rm  ry  __classcell__)rh  s   @r"   r  r    s    :e,.
r#   r  )r   BaseExceptionra   )r   r
   rc   r	   )3
__future__r   rL   r   _C.libtritonr   r   r   backends.compilerr   r	   r
   r   r   r   runtime.autotunerr   runtime.cacher   r   r   runtime.driverr   tools.disasmr   pathlibr   r   	functoolsr   r   r   ptx_prototype_patternrr   ptx_arg_type_patternru   r   r%   rj   	lru_cacher   r   r   r   r   r(  r   r0  r0   r<  r  rh   r#   r"   <module>r     s   "   >  ( 6 ! . U U # #  	  	   Y 	   , 	 

# #L' 'T  1  1F O O,#$L$
 $
NN5b) ) 
d 
S Sr#   