o wZh!@sUdZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddl m Z ddlmZdd lmZmZdd lmZddlZddlZGd d d eZGd ddZeeefZdZiZeed< dddZgdZGddde Z!ddZ"GdddZ#GdddZ$dedefddZ%d ed!ede&fd"d#Z'd$d%Z( & & & ' 'dd(ed)ed*ed!ed+e&d,e&deefd-d.Z)d/ed0ed1ed2ed3eee*fd4e&d,e&d5e$d6e&ddfd7d8Z+d9d:Z,d;d<Z-e.d=Z/d>d?Z0d@dAZ1dBdCZ2dDdEZ3e.dFZ4dGdHZ5e.dIZ6dJdKZ7e.dLZ8dMdNZ9ddOdPZ:dQdRZ;dSdTZdYdZZ?Gd[d\d\Z@Gd]d^d^ZAeAZBiZCeAZDiZEeeeFfed_<iZGe D]UZHeIeHes_JeHJD]E\ZKZLeLdZMeLddZNejOeNvreDPeKejQeNvreEReKd`reMeGeK<neMeEeK<ejSeNvrejQeNvreBPeKeMeCeK<qcqUe.eBTZUe.daeDTdbZVe.dcZWe.ddZXe.deZYe.dfZZ d/ed0ed1ed2ed3eee*fd4e&d,e&d5e$d6e&defdgdhZ[ddidjZ\dkdlZ]dmdnZ^e.doZ_dpdqZ`drdsZa ' t u ` & v & ' & w ' ' ' ddxedye&d!edzed/ed2ed)ed{ed+e&d*ed6e&d4e&d,e&d|e&d5e e$def d}d~ZbdS)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) constants)CUDA_TO_HIP_MAPPINGS)MATH_TRANSPILATIONS)Optional)Iterator)MappingIterable)Enumc@seZdZdZdZdS) CurrentStaterN)__name__ __module__ __qualname__ INITIALIZEDDONErrO/var/www/auris/lib/python3.10/site-packages/torch/utils/hipify/hipify_python.pyr -sr c@seZdZddZddZdS) HipifyResultcCs||_||_d|_dS)N current_state hipified_pathstatus)selfrrrrr__init__2s zHipifyResult.__init__cCsd|jd|jd|jS)NzHipifyResult:: current_state: z, hipified_path : z , status: rrrrr__str__7szHipifyResult.__str__N)r rrrrrrrrr1s rz;// !!! This is a file automatically generated by hipify!!! HIPIFY_FINAL_RESULTZscalar_t)ZDtypeT)! 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_filer5Trie preprocessorfile_specific_replacementfile_add_headerfix_static_global_kernelsextract_argumentsstr2boolr rhipifycs$eZdZfddZddZZS)r cst|||_dSN)superrmessage)rr@ __class__rrrNs  zInputError.__init__cCs d|jS)Nz Input error: )r@rrrrrRs zInputError.__str__)r rrrr __classcell__rrrArr Ks r cCst||ddS)Nignore)errors)open)filenamemoderrrr!Vr!c@s,eZdZdZdZdZdZdZdZdZ dZ d S) r"zzzzzzzzN) r rrHEADERZOKBLUEOKGREENWARNINGFAILENDCZBOLDZ UNDERLINErrrrr"[sr"c@s<eZdZdZdddZddZddZdd d Zd d Zd S)r#z+Context Manager to clean up generated filesFcCs||_t|_g|_dSr>)keep_intermediatessetfiles_to_clean dirs_to_clean)rrOrrrrps zGeneratedFileCleaner.__init__cCs|Sr>rrrrr __enter__uszGeneratedFileCleaner.__enter__cOs6tj|s|jtj|t|g|Ri|Sr>)ospathexistsrQaddabspathrF)rfnargskwargsrrrrFxs zGeneratedFileCleaner.opencCs|tj|\}}|stj|\}}|r#|r#tj|s#|j|ddtj|r+|sendswith).0erGrr z#match_extensions..any)rGrurr{rr$r$crv)Nc3s|] }t|VqdSr>)fnmatch)rypatternfilepathrrr|sz_fnmatch..r~)rpatternsrrr_fnmatchsrrF root_pathincludesignoresout_of_place_onlyis_pytorch_extensionccst|}tj|ddD]o\}}} tj||} | dkr|rwt | swq>| Vq>q dS)NT)topdown.z.gitbuildZ third_partyzthird_party/nvfuser)rPrTwalkrUrelpathremoverartjoinrr$r2r5r1)rrrrurrZ exact_matchesZ abs_dirpathdirs filenamesZ rel_dirpathrGr rel_filepathrrrr%s>     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)rrZignoredz->T)flush) rTrUrXrrr rrr7rprintr) rrrrrrrrrfin_path hipify_resultresultrrrr&s    r&cCsPdd|dD}tdt|dtd|tdt|dddS) NcSsh|]\}}|qSrr)ryZ cuda_callZ _filepathrrr sz compute_stats..unsupported_callsz1Total number of unsupported CUDA function calls: rl, z+ Total number of replaced kernel launches: kernel_launches)rlenr)rrrrrr'sr'cCsd}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)ry_rrr szadd_dim3..r startr(),end  zdim3()rrrange enumeraterstrip) kernel_string cuda_kernelcountclosureZarg_locsindcZ first_arg_rawZsecond_arg_rawZfirst_arg_cleanZsecond_arg_cleanZfirst_arg_dim3Zsecond_arg_dim3Zfirst_arg_raw_dim3Zsecond_arg_raw_dim3rrrr(s8    **    r(z([ ]+)(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)Nrrrrre) kernel_launch kernel_nametemplatez<>rrr r>r<>r#:rrr)risalnum) Z in_kernelposrSTARTZ AT_TEMPLATEZAFTER_TEMPLATEZAT_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.rrrerrzno kernel end found)rrr)findr ra)rZ kernel_endZkernel_positionsZ kernel_startrrrfind_kernel_boundsTs   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) Nrrqz//*z/*"\' rxr)rZ in_commentZprev_cZ new_stringrrrr mask_commentsls8  z,processKernelLaunches..mask_commentsrrrrrrerrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr(rr;rrra)rrrrrZget_kernel_positions output_stringZkernelparamsZ parenthesisrrZend_param_indexZkernel_name_with_templateZcuda_kernel_dim3Znum_klpZ hip_kernelrrrr)s4 ;!  "  r)cCsd}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)rereTr)NN)r) input_stringrrZinside_parenthesisparensrZp_startZp_endrrrr*s&   r*cCt||ddgdS)z%Finds the first balanced parantheses.{}rr*rrrrrr+r+cCr)z!Finds the first balanced bracket.rrrrrrrrr,rr,z\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)rrr)rrfuncrrrr-s r-z:?:?\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)ryextrrrr|s(z#hip_header_magic..ZhipLaunchKernelGGL __global__Z __shared__Nz#include "hip/hip_runtime.h" )rRE_SYNCTHREADSr)rheadersZhasDeviceLogicrrrr.s  r.zGextern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;cCs|}tdd|}|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)" cSs.d|dpdd|dd|ddS) NzHIP_DYNAMIC_SHARED(rrrr rrrrrrrrrs.z'replace_extern_shared..)RE_EXTERN_SHAREDr)rrrrrr/s r/cCstj|rJ|st|s|Stj|\}}tj|\}}|dkr&d}|}|}|dd}|dd}|dd}|dd}|dd}|d krR|dd}|s_||kr_tj|d}|ro||kro|||kro|d }tj|||S) z3 Returns the new name of the hipified file .cu.hipcudaZhipCUDAHIPTHCTHHz caffe2/coreZ_hip)rTrUisabsr1r]splitextrrr)rrdirpathrGrootr orig_filenameZ orig_dirpathrrrr0"s* $       r0cCs>tj|rJ|drdS|drdS|drdSdS)Ntorch/Fthird_party/nvfuser/tools/autograd/templates/TrTrUr startswithrrrrr1js   r1cCsZtj|rJ|dr|drdSdS|drdS|dr$dS|dr+dSdS)Nzaten/zaten/src/ATen/core/FTrrrrrrrrr2vs     r2cCst|r d|vSdS)NsparseFr2lowerrrrrr3s r3cCs<t|rd|vr dSd|vrd|vrdSdSdS)NrTZlinalgZbatchlinearalgebralibblasFrrrrrr4s   r4cCsRtj|rJ|drdStj|}tj|\}}d|vs%|dvo(d|vS)Nzc10/cudaTZgpur.cuhZcudnn)rTrUrrbasenamer)rrGrrrrrr5s   r5c@seZdZdZddZdS)TrieNodezA Trie node whose children are represented as a directory of char: TrieNode. A special char '' represents end of word cCs i|_dSr>)childrenrrrrrs zTrieNode.__init__N)r rrrorrrrrr s r c@sVeZdZdZddZddZddZdd Zd d Ze j d d Z ddZ ddZ dS)r6zCreates 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.cCs&t|_tjdd|_|j|_dS)z,Initialize the trie with an empty root node.F)usedforsecurityN)r rhashlibmd5_hashdigest_digestrrrrrsz Trie.__init__cCsT|j||j|_|j}|D]}|j|t|j|}qd|jd<dS)zAdd a word to the Trie. TrN) rupdateencoderrrr  setdefaultr rwordnoderrrrrWs  zTrie.addcCs|jS)zReturn the root node of Trie. )rrrrrdumpsz Trie.dumpcCs t|S)z Escape a char for regex. )rescape)rrrrrquotes z Trie.quotecCs4|j}|D]}||jvr|j|}qdSd|jvS)zZSearch whether word is present in the Trie. Returns True if yes, else return FalseFr)rr rrrrrs    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)?) r rkeyssorted isinstancer _patternrrar Exceptionr) rrrrZaltccqrrecurseZcconlyrrrrr!s:      z Trie._patterncC||j|jSz#Export the Trie to a regex pattern.r!rrrrrrrz Trie.patterncCr&r'r(rrrrexport_to_regexr)zTrie.export_to_regexN)r rrrorrWrrr functools lru_cacher!rr*rrrrr6s  + r6 PYTORCH_MAPrz(?<=\W)(z)(?=\W)z#include "([^"]+)"z#include <([^>]+)>z"#define THC_GENERIC_FILE "([^"]+)"z\.cu\bc  stjtj|t} |vrd| _d| _tj| _ | St tj |} t dd*} | tkrHd| _d| _tj| _ | WdS| d| } Wdn1s[wY| } tjtjt| }tjtj|stj|ddfd d }rt| } n t| rt|| } nt| rt| } n d d }t|| } d' f dd }t|dd | } t|dd| } t|d| } |dr| dd} | dd} td| } st | } |dr d|vr t!| } t"| } r1| | kr1tjtj|kr1| _d| _tj| _ | S|kr@t#dr@t| } d }tj|rft |dd}|| k}Wdn 1sawY|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]rcSt|dSNr)r-rmrrrpt_repl[rIzpreprocessor..pt_replcst|d|Sr0)PYTORCH_SPECIAL_MAPgetrr1)r3rrpt_special_repl^rz%preprocessor..pt_special_replcSr/r0) CAFFE2_MAPrr1rrrc2_replkrIzpreprocessor..c2_replTc s$ f dd}|S)Nc s|d}tj||ds|dr%|ds% t|dSrtfddDrd}d}rStj}tj tj ||}tj |rS|}|}|durxD]}tj |}tj tj ||}tj |rw|}|}qY|dur|dS|t vrt |  n=|t vrt |}|jtjkrtj|}tj tj t|} | |_|t |< tj| dur| |S||St |j} tj| dur| |S||S|dS)Nr)z ATen/cudazATen/native/cudazATen/native/nested/cudazATen/native/quantized/cudazATen/native/sparse/cudazATen/native/transformers/cudazTHC/rZTHCPc3s|]}|VqdSr>rw)rysr{rrr|r}z>preprocessor..mk_repl..repl..r)rrTrUrrformatr0rdirnamerXrrVrr&rr rrr) r2rk header_dirZheader_filepathZheader_dir_to_checkZheader_path_to_checkheader_include_dirZ header_resultZheader_rel_pathZheader_fout_pathZhipified_header_filepath) rrrrrinclude_current_dirrrrrtemplr{rreplqsn          z+preprocessor..mk_repl..replr)r?r>r@) rrrrrrrrr)r>r?rmk_replps :zpreprocessor..mk_replz#include "{0}"z#include <{0}>Fz#define THC_GENERIC_FILE "{0}"zCMakeLists.txtrrrrrrZ PowKernelz[skipped, no changes])rr.c.cc.cpp.h.hppwz[ok]zFailed to save z with "z ", leaving z unchanged.filez[skipped, no permissions]zkipped, already hipified])T)-rTrUrXrrrrr rrrtrrFreadlineHIPIFY_C_BREADCRUMBseekreadr0rVr;r^RE_PYTORCH_PREPROCESSORrr4r2RE_CAFFE2_PREPROCESSORRE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILErxrr RE_CU_SUFFIXr)r-r.r$writePermissionErrorrr"rLstrerrorrNsysstderr)rrrrrrrrrrrZfinZ output_sourceZorig_output_sourceZ fout_pathr6r8rAZdo_writeZfout_oldZfoutrzr) rrrrrrrr3rrrr75s       <       &r7cst|d6}|}|rtdt|dfdd|}n||}|d|||WddS1s>wYdS)Nr+z\b(z)\bcsSr>r)rreplace_stringrrrsz+file_specific_replacement..r) r!rMrrrrrrLrTtruncate)r search_stringr[strictrkcontentsrrZrr8s &    "r8cCst|d5}|}|ddkr|ddkrd|d}d|d|}|d|||WddS1s=wYdS) NrYrrrerrz #include z )r!rMrLrTr\)rheaderrkr_rrrr9s     "r9cCs|dd}|S)z.crxrra)ryrDrzrrrfr{)rrrurrr)Pathc3sH|]}|rtt|rtt|st|jrt|VqdSr>)is_filerstrr$name)ryrU)rurrrrr|s    zhipify..rT)rO)rrz-Successfully preprocessed all matching files.rH)rTgetcwdrUrVrrWexitrstripshutilcopytreermaprtr%rPrrrapathlibr|extendrglobr#r&r"rKrNrXr'r)rsrtrururrrrvrrrrrrwrrZ all_files_setrkr|r=Zheader_include_dir_pathrrr)rurrrrsrr=Fs\               r=)rrrFFrn)FrprqrrrrrFrTFFFN)crormrrrrWrTrrZcuda_to_hip_mappingsrrtypingrcollections.abcrrr enumr r+r r rdictr~ZHipifyFinalResultrKr__annotations__ZPYTORCH_TEMPLATE_MAP__all__r"r r!r"r#rtboolr$rr%rr&r'r(compilerr)r*r+r,Z RE_ASSERTr-rr.rr/r0r1r2r3r4r5r r6Z CAFFE2_TRIEr7Z PYTORCH_TRIEr-objectr4mappingr itemssrcridstZ meta_dataZ API_CAFFE2rWZ API_SPECIALr5Z API_PYTORCHr*rOrNrPrQrRrSr7r8r9r:Z RE_INCLUDEr;r<r=rrrrs          " ,    #&    H  ^               9  .