̲j!UdZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddl m Z mZmZmZddlmZmZdd lmZddlZddlZGd d eZGd d Ze eefZdZiZeed< dddZgdZ Gdde!Z"dZ#GddZ$GddZ%dedefdZ&dedede'fdZ(d Z) dbd#ed$ed%eded&e'd'e'deefd(Z*d)ed*ed+ed,ed-e eefd.e'd'e'd/e%d0e'ddfd1Z+d2Z,d3Z-ej.d4Z/d5Z0d6Z1d7Z2d8Z3ej.d9Z4d:Z5ej.d;Z6d<Z7ej.d=Z8d>Z9dcd?Z:d@Z;dAZdDZ?GdEdFZ@GdGdHZAeAZBiZCeAZDiZEe eeFfedI<iZGe D]ZHeIeHesJeHJD]\ZKZLeLdZMeLddZNejOeNvr?eDPeKejQeNvreEReKdJreMeGeK<neMeEeK<ejSeNvr#ejQeNvreBPeKeMeCeK<ej.eBTZUej.dKeDTdLZVej.dMZWej.dNZXej.dOZYej.dPZZ d)ed*ed+ed,ed-e eefd.e'd'e'd/e%d0e'defdQZ[dcdRZ\dSZ]dTZ^ej.dUZ_dVZ`dWZa ddd\ed]e'ded^ed)ed,ed$ed_ed&e'd%ed0e'd.e'd'e'd`e'd/ee%def daZbdS)ea The Python Hipify script. ## # Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. # 2017-2018 Advanced Micro Devices, Inc. and # Facebook Inc. All rights reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal # in the Software without restriction, including without limitation the rights # to use, copy, modify, merge, publish, distribute, sublicense, and/or sell # copies of the Software, and to permit persons to whom the Software is # furnished to do so, subject to the following conditions: # # The above copyright notice and this permission notice shall be included in # all copies or substantial portions of the Software. # # THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR # IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, # FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE # AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER # LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. N) constants)CUDA_TO_HIP_MAPPINGS)MATH_TRANSPILATIONS)DictListIteratorOptional)MappingIterable)EnumceZdZdZdZdS) CurrentStaterN)__name__ __module__ __qualname__ INITIALIZEDDONEQC:\Projects\Historical\venv\Lib\site-packages\torch/utils/hipify/hipify_python.pyrr,sK DDDrrceZdZdZdZdS) HipifyResultc0||_||_d|_dS)N current_state hipified_pathstatus)selfrrs r__init__zHipifyResult.__init__1s** rc6d|jd|jd|jS)NzHipifyResult:: current_state: z, hipified_path : z , status: rr!s r__str__zHipifyResult.__str__6sGC1CCCW[WiCCuyvACC DrN)rrrr"r%rrrrr0s7 DDDDDrrz;// !!! This is a file automatically generated by hipify!!! HIPIFY_FINAL_RESULTscalar_t)DtypeT)! InputErroropenfbcolorsGeneratedFileCleanermatch_extensionsmatched_files_iterpreprocess_file_and_save_result compute_statsadd_dim3processKernelLaunchesfind_closure_groupfind_bracket_groupfind_parentheses_groupreplace_math_functionship_header_magicreplace_extern_sharedget_hip_file_pathis_out_of_placeis_pytorch_fileis_cusparse_fileis_special_fileis_caffe2_gpu_filer?Trie preprocessorfile_specific_replacementfile_add_headerfix_static_global_kernelsextract_argumentsstr2boolrrhipifyc$eZdZfdZdZxZS)r*cXt|||_dSN)superr"message)r!rL __class__s rr"zInputError.__init__Ms& !!! rcd|jS)Nz Input error: )rLr$s rr%zInputError.__str__Qs-t|---r)rrrr"r% __classcell__)rMs@rr*r*JsG.......rr*c&t||dS)Nignore)errors)open)filenamemodes rr+r+Us $x 0 0 00rc.eZdZdZdZdZdZdZdZdZ dZ d S) r,zzzzzzzzN) rrrHEADEROKBLUEOKGREENWARNINGFAILENDCBOLD UNDERLINErrrr,r,Zs6 F FGG D D DIIIrr,c4eZdZdZd dZdZdZd dZdZdS) r-z+Context Manager to clean up generated filesFcH||_t|_g|_dSrJ)keep_intermediatessetfiles_to_clean dirs_to_clean)r!ras rr"zGeneratedFileCleaner.__init__os$"4!eerc|SrJrr$s r __enter__zGeneratedFileCleaner.__enter__ts rctj|s7|jtj|t |g|Ri|SrJ)ospathexistsrcaddabspathrS)r!fnargskwargss rrSzGeneratedFileCleaner.openws[w~~b!! 9   # #BGOOB$7$7 8 8 8B(((((((rctj|\}}|s"tj|\}}|r8|r6tj|s||dtj|r|sMtj||jtj |dSdS)NT)exist_ok) rhrisplitrjmakedirsisdirmkdirrdappendrl)r!dnrqparentns rrszGeneratedFileCleaner.makedirs|sGMM"%%  . f--IFA  1a 1v 6 6 1 MM&4M 0 0 0w}}R   ; ; HRLLL   % %bgoob&9&9 : : : : : ; ;rc|jsE|jD]}tj||jdddD]}tj|dSdS)N)rarcrhunlinkrdrmdir)r!typevalue tracebackfds r__exit__zGeneratedFileCleaner.__exit__sn& (   ! '"-       rNF) rrr__doc__r"rfrSrsrrrrr-r-mso55    ))) ;;;;rr-rireturncB|tjdS)N/)replacerhsep)ris r _to_unix_pathrs << $ $$rrT extensionsc:tfd|DS)zz#match_extensions..s188x  ##888888rany)rTrs` rr.r.s& 8888Z888 8 88rc:tfd|DS)Nc3BK|]}tj|VdSrJ)fnmatch)rpatternfilepaths rrz_fnmatch..s/JJgwx11JJJJJJrr)rpatternss` r_fnmatchrs& JJJJJJJ J JJrrF root_pathincludesignoresout_of_place_onlyis_pytorch_extensionc#Kt|}tj|dD]W\}}} tj||} | dkr`d|vr|dd|vr|dd|vr*|d|d| D]} ttj|| } ttj| | } t| |r[t| |sKt| |s| |vr7|s1t| st| s|rt| s| VȐYdS)NT)topdown.z.gitbuild third_partyzthird_party/nvfuser)rbrhwalkrirelpathremovervrjoinrr.r<r?r;)rrrrrr exact_matches abs_dirpathdirs filenames rel_dirpathrTr rel_filepaths rr/r/sMMM+-')T*J*J*J&dIgook9== #  ~~ F###$ G$$$$$ M*** 1222!  H$RW\\+x%H%HIIH(k8)L)LMML8,, !(G44 &h ;; @H=?X?X+!*<88!AST`AaAa! (!1N1N!  routput_directoryr all_filesheader_include_dirsstatship_clang_launch clean_ctx show_progressc ^tjtj||} t t j| } | t| <t||||||||| } |r'd| j vrt| d| j | j d| t| <dS)N)rrignoredz->T)flush) rhrirlrrrrr&rAr printr) rrrrrrrrrfin_path hipify_resultresults rr0r0swrw||,z compute_stats..s\\\'= 9\\\runsupported_callsz1Total number of unsupported CUDA function calls: r, z+ Total number of replaced kernel launches: kernel_launches)rlenr)rrs rr1r1s\\GZA[\\\ XcBS>T>T X X XYYY $))% & &''' ZUCT=U9V9V Z Z Z[[[[[rcd}d}|dddd}dtdD}d||d<t|D]q\}}|dkrne|d kr|dz }n |d kr|dz}|d ks|t|dz kr1|dkr+||d kz||d <|dz }|dkr|dz||d<r||dd|dd dz}||dd|dd }||dd|dd d dd} ||dd|dd d dd} d| d } d| d } || | } || | }|||z| |z}|S)zBadds dim3() to the second and third arguments in the kernel launchr<<>>cg|]}iSrr)r_s r zadd_dim3..s%;%;%;Qb%;%;%;rrstartr(),end  zdim3()rrange enumeraterstrip) kernel_string cuda_kernelcountclosurearg_locsindc first_arg_rawsecond_arg_rawfirst_arg_cleansecond_arg_cleanfirst_arg_dim3second_arg_dim3first_arg_raw_dim3second_arg_raw_dim3s rr2r2sL EG!))%44<z'processKernelLaunches..s(1.Osyy||.O.O.Orc|d|ddddddddd}ddi}d}d}d }d }|}t|d ddz ddD]} |} |||fvr^| d kr#||kr |}||d d<|dxxdz cc<| dkr/|dxxdzcc<|ddkr||kr ||d d<|}||kr |s |dvr>||kr |}||dd<|dkr$d|dd<|d|d |d gcS||kr$||dd<|d|d |d gcSdS)Nrrrrr{) kernel_launch kernel_nametemplatez<>rrrr>r<>r#rr:r)risalnum) in_kernelposrSTART AT_TEMPLATEAFTER_TEMPLATEAT_KERNEL_NAMEr icharstrings rgrab_method_and_templatez7processKernelLaunches..grab_method_and_templates(1'9)EBRSS%'33"$R00  q  s?+G4q8"bAA% a% aA!9D%---3;;!,12J.$KKK1$KKK3;;$KKK1$KKKT{a''V{-B-B34J0!/$$!9$$&&a&)7P*P*P//!/45M*51Avv67M*73"%]!3s:#oJ^```` //67M*73"%]!3s:#oJ^````K% a% arc@d}g}|d|dkr|d|}|d|dz}|dkrtd||||||d|d|dk|S)zKFinds the starting and ending points for all kernel launches in the string.rrr{rrzno kernel end found)rrr)findr*rv)r kernel_endkernel_positions kernel_starts rfind_kernel_boundsz1processKernelLaunches..find_kernel_boundsSs kk%,,22!;;uj99L UL99A=JQ !6777  # #l:.4\:5M.N%P%P Q Q Qkk%,,22 rc0d}d}d}|D]}|dkr3|dkr |dkrd}nb|dkr |dkrd}nS|dkr|dkr|dkrd}n>|dkr|d ks|d krd}n)|dkr|dkr|dkrd}n|dkr|dkr|dkrd}|}|dkr||z }|d z }|S) Nrrz//*z/*"\' rxr)r in_commentprev_c new_stringrs r mask_commentsz,processKernelLaunches..mask_commentsks    " "AR88# !%JJ#XX&C--!%JJ#XX&D..Vs]]!$Jt##99T !#Jt##88# !#Js""88$!#JFRa c! rrrrrrr{rrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr2rrErrv)rrrrrget_kernel_positions output_stringkernelparams parenthesisrrend_param_indexkernel_name_with_templatecuda_kernel_dim3num_klp hip_kernels` rr3r3sA ! !"O"OQW X XF9a9a9a9a9av   0B 2 2==3H3H I IJJM'44))&11kk#ve}55 VAYw/ a?@ vgve}<= %ay/255!!1$*6!9W+=f_>UV[>\_`>`+`$a!#M;??'6'?+B+B5#+N+N+V+VW\^a+b+bccdd*-=ad-C-K-K 5AK(50.2.229'%2F2Fww 4HH !:CB[<[^a>> find_closure_group("(hi)", 0, ["(", ")"]) (0, 3) Fr)r{r{Tr)NN)r) input_stringrrinside_parenthesisparensrp_startp_ends rr4r4s F CNGU L!! ! !  a ( (!U**%)"!  # %( * */A * aKF{{~% q L!! ! ! :rc*t||ddgS)z%Finds the first balanced parantheses.{}rr4r#rs rr5r5 lE#s D D DDrc*t||ddgS)z!Finds the first balanced bracket.rrrr+r,s rr6r6r-rz\bassert[ ]*\(cl|}tD])}||dt|d}*|S)aFIXME: Temporarily replace std:: invocations of math functions with non-std:: versions to prevent linker errors NOTE: This can lead to correctness issues when running tests, since the correct version of the math function (exp/expf) might not get called. Plan is to remove this function once HIP supports std:: math function calls inside device code r)rr)r#rfuncs rr7r7sM!M#\\%--$kkk>QRV>W;Z;Z;Z[[ rz:?:?\b(__syncthreads)\b(\w*\()c|ddg}tfd|DrSdv}|dvz }|dvz }|tduz }|rd|zS) aIf the file makes kernel builtin calls and does not include the cuda_runtime.h header, then automatically add an #include to match the "magic" includes provided by NVCC. TODO: Update logic to ignore cases where the cuda_runtime.h is included by another file. zhip/hip_runtime.hzhip/hip_runtime_api.hc3PK|] }tjd|d|dV!dS)z #include ("z"|)N)research)rextrs rrz#hip_header_magic..sC W W293C33C333] C C W W W W W WrhipLaunchKernelGGL __global__ __shared__Nz#include "hip/hip_runtime.h" )rRE_SYNCTHREADSr4)r#headershasDeviceLogicrs @rr8r8s!M#$;\s]+)\s+(\w+)\s*\[\s*\]\s*;cB|}td|}|S)aMatch extern __shared__ type foo[]; syntax and use HIP_DYNAMIC_SHARED() MACRO instead. https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md#__shared__ Example: "extern __shared__ char smemChar[];" => "HIP_DYNAMIC_SHARED( char, smemChar)" "extern __shared__ unsigned char smem[];" => "HIP_DYNAMIC_SHARED( unsigned char, my_smem)" cd|dpdd|dd|ddS) NzHIP_DYNAMIC_SHARED(rrrrrrrrrs rrz'replace_extern_shared..sF^#))A,,*<"^^syy||^^syyYZ||^^^r)RE_EXTERN_SHAREDr)r#rs rr9r9s1!M$((^^`mooM rctj|rJ|st|s|Stj|\}}tj|\}}|dkrd}|}|}|dd}|dd}|dd}|dd}|dd}|d kr|dd}|s&||kr tj|d}|r||kr||z|kr|d z}tj|||zS) z3 Returns the new name of the hipified file .cu.hipcudahipCUDAHIPTHCTHHz caffe2/core_hip)rhriisabsr;rrsplitextrr)rrdirpathrTrootr5 orig_filename orig_dirpaths rr:r:!sn w}}\**** *  (E(E l33GX  **ID#H e||MLoofe,,Goofe,,GooeU++G << & &D << & &D-||E5)) /G|$;$;',,w..< 7 7TCZMr>sk|$$ |))++ + +4 ++-- - -*l.@.@.B.BBBu4 5rc tj|rJ|drdStj|}tj|\}}d|vs|dvod|vS)Nzc10/cudaTgpur@.cuhcudnn)rhrirIrTbasenamerJ)rrTrr5s rr?r?sw}}\**** *z**tw --H W  h ' 'FAs X  7!7 VgX>UVrceZdZdZdZdS)TrieNodezA Trie node whose children are represented as a directory of char: TrieNode. A special char '' represents end of word ci|_dSrJ)childrenr$s rr"zTrieNode.__init__s  rN)rrrrr"rrrreres-rrec\eZdZdZdZdZdZdZdZe j dZ dZ d Z d S) r@zCreates a Trie out of a list of words. The trie can be exported to a Regex pattern. The corresponding Regex should match much faster than a simple Regex union.ct|_tj|_|j|_dS)z,Initialize the trie with an empty root node.N)rerLhashlibmd5_hashdigest_digestr$s rr"z Trie.__init__s3JJ []] z((** rc.|j||j|_|j}|D]6}|j|t|j|}7d|jd<dS)zAdd a word to the Trie. TrN) rlupdateencodermrnrLrg setdefaultrer!wordnoders rrkzTrie.adds $++--(((z((** y ' 'D M $ $T8:: 6 6 6=&DD  brc|jS)zReturn the root node of Trie. )rLr$s rdumpz Trie.dumps yrc*tj|S)z Escape a char for regex. )r3escape)r!rs rquotez Trie.quotesyrc^|j}|D]}||jvr|j|}dSd|jvS)zZSearch whether word is present in the Trie. Returns True if yes, else return FalseFr)rLrgrss rr4z Trie.searchsKy  Dt}$$}T*uuT]""rc|}d|jvr,t|jdkrdSg}g}d}t|jD]}t |j|t r ||j||j}|| ||zu#t$r+|| |YwxYwd}t|dk } t|dkr]t|dkr||dn.|dd |zdzt|dkr |d} ndd |zd z} |r| r| d z } nd| d } | S) zConvert a Trie into a regular expression pattern Memoized on the hash digest of the trie, which is built incrementally during add(). rrNr[]z(?:|r?z)?) rgrkeyssorted isinstancere_patternrnrvrz Exceptionr) r!rLrmrualtccqrrecursecconlyrs rrz Trie._patterns   3t}'9'9';';#<#<#A#A4  4=--//00  D$--x88 0"mmDM$,?NNGJJtzz$//'9:::: 000IIdjj../////0XX\! r77Q;;2ww!|| 2a5!!!! 3,s2333 s88q==VFFSXXc]]*S0F  * *# )v))) sAC2D DcB||j|jSz#Export the Trie to a regex pattern.rrLrnr$s rrz Trie.pattern}}TY 555rcB||j|jSrrr$s rexport_to_regexzTrie.export_to_regexrrN)rrrrr"rkrwrzr4 functools lru_cacherrrrrrr@r@sSS+++ ! ! ! # # #)))V66666666rr@ PYTORCH_MAPrz(?<=\W)(z)(?=\W)z#include "([^"]+)"z#include <([^>]+)>z"#define THC_GENERIC_FILE "([^"]+)"z\.cu\bc P tjtj|t} |vr!d| _d| _tj| _ | Sttj |} td5} | tkr-d| _d| _tj| _ | cdddS| d| } dddn #1swxYwY| } tjtjt#| }tjtj|s2tj|dfd}rt*| } ntt/| rt*|| } nIt1| rt*| } nd }t2|| } d#f d }t4|d d | } t6|d d| } t8|d| } |drG| dd} | dd} t>d| } stA| } |drd|vrtC| } tE| } rg| | kratjtj|kr!| _d| _tj| _ | S|krtGdr t| z} d }tj|rAt|d5}|| k}dddn #1swxYwY|r  |dd5}|$| dddn #1swxYwY|| _d| _tj| _ | S#tJ$rm}tMtNj(d|d|j)ddtNj* tVj, | _d!| _tj| _ | cYd}~Sd}~wwxYw|| _d"| _tj| _ | S)$z< Executes the CUDA -> HIP conversion on the specified file. Nz[ignored, not to be hipified]zutf-8)encodingz#[ignored, input is hipified output]rcBt|dSNr)rrms rpt_replzpreprocessor..pt_replZs1771::&&rcrt|d|Sr)PYTORCH_SPECIAL_MAPgetr)rrs rpt_special_replz%preprocessor..pt_special_repl]s+"&&qwwqzz771::>>>rcBt|dSr) CAFFE2_MAPrrs rc2_replzpreprocessor..c2_repljs!!''!**--rTc *  f d}|S)Nc  |d}tj| |ds*|drK|ds6t |dSrt fd Drkd}d}rtj}tj tj ||}tj |r|}|}|D]}tj |}tj tj ||}tj |r|}|}||dS|tvrt|   n|tvrt|}|jtjkrtj|}tj tj t |} | |_|t|<tj| | n||St|j} tj| | n||S|dS)Nr)z ATen/cudazATen/native/cudazATen/native/nested/cudazATen/native/quantized/cudazATen/native/sparse/cudazATen/native/transformers/cudazTHC/rFTHCPc3BK|]}|VdSrJr)rsrTs rrz>preprocessor..mk_repl..repl..s/,U,UaQZZ-A-A,U,U,U,U,U,Urr)rrhrircrTformatr:rdirnamerlrrjr&r0rrrrr)rr header_dirheader_filepathheader_dir_to_checkheader_path_to_checkheader_include_dir header_resultheader_rel_pathheader_fout_pathhipified_header_filepathrTrrrrrinclude_current_dirrrrrtempls @rreplz+preprocessor..mk_repl..replps< Aw''**H &'' Ye$$ Y./\\&-A-A Y||$5aggajjBV$W$WXXX$( W,U,U,U,U9,U,U,U)U)U( W! "&&?*,'//(*C*C'+-7??27<"*.ACC*.0gll;KM_.`.`+/1wrw||L_ab?c?c/d/d,7>>*>??C)+O+](||BGOOPhPt4L4L9H*%V%VWWW771:: rr) rrrrrrrrrrrrs `` rmk_replzpreprocessor..mk_replosX9 9 9 9 9 9 9 9 9 9 9 9 9 9 9 t rz#include "{0}"z#include <{0}>Fz#define THC_GENERIC_FILE "{0}"zCMakeLists.txtrDrErFrGrAr` PowKernelz[skipped, no changes])r@ra.c.cc.cpp.h.hppwz[ok]zFailed to save z with "z ", leaving z unchanged.filez[skipped, no permissions]zkipped, already hipified])T)-rhrirlrr&rr rrrrrrSreadlineHIPIFY_C_BREADCRUMBseekreadr:rjrrsRE_PYTORCH_PREPROCESSORrr>r<RE_CAFFE2_PREPROCESSORRE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILErr RE_CU_SUFFIXr3r7r8r.writePermissionErrorrr,rZstrerrorr\sysstderr)rrrrrrrrrrrfin output_sourceorig_output_source fout_pathrrrdo_writefout_oldfoutrrrs` ``````` @@rrArA4swrw||, &2&7 # ;K!L!LMML h ) ) )#S <<>>0 0 0*.M '#HM *6*;M ' ########   ###############' -=?PQ]_s?t?t u uvvI 7>>"'//)44 5 5727??955666'''?????  O/33G]KK < ( ( O377WWMM \ * * O377OOMM . . .266w NNM;;;;;;;;;;;;;;x$''0@$(G(GWWM#''0@%(H(H-XXM'++GG4T,U,UWdeeM)**@%--fe<< %--eU;; $((??  D-mUCC  /**> 80K0K.}== %]33M  - / / GOOH % %)C)C C C&. #6 &2&7 #9!1(M *6*;M '  !'0 #< &2&7 #sm8>D8)D88D<?D<4RR R'T ?S! T !S%%T (S%)#T VA"U?9V?Vct|d5}|}|r0tjdtj|dfd|}n||}|d|||ddddS#1swxYwYdS)Nr+z\b(z)\bcSrJr)rreplace_strings rrz+file_specific_replacement..snrr) r+rr3rryrrrtruncate)r search_stringrstrictrcontentss ` rrBrBs x  !6688  GvBRY}%=%=BBBD\D\D\D\^fggHH'' ~FFH q   sBB::B>B>cJt|d5}|}|ddkr|ddkrd|d}d|d|z}|d|||ddddS#1swxYwYdS) Nrrrr{rr z #include z )r+rrrr)rheaderrrs rrCrCs x  !6688 !9  r c 1 1"]]]F++++x7 q   sA:BBBc2|dd}|S)z.ds)aaaWGOO$57GHHaaarc<g|]}|Srr)rrQrrs rrzhipify..es)]]]66>>"35EFF]]]r)rrrrrr)Pathc3K|]v}|tt|3tt|Pt|jet|VwdSrJ)is_filerstrr.name)rrirrrs rrzhipify..s    TH-- c$ii11 !,=>>  II      rr T)ra)rrz-Successfully preprocessed all matching files.r)rhgetcwdrirjrrexitrstripshutilcopytreermaprr/rbrIrrvpathlibrextendrglobr-r0r,rYr\rr1r&)rrrrrrrrrrrrrrrr all_files_setrrrheader_include_dir_pathrrs` `` ` ` rrGrGEsE"BIKK 7>>+ , , <===   6  %%%,v5,,,aaaaaX`aaa]]]]]U\]]] 7>>* + +=)+;<<<C x0011H3}g..//G'(8807J:K=QSSSTTI NNM   w}}Q 2 -q11A M ! !   Q   1    7==+ , , _&*d+=&>&> # #&*d27<<8HJ\+]+]&^&^ #      "9"?"?"D"D       (DAAA 46"MME&=NYY;qq'(8(IOb(-/?AUW`bo q q q q '/K Kgl Zadaklllle r)rrrFFr)FrrrrrrFrTFFFN)crrrr3r rrhrrcuda_to_hip_mappingsrrtypingrrr r collections.abcr r enumr rrjrrrHipifyFinalResultrr&__annotations__PYTORCH_TEMPLATE_MAP__all__rr*r+r,r-rboolr.rr/r0r1r2compilerr3r4r5r6 RE_ASSERTr7r9r8r>r9r:r;r<r=r>r?rer@ CAFFE2_TRIEr PYTORCH_TRIErobjectrmappingritemssrcrdst meta_data API_CAFFE2rk API_SPECIALr API_PYTORCHrrrrrrrrArBrCrD RE_INCLUDErErFrGrrrr(s0 666666555555111111111111--------     4   DDDDDDDDl*+T)+&+++!+*== s s s........111 &D%%%%%%9s99T9999 KKK  !"'%* ))))) )  ) # ) 08} ))))X++++& + CI +  +#+(++!%++++2 \ \ \   F2:>??WWWt###LEEE EEE BJ( ) )     =>>:2:hii   E-E-E-E-P    WWW\6\6\6\6\6\6\6\6|dff  tvv !# T#v+ ####""G :gw ' ''' 'mmoo " " UAh!""I  y 0 0   S ! ! !$ 11kooc26N6N1+.#C((#& C  1 1i6KS\6\6\ OOC !JsO "$K$?$?$A$ABB$"*%X1M1M1O1O%X%X%XYY"*233"*233 bj!FGGrz)$$ vvvv& v CI v  v#v(vv!-vvvvp     RZ( ) ) +++\DDD T"8$&#"!&$)04RRRRR R  R " RRRRRRRR"R,-R !RRRRRRr