Sjo!UdZddlZddlZddlZddlZddlZddlZddlZddlm Z ddlm Z ddlm Z ddl m Z ddl mZmZdd lmZddlZddlZd ZGd d eZGd dZeeefZdZiZeed< dddZgdZGddeZ dZ!GddZ"GddZ#dedefdZ$dedede%fd Z&d!Z' d^d"ed#ed$eded%e%d&e%de efd'Z(d(ed)ed*ed+ed,eee)fd-e%d&e%d.e#d/e%ddfd0Z*d_d1Z+d2Z,ejZd3Z.d4Z/d5Z0d6Z1d7Z2ejZd8Z3d9Z4ejZd:Z5d;Z6ejZd<Z7d=Z8d`d>Z9de%fd?Z:de%fd@Z;dAZGdDdEZ?GdFdGZ@e@ZAiZBeeeCfedH<e D]AZDeEeDeseFdIeDjD]\ZHZIeAjeHeIeBeH<CejZdJeAjdKZLejZdLZMejZdMZNejZdNZOejZdOZP d(ed)ed*ed+ed,eee)fd-e%d&e%d.e#d/e%defdPZQd`d_dQZRd_dRZSdSZTejZdTZUdUZVdVede%fdWZW dadXedYe%dedZed(ed+ed#ed[ed%e%d$ed/e%d-e%d&e%d\e%d.e#dzdef d]ZXy)ba 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)CUDA_TO_HIP_MAPPINGS)MATH_TRANSPILATIONS)CAFFE2_PATH_MAPPINGS)Iterator)MappingIterable)EnumcBtjd|tdy)Nz-hipify version 2.0.0 no longer uses function ) stacklevel)warningswarn FutureWarning)names _/var/www/html/ia-enhanced/venv/lib/python3.12/site-packages/torch/utils/hipify/hipify_python.py _deprecatedr-s MMA$H-defceZdZdZdZy) CurrentStaterr N)__name__ __module__ __qualname__ INITIALIZEDDONErrrr0s K Drrc eZdZddZdefdZy) HipifyResultreturnNc.||_||_d|_y)N current_state hipified_pathstatus)selfr#r$s r__init__zHipifyResult.__init__5s** rcTd|jd|jd|jS)NzHipifyResult:: current_state: z, hipified_path : z , status: r"r&s r__str__zHipifyResult.__str__:sF01C1C0DDVW[WiWiVjjtuyvAvAuBC DrrN)rrrr'strr*rrrrr4s DDrrz;// !!! 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_fileTrie preprocessorfile_specific_replacementfile_add_headerfix_static_global_kernelsextract_argumentsstr2boolrrhipifyc,eZdZdfd ZdefdZxZS)r1rc2t||||_yN)superr'message)r&rS __class__s rr'zInputError.__init__Qs ! rc d|jS)Nz Input error: )rSr)s rr*zInputError.__str__Ust||n--rr+)rrrr'r,r* __classcell__)rTs@rr1r1Ns..rr1ct||dS)Nignore)errors)open)filenamemodes rr2r2Ys $x 00rc,eZdZdZdZdZdZdZdZdZ dZ y ) r3zzzzzzzzN) rrrHEADEROKBLUEOKGREENWARNINGFAILENDCBOLD UNDERLINErrrr3r3^s* F FGG D D DIrr3c6eZdZdZdd dZdZdZdd dZdZy) r4z+Context Manager to clean up generated filesNc>||_t|_g|_yrQ)keep_intermediatessetfiles_to_clean dirs_to_clean)r&rhs rr'zGeneratedFileCleaner.__init__ss"4!erc|SrQrr)s r __enter__zGeneratedFileCleaner.__enter__xs rctjj|s8|jj tjj |t |g|i|SrQ)ospathexistsrjaddabspathrZ)r&fnargskwargss rrZzGeneratedFileCleaner.open{sHww~~b!    # #BGGOOB$7 8B((((rctjj|\}}|s"tjj|\}}|r4|r2tjj|s|j |dtjj |r|sNtj ||jjtjj|yy)NT)exist_ok) rorpsplitrqmakedirsisdirmkdirrkappendrs)r&dnrxparentns rrzzGeneratedFileCleaner.makedirssGGMM"%  f-IFA av 6 MM&4M 0ww}}R  HHRL    % %bggoob&9 :)1rc|jsS|jD]}tj||jdddD]}tj |yy)N)rhrjrounlinkrkrmdir)r&typevalue tracebackfds r__exit__zGeneratedFileCleaner.__exit__sY&&((  !  ''"-   'rFr+) rrr__doc__r'rmrZrzrrrrr4r4qs5 ) ;rr4rprcB|jtjdS)N/)replacerosep)rps r _to_unix_pathrs << $$rr[ extensionsc,tfd|DS)zz#match_extensions..s8x  #8any)r[rs` rr5r5s 8Z8 88rc,tfd|DS)Nc3JK|]}tj|ywrQ)fnmatch)rpatternfilepaths rrz_fnmatch..sJgwx1Js #r)rpatternss` r_fnmatchrs JJ JJr root_pathincludesignoresout_of_place_onlyis_pytorch_extensionc# Kt|}tj|dD]\}}} tjj ||} | dk(rPd|vr|j dd|vr|j dd|vr"|j d|j d| D]Z} ttjj|| } t| |s9t| |rFt| |s| |vsW| \yw)NT)topdown.z.gitbuild third_partyzthird_party/nvfuser) rirowalkrprelpathremover}rjoinrr5) rrrrrr exact_matches abs_dirpathdirs filenames rel_dirpathr[rs rr6r6sMM+-'')T*J&dIggook9= # ~ F#$ G$$ M* 12! H$RWW\\+x%HIH8,!(G4%h ;x=?X sCD D)D: Doutput_directoryr all_filesheader_include_dirsstatship_clang_launch clean_ctx show_progressc ftjjtjj||} t t j | } | t| <t||||||||| } |r2d| jvr$t| d| j| jd| t| <y)N)r#r$ignoredz->T)flush) rorprsrrrrr-rHr%printr$) rrrrrrrrrfin_path hipify_resultresults rr7r7swwrww||,TUV=W XY $))% &' 8UCT=U9VWX8Y Z[]s A&cd}d}|jddjdd}tdDcgc]}i}}d||d<t|D]b\}}|dkDrnX|dk(r|dz }n |d k(r|dz}|d k(s|t|dz k(s9|dk(s?||d k7z||d <|dz }|dksX|dz||d<d||dd|dd dz}||dd|dd } ||dd|dd jd dj d } ||dd|dd jd dj d } d| d } d| d } |j| | }| j| | }|j|| z||z}|Scc}w)zBadds dim3() to the second and third arguments in the kernel launchr<<>>r startr(),end  zdim3()rrange enumeraterstrip) kernel_string cuda_kernelcountclosure_arg_locsindc first_arg_rawsecond_arg_rawfirst_arg_cleansecond_arg_cleanfirst_arg_dim3second_arg_dim3first_arg_raw_dim3second_arg_raw_dim3s rr9r9s EG!))%4<z'processKernelLaunches..s!1syy|nB.Orc|d|ddddddddd}ddi}d}d}d }d }|}t|d ddz ddD]} |} |||fvrJ| d k(r||k(r |}||d d<|dxxdz cc<| dk(r$|dxxdzcc<|ddk(r||k(r ||d d<|}||k7s^ |js |dvr-||k7r |}||dd<|dk(sd|dd<|d|d |d gcS||k(s||dd<|d|d |d gcSy)Nrrrrr) kernel_launch kernel_nametemplatez<>rrr r>r<>r#rr:r)risalnum) in_kernelposrSTART AT_TEMPLATEAFTER_TEMPLATEAT_KERNEL_NAMEr%icharstrings rgrab_method_and_templatez7processKernelLaunches..grab_method_and_templates(1'9)EBRS%'3"$R0 q  s?+G4q8"bA% aA!9D%--3;!,12J.$K1$K3;$K1$KT{a'V{-B34J0!/$!9$$&&)7P*P/!/45M*51Av67M*73"%]!3s:#oJ^``/67M*73"%]!3s:#oJ^``K% arcd}g}|jd|dk7re|jd|}|jd|dz}|dkr td|j|||||d|jd|dk7re|S)zKFinds the starting and ending points for all kernel launches in the string.rrrrrzno kernel end found)rrr)findr1r})r kernel_endkernel_positions kernel_starts rfind_kernel_boundsz1processKernelLaunches..find_kernel_boundsQs kk%,2!;;uj9L UL9A=JQ !677  # #l:.4\:.N%P Qkk%,2 rcd}d}d}|D]z}|dk(r,|dk(r|dk(rd}nT|dk(r|dk(rd}nG|dk(rB|dk7r=|dk7r8d}n5|dk(r |d k(s|d k(r&d}n#|dk(r |dk(r|dk(rd}n|dk(r |dk(r|dk7rd}|}|dk(r||z }v|d z }||S) Nr!rz//*z/*"\' rxr)r in_commentprev_c new_stringrs r mask_commentsz,processKernelLaunches..mask_commentsis   "AR8# !%J#X&C-!%J#X&D.Vs]!$Jt#9T !#Jt#8# !#Js"8$!#JFRa c! 3 "4rrrrrrrrrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr9rrLrr})rrrr rget_kernel_positions output_stringkernelparams parenthesisrrend_param_indexkernel_name_with_templatecuda_kernel_dim3num_klp hip_kernels` rr:r:s ! !"OQW XF9av 0B 2=3H IJM'4)&1kk#ve}5 VAYw/ a@ vgve}= %ay/25!1$*6!9W+=f_>UV[>\_`>`$a!#M;?'6'?+B+B5#+N+V+VW\^a+bcd*-=a-C-K-K 5AK(50.229'%2Fww 4H !:CB[<[^a>> find_closure_group("(hi)", 0, ["(", ")"]) (0, 3) Fr)rrTr)NN)r) input_stringrrinside_parenthesisparensrp_startp_ends rr;r;s F CNGU L! !  a (!U*%)"!  # %( */A aKF{~% q L! ! rc"t||ddgS)z%Finds the first balanced parentheses.{}rr;r%rs rr<r< lE#s DDrc"t||ddgS)z!Finds the first balanced bracket.rrrr-r.s rr=r=r/rz\bassert[ ]*\(c^|}tD]!}|j|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 rr>r>sJ!M#\%--$qk>QRV>W=XXY;Z[ \ rz:?:?\b(__syncthreads)\b(\w*\()c|ddg}tfd|DrSdv}|dvz }|dvz }|tjduz }|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.hc3XK|]!}tjd|d|d#yw)z #include ("z"|)N)research)rextrs rrz#hip_header_magic..s+ W299 C5C53] C Ws'*hipLaunchKernelGGL __global__ __shared__Nz#include "hip/hip_runtime.h" )rRE_SYNCTHREADSr6)r%headershasDeviceLogicrs @rr?r?s!M#$;\s]+)\s+(\w+)\s*\[\s*\]\s*;c8|}tjd|}|S)a Match 'extern __shared__ type foo[];' syntax and use HIP_DYNAMIC_SHARED() MACRO instead. See: https://github.com/ROCm/hip/blob/master/docs/markdown/hip_kernel_language.md#__shared__ Examples: "extern __shared__ char smemChar[];" => "HIP_DYNAMIC_SHARED( char, smemChar)" "extern __shared__ unsigned char smem[];" => "HIP_DYNAMIC_SHARED( unsigned char, my_smem)" c|d|jdxsdd|jdd|jddS) NzHIP_DYNAMIC_SHARED(rr!rr rrrrrs rrz'replace_extern_shared..s<)#))A,*<")=Qsyy|nBsyyYZ|n\]^r)RE_EXTERN_SHAREDr)r%rs rr@r@s(!M$((^`moM rctjj|r td|s t |s|Stjj |\}}tjj |\}}|dk(rd}|}|}|jdd}|jdd}|jdd }|jdd}|jdd}|d k7r|jdd }|s%||k(r tjj|d}|r||k(r ||z|k(r|d z}tjj|||zS) z3 Returns the new name of the hipified file $rel_filepath must be a relative path.cu.hipcudahipCUDAHIPTHCTHHz caffe2/core_hip) rorpisabsAssertionErrorrBrysplitextrr) rel_filepathrdirpathr[rootr7 orig_filename orig_dirpaths rrArA"s5  ww}}\"CDD  (E l3GX  *ID#H e|MLoofe,Goofe,GooeU+G << &D << &D-||E5) G|$;'',,w.< 7TCZMUVrceZdZdZddZy)TrieNodezA Trie node whose children are represented as a directory of char: TrieNode. A special char '' represents end of word Nci|_yrQ)childrenr)s rr'zTrieNode.__init__s  rr+)rrrrr'rrrriris rricbeZdZdZd dZd dZdZdZdZe jdZ d Z d Z y) rGzCreates 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.Nct|_tjd|_|jj |_y)z,Initialize the trie with an empty root node.F)usedforsecurityN)rirQhashlibmd5_hashdigest_digestr)s rr'z Trie.__init__s/J [[7 zz((* rc>|jj|j|jj|_|j }|D]5}|j j|t|j |}7d|j d<y)zAdd a word to the Trie. Tr!N) rqupdateencoderrrsrQrk setdefaultrir&wordnoders rrrzTrie.addsz $++-(zz((* yy 'D MM $ $T8: 6==&D '! brc|jS)zReturn the root node of Trie. )rQr)s rdumpz Trie.dumps yyrc,tj|S)z Escape a char for regex. )r5escape)r&rs rquotez Trie.quotesyyrc|j}|D]!}||jvr|j|}!yd|jvS)zZSearch whether word is present in the Trie. Returns True if yes, else return FalseFr!)rQrkrxs rr6z Trie.searchsJyy Dt}}$}}T*  T]]""rcN|}d|jvr't|jjdk(ryg}g}d}t|jjD]o}t |j|t rN |j |j||j}|j|j||znd}qt|dkD } t|dkDrIt|dk(r|j|dn&|jddj|zdzt|dk(r|d} nddj|zd z} |r| r| d z } | Sd| d } | S#t$r$|j|j|YBwxYw) zConvert a Trie into a regular expression pattern Memoized on the hash digest of the trie, which is built incrementally during add(). r!rNr[]z(?:|r?z)?) rkrkeyssorted isinstanceri_patternrsr}r Exceptionr) r&rQrrrzaltccqrrecursecconlyrs rrz Trie._patterns  3t}}'9'9';#<#A  4==--/0 D$---x80"mmDMM$,?NGJJtzz$/'9: X\! r7Q;2w!| 2a5! 3,s23 s8q=VFSXXc]*S0F #  vhb) -!0IIdjj./0sA E77)F$#F$cN|j|j|jSz#Export the Trie to a regex pattern.rrQrsr)s rrz Trie.pattern }}TYY 55rcN|j|j|jSrrr)s rexport_to_regexzTrie.export_to_regexrrr+)rrrrr'rrr|rr6 functools lru_cacherrrrrrrGrGsFS+ ! #))V66rrG PYTORCH_MAPz=Expected each mapping in CUDA_TO_HIP_MAPPINGS to be a Mappingz(?<=\W)(z)(?=\W)z#include "([^"]+)"z#include <([^>]+)>z"#define THC_GENERIC_FILE "([^"]+)"z\.cu\bc x tjjtjj|t |}t } |vr%d| _d| _tj| _ | St tjj|} td5} | jtk(r.d| _d| _tj| _ | cdddS| jd| j!} ddd } tjjtjjt#| }tjj%tjj'|s.j)tjj'|d}t*j-|| } t/j0D]O\}}d|vr4d|vr0d t3j4|z}t3j,||| } >| j7||} Qd#f d }t8j-|d d | } t:j-|d d| } t<j-|d| } |j?dr:| j7dd} | j7dd} t@j-d| } s tC| } |j?drd|vr tE| } tG| } ri| | k(rdtjj'tjj'|k(r%| _d| _tj| _ | S|k7rtIdr t| z} d }tjj%|r)t|d5}|j!| k7}ddd|rT j|dd5}|jK| ddd|| _d| _tj| _ | S|| _d"| _tj| _ | S#1swYnxYw#1swYxYw#1swYlxYw#tL$r~}tOtPjRd|d|jTddtPjV tXjZ | _d!| _tj| _ | cYd}~Sd}~wwxYw)$z< Executes the CUDA -> HIP conversion on the specified file. Nz[ignored, not to be hipified]zutf-8)encodingz#[ignored, input is hipified output]rc2t|jdS)Nr)rr)ms rpt_replzpreprocessor..pt_replXs1771:&&rz/hip/z (?preprocessor..mk_repl..repl..s,UaQZZ-A,Urr)rrorprgrXformatrArdirnamersrrqr-r7r#rrrr$r)rr header_dirheader_filepathheader_dir_to_checkheader_path_to_checkheader_include_dir header_resultheader_rel_pathheader_fout_pathhipified_header_filepathr[rrrrrinclude_current_dirrrrrtempls @rreplz+preprocessor..mk_repl..replqs Aww''*H &'e$Q\\&-A||$5aggajBV$WXX$,U9,U)U! "&&*,''//(*C'+-77??277<"*.AC*.0ggll;KM_.`+/1wwrww||L_ab?c/d,77>>*>?)+O+]+](||M"''//^v_CBZGVXb3d%eff771: rr) rrrrrrrrrrrrs `` rmk_replzpreprocessor..mk_replps9 9 t rz#include "{0}"z#include <{0}>Fz#define THC_GENERIC_FILE "{0}"zCMakeLists.txtrGrHrIrJrDrd PowKernelz[skipped, no changes])rCre.c.cc.cpp.h.hppwz[ok]zFailed to save z with "z ", leaving z unchanged.filez[skipped, no permissions]zkipped, already hipified])T).rorprsrrr-r$r%rrr#rrZreadlineHIPIFY_C_BREADCRUMBseekreadrArqrrzRE_PYTORCH_PREPROCESSORrritemsr5r~rRE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILEr RE_CU_SUFFIXr:r>r?r5writeOSErrorrr3rastrerrorrcsysstderr)rrrrrrrrrrrOfin output_sourceorig_output_source fout_pathr cuda_pathhip_pathrrdo_writefout_oldfoutrrs` ``````` @rrHrH1swwrww||, &2&7&7 # ;K!LML h )#S <<>0 0*.M '#HM *6*;*;M ' ##   #' -=?PQ]_s?t uvI 77>>"''//)4 5277??956',//GM 499; G 8 h 7)#;$bii &::GFF7HmDM*11)XFM G;;x$''0@$(GWM#''0@%(H-XM'++G4T,UWdeM)*%--fe< %--eU; $((?  -mUC  /* 80K.}= %]3M  - / GGOOH %)C C&. #6 &2&7&7 #9!1(M *6*;*;M '  !sU:.snrr) r2rr5rr~rrrtruncate)r search_stringrstrictrcontentss ` rrIrIs x !668 vvRYY}%=$>cBD\^fgH'' ~FH q   s B B  B)ct|d5}|j}|ddk7r|ddk7rd|d}d|d|z}|jd|j||j dddy#1swYyxYw) Nrrrrrr z #include z )r2rrrr)rheaderrrs rrJrJs x !668 !9 r c 1]Fxs+x7 q   s A"A88Bc*|jdd}|S)z||dk(r|dxxdz cc<n(||dk(r ||dz dk7r|ddkDr |dxxdzcc<|ddk(r|ddk(r|j||d  |S|ddk(r)|ddk(r!||d k(r|j||d |dz}|dz }|t|kr|S) a2 Return the list of arguments in the upcoming function parameter closure. Example: string (input): '(blocks, threads, 0, THCState_getCurrentStream(state))' arguments (output): [{'start': 1, 'end': 7}, {'start': 8, 'end': 16}, {'start': 17, 'end': 19}, {'start': 20, 'end': 53}] r)rrrrrrr-rr)rr})rr argumentsclosurescurrent_positionargument_start_poss rrLrL sjI  H)A- S[ ( " #s * SMQ M $ % , SMQ M $ % , SMQ M $ % ,8H18L1MQT1TYabeYfijYj SMQ M C=A (3-1"4   '9BRS T   C=A (3-1"4@P9QUX9X   '9BRS T!1A!5 A+ S[ (. rvcx|jdvry|jdvrytjd)zArgumentParser doesn't support type=bool. Thus, this helper method will convert from possible string types to True / False.)yestruety1T)nofalserr0FzBoolean value expected.)r^argparseArgumentTypeError)rs rrMrM9s: wwy22 4 4(()BCCrproject_directory show_detailedheader_extensions extra_fileshipify_extra_files_onlyc d |dk(rtj}tjj|s t dt j d|s|jd|dz}||k7r@Dcgc]}|j||c} Dcgc]}|j||c} tjj|stj||tttttt  tt| ||| }t|}|D]W}tjj!|s tjj#||}||vsG|j%|Yddlm}|D]{}tjj!|r ||}n&|tjj#||}|j+ fd |j-d D}| t/d }ggd }| s|n|D]}t1|||||| | || t t2j4dzt2j6zt j8|r t;|t<Scc}wcc}w)Nr!z,The project folder specified does not exist.rr_amd)rrrrrr)Pathc3K|]`}|jstt|r8tt|s#t|jr t|bywrQ)is_filerr,r5r)rrprrrs rrzhipify..sP  TH-c$i1 ,=> I s A)AA)r T)rh)rrz-Successfully preprocessed all matching files.r)rogetcwdrprqrrexitrstriprshutilcopytreermaprr6rirLrr}pathlibrextendrglobr4r7r3r`rcrr8r-)rrrrrrrrrrrrrrrincluderXr all_files_setrrrheader_include_dir_pathrrs ` ` ` rrNrNDsk"BIIK 77>>+ , <=     %,v5,,X`aWGOO$57GHaU\]66>>"35EF] 77>>* +)+;<C x01H3}g./G'(8807J:K=QSTI NM  ww}}Q -q1A M !   Q   1   77==+ ,&*+=&> #&*277<<8HJ\+]&^ # "9"?"?"D   (DA 46"ME&=Y;q'(8(IOb(-/?AUW`bo qq '//K Kgll Zadakakle ib]s =J(J-)rrrFFr+r)F)rCrerrrrz.inr)rerrr!r)r rFrTFFFN)Yrrrr5rrrorcuda_to_hip_mappingsrrrcollections.abcrrr enumr rrorrrdictr,HipifyFinalResultrr-__annotations__PYTORCH_TEMPLATE_MAP__all__rr1r2r3r4rboolr5rr6rr7r8r9compilerr:r;r<r= RE_ASSERTr>r;r?r@r@rArBrCrDrErFrirG PYTORCH_TRIErobjectmappingr TypeErrorrsrcdstrrrrrrrrrHrIrJrK RE_INCLUDErLrMrNrrrrs0 656$-g 4 DDl*+T)+&+!+*= s..1 &D%%%9s99T9 K  !"'%* #### #  # # # 08} #L++++& + CI +  +#+(++!%+2 \ F2::>?Wt#LE E BJJ( )   =>:2::hi"F-R T T. T  W\6\6|v !# T#v+ ##G gw 'WXXMMOS C%"**1M1M1O0PPW%XY"**23"**23 bjj!FGrzz)$ zzzz& z CI z  z#z(zz!-zx  RZZ( ) )XDDD T"8$&#"!&$)-1SSSS S  S " SSSSSSSS"S$d*S !Sr