o ki)!@s"UdZddlZddlZddlZddlZddlZddlZddlZddlm Z ddlm Z ddlm Z ddl m Z ddl mZmZdd lmZddlZddlZd d ZGd d d eZGdddZeeefZdZiZeed< dddZgdZGdddeZ ddZ!GdddZ"GdddZ#dedefdd Z$d!ed"ede%fd#d$Z&d%d&Z' ' ' ' ( (dd)ed*ed+ed"ed,e%d-e%de efd.d/Z(d0ed1ed2ed3ed4eee)fd5e%d-e%d6e#d7e%ddfd8d9Z*dd:d;Z+dZ.d?d@Z/dAdBZ0dCdDZ1dEdFZ2e-dGZ3dHdIZ4e-dJZ5dKdLZ6e-dMZ7dNdOZ8ddPdQZ9de%fdRdSZ:de%fdTdUZ;dVdWZGd\d]d]Z?Gd^d_d_Z@e@ZAiZBeeeCfed`<e D] ZDeEeDesfeFdaeDGD]\ZHZIeAJeHeIeBeH<qjqZe-dbeAKdcZLe-ddZMe-deZNe-dfZOe-dgZP d0ed1ed2ed3ed4eee)fd5e%d-e%d6e#d7e%defdhdiZQdddjdkZRddldmZSdndoZTe-dpZUdqdrZVdsede%fdtduZW ( v w x ' y ' ( ' z ( ( ( dd{ed|e%d"ed}ed0ed3ed*ed~ed,e%d+ed7e%d5e%d-e%de%d6e#dBdef ddZXdS)a 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)EnumcCstjd|tdddS)Nz-hipify version 2.0.0 no longer uses function ) stacklevel)warningswarn FutureWarning)namerj/var/www/addictedbytheproject.nl/epg/venv/lib/python3.10/site-packages/torch/utils/hipify/hipify_python.py _deprecated-src@seZdZdZdZdS) CurrentStaterr N)__name__ __module__ __qualname__ INITIALIZEDDONErrrrr0src@s$eZdZdddZdefddZdS) HipifyResultreturnNcCs||_||_d|_dS)N current_state hipified_pathstatus)selfrrrrr__init__5s zHipifyResult.__init__cCsd|jd|jd|jS)NzHipifyResult:: current_state: z, hipified_path : z , status: rr rrr__str__:szHipifyResult.__str__rN)rrrr!strr#rrrrr4s rz;// !!! 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_argumentsstr2boolrrhipifycs,eZdZdfdd ZdefddZZS)r*rNcst|||_dSN)superr!message)r rJ __class__rrr!Qs  zInputError.__init__cCs d|jS)Nz Input error: )rJr"rrrr#Us zInputError.__str__r$)rrrr!r%r# __classcell__rrrKrr*Nsr*cCst||ddS)Nignore)errors)open)filenamemoderrrr+Yr+c@s,eZdZdZdZdZdZdZdZdZ dZ d S) r,zzzzzzzzN) rrrHEADEROKBLUEOKGREENWARNINGFAILENDCBOLD UNDERLINErrrrr,^sr,c@s@eZdZdZddddZddZd d Zddd d Zd dZdS)r-z+Context Manager to clean up generated filesFrNcCs||_t|_g|_dSrH)keep_intermediatessetfiles_to_clean dirs_to_clean)r r\rrrr!ss zGeneratedFileCleaner.__init__cCs|SrHrr"rrr __enter__xszGeneratedFileCleaner.__enter__cOs6tj|s|jtj|t|g|Ri|SrH)ospathexistsr^addabspathrP)r fnargskwargsrrrrP{s zGeneratedFileCleaner.opencCs|tj|\}}|stj|\}}|r#|r#tj|s#|j|ddtj|r+|sz#match_extensions..any)rQrrrrr.sr.cr)Nc3s|] }t|VqdSrH)fnmatch)rpatternfilepathrrrsz_fnmatch..r)rpatternsrrr_fnmatchsrrF root_pathincludesignoresout_of_place_onlyis_pytorch_extensionc cst|}tj|ddD]T\}}} tj||} | dkrq dS)NT)topdown.z.gitbuild third_partyzthird_party/nvfuser) r]rawalkrbrelpathremovernrjoinrr.) rrrrrr exact_matches abs_dirpathdirs filenames rel_dirpathrQrrrrr/s2    r/output_directoryr all_filesheader_include_dirsstatship_clang_launch clean_ctx show_progressc Csttjtj||} ttj| d} | t| <t||||||||| } |r4d| j vr4t | d| j | j dd| t| <dS)N)rrignoredz->T)flush) rarbrerrrrr&rArprintr) rrrrrrrrrfin_path hipify_resultresultrrrr0s    r0cCsPdd|dD}tdt|dtd|tdt|dddS) NcSsh|]\}}|qSrr)r cuda_call _filepathrrr sz compute_stats..unsupported_callsz1Total number of unsupported CUDA function calls: ry, z+ Total number of replaced kernel launches: kernel_launches)rlenr)rrrrrr1sr1cCsd}d}|dddd}ddtdD}d||d<t|D]E\}}|d kr+n<|d kr4|d 7}n|d kr<|d 8}|d ksH|t|d krf|dkrf||d k||d <|d 7}|dkrf|d ||d<q!||dd|dd d }||d d|d d }||dd|dd ddd} ||d d|d d ddd} d| d } d| d } || | } || | }|||| |}|S)zBadds dim3() to the second and third arguments in the kernel launchr<<>>cSsg|]}iqSrr)r_rrr szadd_dim3..r startr(),end  zdim3()r~range enumeraterstrip) kernel_string cuda_kernelcountclosurearg_locsindc first_arg_rawsecond_arg_rawfirst_arg_cleansecond_arg_cleanfirst_arg_dim3second_arg_dim3first_arg_raw_dim3second_arg_raw_dim3rrrr2s8    **    r2z([ ]+)(detail?)::[ ]+\\\n[ ]+csTtddfdd}dd}dd}t||}}|D]}||}d |d } |d d | d } |d |d } |d d dkrPd nd } |d d || d d } t| | }ttd |ddd dd}d|d dddd|ddddd| d | d}|| |}|d|q"|S)zK Replace the CUDA style Kernel launches with the HIP style kernel launches.cSs|d|ddS)Nrr z::groupinprrrsz'processKernelLaunches..c sz|d|ddddddddd}ddi}d}d}d }d }|}t|d ddddD]}|} |||fvro| d krQ||krI|}||d d<|dd7<| dkro|dd8<|ddkro||kro||d d<|}||kr|s|dvr||kr|}||dd<|dkrd|dd<|d|d |d gSq-||kr||dd<|d|d |d gSq-dS)Nrrrrrr) kernel_launch kernel_nametemplatez<>rrr r>r<>r#rr:r)risalnum) in_kernelposrSTART AT_TEMPLATEAFTER_TEMPLATEAT_KERNEL_NAMEricharstringrrgrab_method_and_templatesJ      z7processKernelLaunches..grab_method_and_templatecSsrd}g}|d|dkr7|d|}|d|d}|dkr"td||||||d|d|dks |S)zKFinds the starting and ending points for all kernel launches in the string.rrrrrrzno kernel end found)rrr)findr*rn)r kernel_endkernel_positions kernel_startrrrfind_kernel_boundsQs   z1processKernelLaunches..find_kernel_boundscSsd}d}d}|D]f}|dkr3|dkr|dkrd}nF|dkr$|dkr$d}n;|dkr2|dkr2|dkr2d}n,|dkrB|d ks?|d krAd}n|dkrQ|dkrP|dkrPd}n|dkr_|dkr_|dkr_d}|}|dkrj||7}q|d 7}q|S) Nrr}z//*z/*"\' rxr)r in_commentprev_c new_stringrrrr mask_commentsis8  z,processKernelLaunches..mask_commentsrrrrrrrrrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr2rrEr~rn)rrrrrget_kernel_positions output_stringkernelparams parenthesisrrend_param_indexkernel_name_with_templatecuda_kernel_dim3num_klp hip_kernelrrrr3s4 ;!  "  r3cCsd}d}|}d\}}|t|krJ|||dkr(|dur#d}d}|}n|d7}n|||dkr@|r@|d8}|dkr@|}||fS|d7}|t|ksdS)aGeneralization for finding a balancing closure group if group = ["(", ")"], then finds the first balanced parentheses. if group = ["{", "}"], then finds the first balanced bracket. Given an input string, a starting position in the input string, and the group type, find_closure_group returns the positions of group[0] and group[1] as a tuple. Example: >>> find_closure_group("(hi)", 0, ["(", ")"]) (0, 3) Fr)rrrrTr)NN)r) input_stringrrinside_parenthesisparensrp_startp_endrrrr4s&   r4cCt||ddgdS)z%Finds the first balanced parentheses.{}rr4r rrrrr5r5cCr)z!Finds the first balanced bracket.rrrrrrrrr6rr6z\bassert[ ]*\(cCs.|}tD]}||dt|d}q|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 rfuncrrrr7s r7z:?:?\b(__syncthreads)\b(\w*\()csh|ddg}tfdd|DrSdv}|dv7}|dv7}|tdu7}|r2d |S) 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.hc3s*|]}td|d|dVqdS)z #include ("z"|)N)research)rextrrrrs(z#hip_header_magic..hipLaunchKernelGGL __global__ __shared__Nz#include "hip/hip_runtime.h" )rRE_SYNCTHREADSr)r headershasDeviceLogicrrrr8s  r8zGextern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;cCs|}tdd|}|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)" cSs.d|dpdd|dd|ddS) NzHIP_DYNAMIC_SHARED(rrrr rrrrrrrrrs.z'replace_extern_shared..)RE_EXTERN_SHAREDr)r rrrrr9s  r9cCstj|r td|st|s|Stj|\}}tj|\}}|dkr(d}|}|}|dd}|dd}|dd }|dd}|dd}|d krT|dd }|sa||kratj|d}|rq||krq|||krq|d }tj|||S) z3 Returns the new name of the hipified file $rel_filepath must be a relative path.cu.hipcudahipCUDAHIPTHCTHHz caffe2/core_hip) rarbisabsAssertionErrorr;rjsplitextr~r) rel_filepathrdirpathrQrootr orig_filename orig_dirpathrrrr:"s,  $       r:cCsBtj|r td|drdS|drdS|drdSdS)Nr#torch/Fthird_party/nvfuser/tools/autograd/templates/T)rarbr-r. startswithr0rrrr;ks    r;cCstdtj|rtd|dr|drdSdS|dr#dS|dr*dS|d r1dS|d r8dS|d r?dSdS) Nr<r#zaten/zaten/src/ATen/core/FTr5r6zthird_party/fbgemm/zthird_party/mslk/r7)rrarbr-r.r8r9rrrr<xs$        r<cCs tdt|rd|vSdS)Nr=sparseFrr<lowerr9rrrr=s r=cCsDtdt|r d|vrdSd|vr d|vrdSdSdS)Nr>r:TlinalgbatchlinearalgebralibblasFr;r9rrrr>s   r>cCs^tdtj|rtd|drdStj|}tj|\}}d|vs+|dvo.d|vS)Nr?r#zc10/cudaTgpur$.cuhcudnn)rrarbr-r.r8basenamer/)r0rQrrrrrr?s   r?c@seZdZdZdddZdS)TrieNodezA Trie node whose children are represented as a directory of char: TrieNode. A special char '' represents end of word rNcCs i|_dSrH)childrenr"rrrr!s zTrieNode.__init__r$)rrrr|r!rrrrrDsrDc@sZeZdZdZdddZdddZdd Zd d Zd d Ze j ddZ ddZ ddZ dS)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.rNcCs&t|_tjdd|_|j|_dS)z,Initialize the trie with an empty root node.F)usedforsecurityN)rDr2hashlibmd5_hashdigest_digestr"rrrr!sz Trie.__init__cCsT|j||j|_|j}|D]}|j|t|j|}qd|jd<dS)zAdd a word to the Trie. TrN) rIupdateencoderJrKr2rE setdefaultrDr wordnoderrrrrds  zTrie.addcCs|jS)zReturn the root node of Trie. )r2r"rrrdumpsz Trie.dumpcCs t|S)z Escape a char for regex. )rescape)r rrrrquotes z Trie.quotecCs4|j}|D]}||jvr|j|}qdSd|jvS)zZSearch whether word is present in the Trie. Returns True if yes, else return FalseFr)r2rErOrrrrs    z Trie.searchc CsB|}d|jvrt|jdkrdSg}g}d}t|jD]4}t|j|trQz||j||j}|| ||Wqt yP|| |Yqwd}qt|dk } t|dkr{t|dkro||dn |dd |dt|dkr|d} n dd |d } |r| r| d 7} | Sd| 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)?) rErkeyssorted isinstancerD_patternrKrnrT Exceptionr) r r2rJrQaltccqrrecursecconlyrrrrr\s:      z Trie._patterncC||j|jSz#Export the Trie to a regex pattern.r\r2rKr"rrrr z Trie.patterncCrcrdrer"rrrexport_to_regexrfzTrie.export_to_regexr$)rrrr|r!rdrRrTr functools lru_cacher\rrgrrrrr@s    + r@ 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  stjtj|t|}t} |vr#d| _d| _tj | _ | Sttj |} t dd*} | tkrLd| _d| _tj | _ | WdS| d| } Wdn1s_wY| } tjtjt| }tjtj|stj|dd}t|| } tD]!\}}d |vrd |vrd t|}t||| } q| ||} qd%f d d }t|dd | } t|dd| } t|d| } |dr| dd} | dd} t d| } st!| } |dr d|vr t"| } t#| } r2| | kr2tjtj|kr2| _d| _tj | _ | S|krAt$drAt| } d }tj|rgt |dd}|| k}Wdn 1sbwY|rz+j |ddd }|%| Wdn 1swY|| _d| _tj | _ | WSt&y}z*t't(j)d|d|j*d d!t(j+ t,j-d"| _d#| _tj | _ | WYd}~Sd}~ww|| _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]rcSst|dS)Nr)rjr)mrrrpt_replVrSzpreprocessor..pt_replz/hip/z (?preprocessor..mk_repl..repl..r)rrarbrCr8formatr:rdirnamererrcr&r0rrrrrr) rlrx header_dirheader_filepathheader_dir_to_checkheader_path_to_checkheader_include_dir header_resultheader_rel_pathheader_fout_pathhipified_header_filepath) rrrrrinclude_current_dirrrrrtemplrrreplosj          z+preprocessor..mk_repl..replr)r|r{r} rrrrrrrrr)r{r|rmk_replns :zpreprocessor..mk_replz#include "{0}"z#include <{0}>Fz#define THC_GENERIC_FILE "{0}"zCMakeLists.txtr(r)r*r+r%r@ 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).rarbrerrr&rrrrrrrPreadlineHIPIFY_C_BREADCRUMBseekreadr:rcrqrkRE_PYTORCH_PREPROCESSORrritemsrrSr~RE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILEr RE_CU_SUFFIXr3r7r8r.writeOSErrorrr,rWstrerrorrYsysstderr)rrrrrrrrrrr0fin output_sourceorig_output_source fout_pathrm cuda_pathhip_pathrrdo_writefout_oldfoutrrr~rrA/s      <       &rAcst|d6}|}|rtdt|dfdd|}n||}|d|||WddS1s>wYdS)Nr+z\b(z)\bcsSrHr)rreplace_stringrrrsz+file_specific_replacement..r) r+rrrrSr~rrtruncate)r search_stringrstrictrxcontentsrrrrBs &    "rBcCst|d5}|}|ddkr|ddkrd|d}d|d|}|d|||WddS1s=wYdS) Nrrrrrrrz #include z )r+rrrr)rheaderrxrrrrrCs     "rCcCs|dd}|S)z.crrr)rrNrrrrbr)rrrrrr)Pathc3sH|]}|rtt|rtt|st|jrt|VqdSrH)is_filerr%r.r)rrb)rrrrrr~s    zhipify..rT)r\)rrz-Successfully preprocessed all matching files.r)ragetcwdrbrcrrexitrstripshutilcopytreermaprr/r]r-rrnpathlibrextendrglobr-r0r,rVrYrr1r&)rrrrrrrrrrrrrrrr all_files_setrxrrvheader_include_dir_pathrrr)rrrrrrrGBs\               rG)rrrFFr$r{)FrrrrrrFrTFFFN)Yr|rrrrrrar cuda_to_hip_mappingsrrrcollections.abcrrrenumr rhrGrrrdictr%HipifyFinalResultrr&__annotations__PYTORCH_TEMPLATE_MAP__all__r]r*r+r,r-rboolr.rr/rr0r1r2compilerr3r4r5r6 RE_ASSERTr7rr8r"r9r:r;r<r=r>r?rDr@ PYTORCH_TRIErjobjectmappingr[ TypeErrorrsrcdstrdrgrrrrrrArBrCrD RE_INCLUDErErFrGrrrrsr         " &    #&    I    ^            =  ,