K iQddlmZmZmZddlmZmZmZmZddl m Z ddl m Z ddl mZmZmZddlmZddlZddlZddlZddlZddlZddlmZd efd Zd Zd Ze d GddZGddeZy)) BaseBackend GPUTargetLanguage)irpassesllvmamd)knobs) dataclass)AnyDictTuple) ModuleTypeN)PathtargetcdS)Ncy)Nrr)lhs_typerhs_types b/mnt/ssd/data/python-lab/Trading/venv/lib/python3.12/site-packages/triton/backends/amd/compiler.pyz"get_min_dot_size..srrs rget_min_dot_sizers  0/rctjj|dk(xs |dk(xr|duStjjS)Ngfx942gfx950T)r r use_block_pingpong)archuse_async_copys ris_pingpong_schedule_enabledr$sI--5 H  M!1!Ln6LX;@99;W;WXrcttjj|dk(StjjS)Nr)r r use_in_thread_transposer"s ris_in_thread_transpose_enabledr(s.!&!B!B!JDH qPUPYPYPqPqqrT)frozenc`eZdZUdZeed<dZeed<dZeed<dZeed<dZ e ed <d Z e ed <d Z eed <dZeed<dZeed<dZeeed<dZeeed<dZeed<dZeeed<dZeed<d Zeed<dZeed<dZeed<d Zeed<dZeed<d Zeed!<d"Zeed#<d$Zeed%<d&Z d'Z!y)( HIPOptions num_warpsr waves_per_eu num_stagesnum_ctasN extern_libsr cluster_dimsFdebugTsanitize_overflowr")fp8e4nvfp8e5fp8e5b16fp8e4b8supported_fp8_dtypesr!deprecated_fp8_dot_operand_dtypesieeedefault_dot_input_precision)r<allowed_dot_input_precisionsenable_fp_fusionlaunch_cooperative_gridrmatrix_instr_nonkdimkpackallow_flush_denormmax_num_imprecise_acc_defaulthip backend_nameinstrumentation_modenone schedule_hintct|jdd}|dk\rdnd}tj|d||jdkDr|j|jdz zdk(sJd |jd k(rI|j dk7r:t jd |j d tj|d dttjdz }|jint|j}dD]}t||dz ||<tj|dt|jy)N @ warp_sizerrznum_warps must be a power of 2r zckpack is deprecated starting from gfx950 and will be removed in later releases. So for now kpack = z7 will be overwritten to 1 to make transitioning easier.rBlib)ocmlocklz.bcr2)intr"object __setattr__r-rBwarningswarnr__file__parentr2dictstrtupleitems)self gfx_majorrQdefault_libdirr2rRs r __post_init__zHIPOptions.__post_init__GsG !B( #r/Br 4i8~~!t~~!9K'LQR&R 0/ 0R II ! a MMuvzwAwAvBByz    tWa 0h..6 ,,4b$t?O?O:P # AC">se3K#?@K  A4k6G6G6I0JKrc dj|jjDcgc] \}}|d|c}}}tj|j dj Scc}}w)N_-zutf-8)join__dict__r_hashlibsha256encode hexdigest)r`namevalkeys rhashzHIPOptions.hashZs]hh9L9L9NOID#4&#OP~~cjj12<<>>PsA4 )"__name__ __module__ __qualname__r-rU__annotations__r.r0r1r2r\r3r^r4boolr5r"r]r:rr;r=r>r?r@rArBrCrDrFrHrJrcrprrrr+r+sIsL#JHcK#L%#E4"t"D# (S%*R46%uSz6'--/9 %*9!d!$)T) !#!E3N$$)*!3*L# "#" M3L&?rr+cFeZdZdZedefdZdeddffd ZdefdZ de fdZ dZ d Z deeeffd Zd Zed Zed ZedZedZedZedZedZedZedZdZej:dZxZS) HIPBackendNrc |jdk(S)NrE)backendrs rsupports_targetzHIPBackend.supports_targetbs~~&&rreturncjt||t|jtsJd|_y)Nhsaco)super__init__ isinstancer"r] binary_ext)r`r __class__s rrzHIPBackend.__init__fs+  &++s+++!rc d|jS)Nhip:r'r`optionss rget_target_namezHIPBackend.get_target_nameksgll^$$rc\dtjjxs|jji}|j dddkDr t d|jjdk(rBttj}|jdhtt||d<d|vr%tttj|d<|jjd k(rCttj}|jd d htt||d <d |vrtjj |d <|jtj"j%Dcic]}||vr |||||c}tdi|Scc}w)Nr"r1rz'num_ctas > 1 not supported for AMD GPUsrtf32r>r:r r8r9r;r?r)r runtime override_archrr"get ValueErrorsetr+r>updater^sortedr:r;languagedefault_fp_fusion__dataclass_fields__keys)r`optsargsr>r;ks r parse_optionszHIPBackend.parse_optionsnsu 33Gt{{7G7GH 88J "Q &FG G ;;  x '+.z/V/V+W ( ( / / 938@\9]3^D/ 0 ! -+0 8W8W1X+YD' ( ;;  x '03J4`4`0a - - 4 4j)5L M8=fEf>g8hD4 5 T )',~~'G'GD# $ )H)H)M)M)O;AT d1g&9QZ; <!D!!;sF)c|j|j|j|jd|jd|jdfS)Nrrr/)r-r1sharedr3)r`metadatas r pack_metadatazHIPBackend.pack_metadatasO       OO  ! !! $  ! !! $  ! !! $   rc0dt|jiS)N min_dot_size)rrrs rget_codegen_implementationz%HIPBackend.get_codegen_implementations 0 =>>rcddlm}d|iS)Nr) libdeviceztriton.language.extra.libdevice)triton.language.extra.hipr)r`rs rget_module_mapzHIPBackend.get_module_maps719==rctj|tjr tjj|yyN)r load_dialectsrwinstrumentation)r`ctxs rrzHIPBackend.load_dialectss2 #  % %  & & 4 4S 9 &rcddl}d}t|dr|j|kSt||jr-t|dr!|j j |kSy)Nri ptr_rangeuntyped_storageF)torchhasattrrrTensorrsize)argr MAX_INT_32s r is_within_2gbzHIPBackend.is_within_2gbs] 3 $==?j0 0 c5<< (WS:K-L&&(--/:= =rcHtj|}d|vr|ddggz }|S)NSztt.pointer_rangerO)r parse_attr)descrets rrzHIPBackend.parse_attrs1$$T* $; ',- -C rc tj||fi|}tjjr|dk(rt j |r|dz }|S)Ntensorr)rget_arg_specializationr r use_buffer_opsrwr)rtykwargsrs rrz!HIPBackend.get_arg_specializationsJ00bCFC 99 # #h:;S;STW;X 3JC rctj|j}|jtj j |tjj|tjj|tj j|tjj|tjj|tj j|tjj|tj j|tjj!||j#||Sr)r pass_managercontext enable_debugrcommon add_inlinerttiradd_rewrite_tensor_pointer(add_rewrite_tensor_descriptor_to_pointeradd_canonicalizer add_combineadd_reorder_broadcastadd_cseadd_triton_licmadd_symbol_dceadd_loop_unrollrun)modrrpms r make_ttirzHIPBackend.make_ttirs __S[[ )  !!"% ..r2 <>#w/   *    ( (&O x):):GLL/Z Hgll3 Hc* %%h0H%P %%h0QSWX %%h0H%P %%h0H'J[J[_aJab%224PbB >? A8Bw?P?PQXQbQb?b>c:de A0W5I5I4JL+2+E+E6  A1=A    ( ( F ( ( 2 F # # %   Q(    ( (!(^22U:NN\12NY./NY./E  ! !(E 2  .5.A.AiltTSEXEXYacgEhTiEi5zA~%%h6 Xt'7'7r2wOgOgh  $ $W\\ 2 F ! !"< = F ! !"< = F ! !"< = 99 * *  3 3CF ;!--l;+.+;+;<]+^+cbc'(,/,<,<=c,d,ihi() $$X.   *8}wQ@js8^ ^ ^%3^%c tjd|}t|dk(sJ|d|d<g}|jdk(r|j dd|j vrdnd }t j|tj|j |||jd }tjjrtd t||S) Nz3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)rrrm attentionzsink-insts-to-avoid-spillsgfx11z -real-true16rGFz!// -----// AMDGCN Dump //----- //)refindallr!rJappendr"rtranslate_to_asmr rr?r dump_amdgcnprint)rrrnamesflagsfeaturesamdgcns r make_amdgcnzHIPBackend.make_amdgcns  QSVW5zQ 8  K / LL5 6%, %<>"&&sC,=,=w||XW\^e^v^v',. 99 5 6 &M rcld}tjjrd}tj||j |}t j5}t j5}t|jd5}|j|dddtj|j|jdddt|jd5}|j} dddddd S#1swYtxYw#1swYNxYw#1swY+xYw#1swY SxYw)NrGrwbrb) r r rr assemble_amdgcnr"tempfileNamedTemporaryFileopenrmwrite link_hsacoread) rrrr.r}tmp_outtmp_infd_infd_outrs r make_hsacozHIPBackend.make_hsacos    ( (&O##CG  ( ( * $g,,. :&&++t,'KK&'v{{GLL9 :gllD) $Vkkm $  $ '' : : $ $  $ sTD))DD2DD)"D3D)D DD D)D& "D))D3c|tjk(rfd|d<fd|d<n|tjk(r fd|d<fd|d<fd|d <fd |d <y) Nc*j||Sr)rrrrr`s rrz'HIPBackend.add_stages..s4>>#xQX3Yrrc*j||Sr)rrTs rrz'HIPBackend.add_stages..sDOOCSZ4[rttgirc*j||Sr)rrTs rrz'HIPBackend.add_stages..sD4G4GXW^4_rc*j||Sr)r4rTs rrz'HIPBackend.add_stages..st~~c8W/Urllirc*j||Sr)rBrTs rrz'HIPBackend.add_stages..s1A1A#xQX1YrrAc*j||Sr)rQrTs rrz'HIPBackend.add_stages..sXw0Wrr})rTRITONGLUON)r`stagesrrs` ` r add_stageszHIPBackend.add_stagessR x &YF6N[F7O  '_F7OUvYxWwrc|jSrr)r`s rrpzHIPBackend.hashs++r) rqrrrsr staticmethodrrzrr]rr rrrr rrrrrrrrrr4rBrQr_ functools lru_cacherp __classcell__)rs@rrwrw_sMO' ''"y"T" %#%"S"4 ?>S*_ 5> :   88t  AAF.  XY  rrw)triton.backends.compilerrrrtriton._C.libtritonrrrr tritonr dataclassesr typingr r rtypesrrirGr8rbrXpathlibrrr$r(r+rwrrrrlswEE55!## 0Y0 X r $=?=?=?@n n r