o 闦i!@sUdZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddl m Z mZmZmZddlmZmZdd lmZddlZddlZGd d d eZGd d d Ze eefZdZiZeed< dddZgdZ Gddde!Z"ddZ#GdddZ$GdddZ%dedefddZ&ded ede'fd!d"Z(d#d$Z) % % % & &d~d'ed(ed)ed ed*e'd+e'deefd,d-Z*d.ed/ed0ed1ed2e eefd3e'd+e'd4e%d5e'ddfd6d7Z+d8d9Z,d:d;Z-e.d<Z/d=d>Z0d?d@Z1dAdBZ2dCdDZ3e.dEZ4dFdGZ5e.dHZ6dIdJZ7e.dKZ8dLdMZ9ddNdOZ:dPdQZ;dRdSZdXdYZ?GdZd[d[Z@Gd\d]d]ZAeAZBiZCeAZDiZEe eeFfed^<iZGe D]UZHeIeHes_JeHJD]E\ZKZLeLdZMeLddZNejOeNvreDPeKejQeNvreEReKd_reMeGeK<neMeEeK<ejSeNvrejQeNvreBPeKeMeCeK<qcqUe.eBTZUe.d`eDTdaZVe.dbZWe.dcZXe.ddZYe.deZZ d.ed/ed0ed1ed2e eefd3e'd+e'd4e%d5e'defdfdgZ[ddhdiZ\djdkZ]dldmZ^e.dnZ_dodpZ`dqdrZa & s t _ % u % & % v & & & ddwedxe'd edyed.ed1ed(edzed*e'd)ed5e'd3e'd+e'd{e'd4ee%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)DictListIteratorOptional)MappingIterable)Enumc@seZdZdZdZdS) CurrentStaterN)__name__ __module__ __qualname__ INITIALIZEDDONErr^/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/torch/utils/hipify/hipify_python.pyr ,sr c@seZdZddZddZdS) HipifyResultcCs||_||_d|_dS)N current_state hipified_pathstatus)selfrrrrr__init__1s zHipifyResult.__init__cCsd|jd|jd|jS)NzHipifyResult:: current_state: z, hipified_path : z , status: rrrrr__str__6szHipifyResult.__str__N)rrrrrrrrrr0s 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_filer9Trie preprocessorfile_specific_replacementfile_add_headerfix_static_global_kernelsextract_argumentsstr2boolr rhipifycs$eZdZfddZddZZS)r$cst|||_dSN)superrmessage)rrD __class__rrrMs  zInputError.__init__cCs d|jS)Nz Input error: )rDrrrrrQs zInputError.__str__)rrrrr __classcell__rrrErr$Js r$cCst||ddS)Nignore)errors)open)filenamemoderrrr%Ur%c@s,eZdZdZdZdZdZdZdZdZ dZ d S) r&zzzzzzzzN) rrrHEADEROKBLUEOKGREENWARNINGFAILENDCBOLD UNDERLINErrrrr&Zsr&c@s<eZdZdZdddZddZddZdd d Zd d Zd S)r'z+Context Manager to clean up generated filesFcCs||_t|_g|_dSrB)keep_intermediatessetfiles_to_clean dirs_to_clean)rrVrrrros zGeneratedFileCleaner.__init__cCs|SrBrrrrr __enter__tszGeneratedFileCleaner.__enter__cOs6tj|s|jtj|t|g|Ri|SrB)ospathexistsrXaddabspathrJ)rfnargskwargsrrrrJws zGeneratedFileCleaner.opencCs|tj|\}}|stj|\}}|r#|r#tj|s#|j|ddtj|r+|sz#match_extensions..any)rKr|rrrr(r(cr})Nc3s|] }t|VqdSrB)fnmatch)rpatternfilepathrrrsz_fnmatch..r)rpatternsrrr_fnmatchsrrF root_pathincludesignoresout_of_place_onlyis_pytorch_extensionccst|}tj|ddD]o\}}} tj||} | dkr|rwt | swq>| Vq>q dS)NT)topdown.z.gitbuild third_partyzthird_party/nvfuser)rWr[walkr\relpathremoverhr{joinrr(r6r9r5)rrrr|rr exact_matches abs_dirpathdirs filenames rel_dirpathrKr 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)rrignoredz->T)flush) r[r\r_rrr rr r;rprintr) rrrrrrrrrfin_path hipify_resultresultrrrr*s    r*cCsPdd|dD}tdt|dtd|tdt|dddS) NcSsh|]\}}|qSrr)r cuda_call _filepathrrr sz compute_stats..unsupported_callsz1Total number of unsupported CUDA function calls: rs, 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)r_rrr szadd_dim3..rstartr(),end  zdim3()ryrange enumeraterstrip) kernel_string cuda_kernelcountclosurearg_locsindc first_arg_rawsecond_arg_rawfirst_arg_cleansecond_arg_cleanfirst_arg_dim3second_arg_dim3first_arg_raw_dim3second_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)Nrrz::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)Nrrrrrl) kernel_launch kernel_nametemplatez<>rrrr>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.rrrlrrzno kernel end found)rrr)findr$rh)r kernel_endkernel_positions kernel_startrrrfind_kernel_boundsSs   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) Nrrxz//*z/*"\' rxr)r in_commentprev_c new_stringrrrr mask_commentsks8  z,processKernelLaunches..mask_commentsrrrrrrlrrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr,rr?ryrh)rrrrrget_kernel_positions output_stringkernelparams parenthesisrrend_param_indexkernel_name_with_templatecuda_kernel_dim3num_klp 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)rlrlTr)NN)r) input_stringrrinside_parenthesisparensrp_startp_endrrrr.s&   r.cCt||ddgdS)z%Finds the first balanced parantheses.{}rr.r rrrrr/r/cCr)z!Finds the first balanced bracket.rrrrrrrrr0rr0z\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)rry)r rfuncrrrr1s r1z:?:?\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 headershasDeviceLogicrrrr2s  r2zGextern\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(rrrrrrrrrrrrrs.z'replace_extern_shared..)RE_EXTERN_SHAREDr)r rrrrr3s r3cCstj|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.hipcudahipCUDAHIPTHCTHHz caffe2/core_hip)r[r\isabsr5rdsplitextryr)rrdirpathrKrootr orig_filename orig_dirpathrrrr4!s* $       r4cCs>tj|rJ|drdS|drdS|drdSdS)Ntorch/Fthird_party/nvfuser/tools/autograd/templates/Tr[r\r) startswithrrrrr5is   r5cCsZtj|rJ|dr|drdSdS|drdS|dr$dS|dr+dSdS)Nzaten/zaten/src/ATen/core/FTr/r0r1r2r4rrrr6us     r6cCst|r d|vSdS)NsparseFr6lowerr4rrrr7s r7cCs<t|rd|vr dSd|vrd|vrdSdSdS)Nr5TlinalgbatchlinearalgebralibblasFr6r4rrrr8s   r8cCsRtj|rJ|drdStj|}tj|\}}d|vs%|dvo(d|vS)Nzc10/cudaTgpur .cuhcudnn)r[r\r)r3basenamer*)rrKrrrrrr9s   r9c@seZdZdZddZdS)TrieNodezA Trie node whose children are represented as a directory of char: TrieNode. A special char '' represents end of word cCs i|_dSrB)childrenrrrrrs zTrieNode.__init__N)rrrrvrrrrrr?s r?c@sVeZdZdZddZddZddZdd Zd d Ze j d d Z 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.cCs"t|_t|_|j|_dS)z,Initialize the trie with an empty root node.N)r?r,hashlibmd5_hashdigest_digestrrrrrs z Trie.__init__cCsT|j||j|_|j}|D]}|j|t|j|}qd|jd<dS)zAdd a word to the Trie. TrN) rCupdateencoderDrEr,r@ setdefaultr?rwordnoderrrrr^s  zTrie.addcCs|jS)zReturn the root node of Trie. )r,rrrrdumpsz 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)r,r@rIrrrrs    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?_patternrErhrN Exceptionr) rr,rDrKaltccqrrecursecconlyrrrrrVs:      z Trie._patterncC||j|jSz#Export the Trie to a regex pattern.rVr,rErrrrrz Trie.patterncCr]r^r_rrrrexport_to_regexr`zTrie.export_to_regexN)rrrrvrr^rLrNr functools lru_cacherVrrarrrrr:s  + r: 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)rdrmrrrpt_replZrMzpreprocessor..pt_replcst|d|Srg)PYTORCH_SPECIAL_MAPgetrrh)rjrrpt_special_repl]rz%preprocessor..pt_special_replcSrfrg) CAFFE2_MAPrrhrrrc2_repljrMzpreprocessor..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/r&THCPc3s|]}|VqdSrBr~)rsrrrrrz>preprocessor..mk_repl..repl..r)rr[r\r>r3formatr4rdirnamer_rr]r r*rr rrr) rirr header_dirheader_filepathheader_dir_to_checkheader_path_to_checkheader_include_dir header_resultheader_rel_pathheader_fout_pathhipified_header_filepath) rrrrrinclude_current_dirrrrrtemplrrreplpsn          z+preprocessor..mk_repl..replr)r~r}r) rrrrrrrrr)r}r~rmk_replos :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 r<.c.cc.cpp.h.hppwz[ok]zFailed to save z with "z ", leaving z unchanged.filez[skipped, no permissions]zkipped, already hipified])T)-r[r\r_rr rrr rrr{rrJreadlineHIPIFY_C_BREADCRUMBseekreadr4r]rsreRE_PYTORCH_PREPROCESSORrr8r6RE_CAFFE2_PREPROCESSORRE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILErry RE_CU_SUFFIXr-r1r2r(writePermissionErrorrr&rQstrerrorrSsysstderr)rrrrrrrrrrrfin output_sourceorig_output_source fout_pathrmrordo_writefout_oldfoutrr) rrrrrrrrjrrrr;4s       <       &r;cst|d6}|}|rtdt|dfdd|}n||}|d|||WddS1s>wYdS)Nr+z\b(z)\bcsSrBr)rreplace_stringrrrsz+file_specific_replacement..r) r%rrrrMryrrtruncate)r search_stringrstrictrrcontentsrrrr<s &    "r<cCst|d5}|}|ddkr|ddkrd|d}d|d|}|d|||WddS1s=wYdS) Nrrrrlrrz #include z )r%rrrr)rheaderrrrrrrr=s     "r=cCs|dd}|S)zs r>z #include .*\ncCs:g}ddd}|}|d}|t|kr||dkr"|dd7<n:||dkr1|dd8<n+||dkr@|dd7<n||dkr\||ddkr\|ddkr\|dd8<|ddkrs|ddkrs|||d  |S|ddkr|ddkr||d kr|||d |d}|d7}|t|ks|S) ad 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)rrh)rr argumentsclosurescurrent_positionargument_start_posrrrr? s2     ( $ r?cCs*|dvrdS|dvrdStd)zArgumentParser doesn't support type=bool. Thus, this helper method will convert from possible string types to True / False.)yestruety1T)nofalserrrk0FzBoolean value expected.)r7argparseArgumentTypeError)vrrrr@:s   r@r r<rrrrz.inrr<rrrTproject_directory show_detailedheader_extensions extra_fileshipify_extra_files_onlyc sdkrttjstdtds"ddkr:fddDfddDtjsFt t t t t t t t t ||| d }t|}|D]}tj|sutj|}||vr~||qfd d lm}|D]'}tj|r||}n |tj|}|fd d |dDq|durtdd}ggd}| s|n|D]}t||||| | || qttjdtjtjd|rt|tS)Nrz,The project folder specified does not exist.rrx_amdcg|]}|qSrr)rincluderrrrrdzhipify..crrr)rrHrrrrer)rrr|rrr)Pathc3sH|]}|rtt|rtt|st|jrt|VqdSrB)is_filerstrr(name)rr\)rrrrrrs    zhipify..rT)rV)rrz-Successfully preprocessed all matching files.r)r[getcwdr\r]rrexitrstripshutilcopytreermapr{r)rWr)rrhpathlibrextendrglobr'r*r&rPrSrr+r )rrr|rrrrrrrrrrrrr all_files_setrrrrxheader_include_dir_pathrrr)rrrrrrrAEs\               rA)rrrFFru)FrrrrrrFrTFFFN)crvrrrrrr[rrcuda_to_hip_mappingsrrtypingrrrr collections.abcr r enumr rbrAr rrHipifyFinalResultrr __annotations__PYTORCH_TEMPLATE_MAP__all__rWr$r%r&r'r{boolr(rr)r*r+r,compilerr-r.r/r0 RE_ASSERTr1rr2rr3r4r5r6r7r8r9r?r: CAFFE2_TRIErn PYTORCH_TRIErdobjectrkmappingrUitemssrcrpdst meta_data API_CAFFE2r^ API_SPECIALrl API_PYTORCHrarrrrrrr;r<r=r> RE_INCLUDEr?r@rArrrrs        " ,    #&    H  ^               9  .