U ʹj!@sUdZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddl m Z mZmZmZddlmZmZdd lmZGd d d eZGd d d Ze eefZdZiZeed<dddZddddddddddddddd d!d"d#d$d%d&d'd'd(d)d*d+d,d-d.d d d/g!ZGd0ddeZ d1dZ!Gd2ddZ"Gd3ddZ#eee$d4d5dZ%d6d7Z&dgeeeee$e$eed:d;dZ'eeeee eefe$e$e#e$dd< d=dZ(d>dZ)d?dZ*e+d@Z,dAdZ-dBdZ.dCdZ/dDdZ0e+dEZ1dFdZ2e+dGZ3dHd Z4e+dIZ5dJd!Z6dhdKd"Z7dLd#Z8dMd$Z9dNd%Z:dOd&Z;dPd'ZiZ?e=Z@iZAe eeBfedR<iZCe D]ZDeEeDestFeDGD]\ZHZIeIdZJeIddZKejLeKkre@MeHejNeKkreAOeHdSreJeCeH<neJeAeH<ejPeKkrejNeKkre>MeHeJe?eH<qqe+e>QZRe+dTe@QdUZSe+dVZTe+dWZUe+dXZVe+dYZWeeeee eefe$e$e#e$ed< dZd)ZXdid[d*ZYd\d+ZZd]d,Z[e+d^Z\d_d-Z]d`d.Z^djee$eeeeeee$ee$e$e$e$ee#ededfd/Z_dS)ka 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_/var/www/html/bwave/python/venv/lib/python3.8/site-packages/torch/utils/hipify/hipify_python.pyr )sr c@seZdZddZddZdS) HipifyResultcCs||_||_d|_dS)N) current_state hipified_pathstatus)selfrrrrr__init__.szHipifyResult.__init__cCsd|j|j|jS)Nz@HipifyResult:: current_state: {}, hipified_path : {}, status: {})formatrrrrrrr__str__3szHipifyResult.__str__N)rrrrrrrrrr-srz;// !!! 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_fileTrie preprocessorfile_specific_replacementfile_add_headerfix_static_global_kernelsextract_argumentsstr2boolhipifycs$eZdZfddZddZZS)r"cst|||_dSN)superrmessage)rrB __class__rrrKs zInputError.__init__cCs d|jS)Nz Input error: )rBrrrrrOszInputError.__str__)rrrrr __classcell__rrrCrr"Hs cCst||ddS)Nignore)errors)open)filenamemoderrrr#Ssc@s,eZdZdZdZdZdZdZdZdZ dZ d S) r$zzzzzzzzN) rrrHEADERZOKBLUEOKGREENWARNINGZFAILENDCZBOLDZ UNDERLINErrrrr$Xsc@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)rrOrrrrmszGeneratedFileCleaner.__init__cCs|Sr@rrrrr __enter__rszGeneratedFileCleaner.__enter__cOs0tj|s |jtj|t|f||Sr@)ospathexistsrQaddabspathrH)rfnargskwargsrrrrHus zGeneratedFileCleaner.opencCsxtj|\}}|s$tj|\}}|rF|rFtj|sF|j|ddtj|rV|stt||jtj |dS)NT)exist_ok) rTrUsplitrVmakedirsisdirmkdirrRappendrX)rdnr\parentnrrrr^zs zGeneratedFileCleaner.makedirscCs@|js<|jD]}t|q |jdddD]}t|q,dS)N)rOrQrTunlinkrRrmdir)rtypevalue tracebackfdrrr__exit__s   zGeneratedFileCleaner.__exit__N)F)F) rrr__doc__rrSrHr^rmrrrrr%ks   )rI extensionsreturncstfdd|DS)zsz#match_extensions..any)rIrorrurr&scstfdd|DS)Nc3s|]}t|VqdSr@)fnmatch)rspatternfilepathrrrvsz_fnmatch..rw)r|patternsrr{r_fnmatchsr~rF) root_pathincludesignoresroout_of_place_onlyis_pytorch_extensionrpccst|}tj|ddD]\}}} tj||} | dkrvd|krH|dd|krZ|dd|krv|d|d| D]p} tj|| } tj| | } t| |rzt| |szt | |s| |krz|st | st | sqz|rt | sqz| VqzqdS)NT)topdown.z.gitbuildZ third_partyzthird_party/nvfuser) rPrTwalkrUrelpathremoverajoinr~r&r4r7r3)rrrrorrZ exact_matchesZ abs_dirpathdirs filenamesZ rel_dirpathrIr| rel_filepathrrrr's8     ) output_directoryr| all_filesheader_include_dirsstatship_clang_launchr clean_ctx show_progressrpc Csttjtj||} ttj| d} | t| <t||||||||| } |rhd| j krht | d| j | j dd| t| <dS)N)rrZignoredz->T)flush) rTrUrXrrr rr r9rprintr) rr|rrrrrrrfin_path hipify_resultresultrrrr(s$  cCsPdd|dD}tdt|dtd|tdt|dddS) NcSsh|] \}}|qSrr)rsZ 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)scCsd}d}|dddd}ddtdD}d||d<t|D]\}}|d krVq|d krh|d 7}n|d krx|d 8}|d ks|t|d krB|dkrB||d k||d <|d 7}|dkrB|d ||d<qB||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)rs_rrr szadd_dim3..rstartr(),end  zdim3()replacerange 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*s6    **    z([ ]+)(detail?)::[ ]+\\\n[ ]+csVtddfdd}dd}dd}t||}}|D] }||}d |d } |d d | d } |d |d } |d d dkrd nd } |d d || d d } t| | }ttd |ddd dd}d|d dddd|ddddd| d | d}|| |}|d|qD|S)zK Replace the CUDA style Kernel launches with the HIP style kernel launches.cSs|d|ddS)Nrrz::groupZinprrrz'processKernelLaunches..c s|d|ddddddddd}ddi}d}d}d }d }|}t|d ddddD]"}|} |||fkr| d kr||kr|}||d d<|dd7<| dkr|dd8<|ddkr||kr||d d<|}||krZ|s|dkrP||kr|}||dd<|dkr~d|dd<|d|d |d gSqZ||krZ||dd<|d|d |d gSqZdS)Nrrrrre) kernel_launch kernel_nametemplatez<>rrrr>r<>rr:r#r)risalnum) Z in_kernelposrZSTARTZ AT_TEMPLATEZAFTER_TEMPLATEZAT_KERNEL_NAMEricharstringrrgrab_method_and_templatesD        z7processKernelLaunches..grab_method_and_templatecSsdd}g}|d|dkr`|d|}|d|d}|dkrDtd||||||dq|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_boundsMs  z1processKernelLaunches..find_kernel_boundscSsd}d}d}|D]}|dkrf|dkr2|dkr2d}q|dkrH|dkrHd}q|dkr|dkr|dkrd}nX|dkr|d ks~|d krd}n:|dkr|dkr|dkrd}n|dkr|dkr|dkrd}|}|dkr||7}q|d 7}q|S) Nr/z//*z/*"\' rxr)rZ in_commentZprev_cZ new_stringrrrr mask_commentses2  z,processKernelLaunches..mask_commentsrrrrrrerrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr*rr=rra)rrrrrZget_kernel_positions output_stringZkernelparamsZ parenthesisrrZend_param_indexZkernel_name_with_templateZcuda_kernel_dim3Znum_klpZ hip_kernelrrrr+ s> ;!   "  cCsd}d}|}d\}}|t|kr|||dkrP|dkrFd}d}|}q|d7}n0|||dkr|r|d8}|dkr|}||fS|d7}qdS)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_parenthesisZparensrZp_startZp_endrrrr,s$   cCst||ddgdS)z%Finds the first balanced parantheses.{}rr,rrrrrr-scCst||ddgdS)z!Finds the first balanced bracket.rrrrrrrrr.sz\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)rrfuncrrrr/s z:?:?\b(__syncthreads)\b(\w*\()csh|ddg}tfdd|Dr&Sdk}|dk7}|dk7}|tdk 7}|rdd |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)rsextrrrrvsz#hip_header_magic..ZhipLaunchKernelGGL __global__Z __shared__Nz#include "hip/hip_runtime.h" )rxRE_SYNCTHREADSr)rheadersZhasDeviceLogicrrrr0s  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(rrrrrrrrrrrrrrz'replace_extern_shared..)RE_EXTERN_SHAREDr)rrrrrr1 s cCstj|rt|s t|s |Stj|\}}tj|\}}|dkrLd}|}|}|dd}|dd}|dd}|dd}|dd}|d kr|dd}|s||krtj|d}|r||kr|||kr|d }tj|||S) z3 Returns the new name of the hipified file .cu.hipcudaZhipCUDAHIPTHCTHHz caffe2/coreZ_hip) rTrUisabsAssertionErrorr3r]splitextrr)rrdirpathrIrootr orig_filenameZ orig_dirpathrrrr2s* $       cCs>tj|rt|drdS|dr,dS|dr:dSdS)Ntorch/Fthird_party/nvfuser/tools/autograd/templates/TrTrUrr startswithrrrrr3cs   cCsZtj|rt|dr,|dr(dSdS|dr:dS|drHdS|drVdSdS)Nzaten/zaten/src/ATen/core/FTrrrrrrrrr4os     cCst|rd|kSdS)NsparseFr4lowerrrrrr5~s cCs<t|r8d|krdSd|kr8d|kr4dSdSdS)NrTZlinalgZbatchlinearalgebralibblasFrrrrrr6s   cCsRtj|rt|drdStj|}tj|\}}d|ksJ|dkoPd|kS)Nzc10/cudaTZgpur.cuhZcudnn)rTrUrrrbasenamer)rrIrrrrrr7s   c@s@eZdZdZddZddZddZdd Zd d Zd d Z dS)r8zRegex::Trie in Python. Creates 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 i|_dSr@datarrrrrsz Trie.__init__cCs<|j}|D]$}||kr||p i||<||}q d|d<dS)Nrrr)rwordrefrrrrrWs  zTrie.addcCs|jSr@rrrrrdumpsz Trie.dumpcCs t|Sr@)rescape)rrrrrquotesz Trie.quotec Cs<|}d|kr t|dkr dSg}g}d}t|D]d}t||trz&|||}||||Wqtk r|||YqXq8d}q8t|dk }t|dkrt|dkr||dn|dd |dt|dkr|d} ndd |d} |r8|r,| d 7} n d| d } | S) Nrrr[]z(?:|r?z)?) rkeyssorted isinstancedict_patternrar Exceptionr) rZpDatarZaltccqrZrecurseZcconlyrrrrrs6     z Trie._patterncCs||Sr@)rrrrrrrzsz Trie.patternN) rrrrnrrWrrrrzrrrrr8s% PYTORCH_MAPrz(?<=\W)(z)(?=\W)z#include "([^"]+)"z#include <([^>]+)>z"#define THC_GENERIC_FILE "([^"]+)"z\.cu\bc  s|tjtj|t} |kr>d| _d| _tj| _ | Stj |} t ddH} | t krd| _d| _tj| _ | W5QRS| d| } W5QRX| } tjtjt| }tjtj|stj|ddfd d }rt| } nDt| r,t|| } n,t| rDt| } nd d }t|| } d' f dd }t|dd | } t|dd| } t|d| } |dr| dd} | dd} td| } st| } |drd|krt | } t!| } rR| | krRtjtj|krR| _d| _tj| _ | S|krpt"drpt | } d }tj|rt |dd}|| k}W5QRX|r`z@j |ddd}|#| W5QRX|| _d| _tj| _ | WSt$k r\}zPt%t&j'd |d!|j(d"d#t&j) t*j+d$| _d%| _tj| _ | WYSd}~XYnXn|| _d&| _tj| _ | SdS)(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|dSNr)rrmrrrpt_repl-szpreprocessor..pt_replcst|d|Sr)PYTORCH_SPECIAL_MAPgetrr)rrrpt_special_repl0sz%preprocessor..pt_special_replcSst|dSr) CAFFE2_MAPrrrrrc2_repl=szpreprocessor..c2_replTc s$ f dd}|S)Nc s|d}tj|\}|ds8|drN|dsN t|dSrtfddDrd}d}rtj}tj tj ||}tj |r|}|}|dkrD]<}tj |}tj tj ||}tj |r|}|}q|dkr |dS|t kr0t |  nz|t krt |}|jtjkrtj|} tj tj t| } | |_|t |< tj| dk r| n||St |j} tj| dk r| n||S|dS)Nr)z ATen/cudazATen/native/cudazATen/native/nested/cudazATen/native/quantized/cudazATen/native/sparse/cudazATen/native/transformers/cudazTHC/rZTHCPc3s|]}|VqdSr@rq)rssrurrrvRsz>preprocessor..mk_repl..repl..r)rrTrUr]rrr2rxdirnamerXrrVr r(rr rrr) rrkrZ header_dirZheader_filepathZheader_dir_to_checkZheader_path_to_checkheader_include_dirZ header_resultZheader_rel_pathZheader_fout_pathZhipified_header_filepath) rrrrrinclude_current_dirrrrrtemplrurreplCsr          z+preprocessor..mk_repl..replr)r"r!r#) rrrrrrrrr)r!r"rmk_replBs :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),rTrUrXrr rrr rrrrHreadlineHIPIFY_C_BREADCRUMBseekreadr2rVrr^RE_PYTORCH_PREPROCESSORrr6r4RE_CAFFE2_PREPROCESSORRE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILErrr RE_CU_SUFFIXr+r/r0r&writePermissionErrorrr$rMstrerrorrNsysstderr)rr|rrrrrrrrrZfinZ output_sourceZorig_output_sourceZ fout_pathrrr$Zdo_writeZfout_oldZfoutrtr) rrrrrrrrrrrr9s       <     &c stt|d`}|}|r>tdt|dfdd|}n ||}|d|||W5QRXdS)Nr+z\b(z)\bcsSr@r)rreplace_stringrrrrz+file_specific_replacement..r) r#r0rrrrr/r7truncate)r|Z search_stringr>strictrkcontentsrr=rr:s &   c Csrt|d^}|}|ddkr8|ddkr8d|d}d|d|}|d|||W5QRXdS) Nr<rrrerrz #include z )r#r0r/r7r?)r|headerrkrArrrr;s    cCs|dd}|S)z s   rrr%r&r'r(z.inr)rr(r)rT)project_directory show_detailedroheader_extensionsrrr extra_filesrrrrrhipify_extra_files_onlyrrpc sdkrttjs.tdtdsDddkrtfdd|D}fdd| D} tjst t t || ||| d }t |}|D]0}tj |stj|}||kr||qd d lm}|D]}tj |r||}n|tj|}|d D]L}|r*tt||r*tt|| s*t|j|r*|t|q*q|dkrtd d}ggd}| s|n|D]}t||||| | || qttjdtjtjd|rt|tS)Nrz,The project folder specified does not exist.rrZ_amdcsg|]}|qSrrC)rsincluderrUrrr7szhipify..csg|]}|qSrrC)rsrFr[rrr8s)rrrorrr)PathrT)rO)rrz-Successfully preprocessed all matching files.r+) rTgetcwdrUrVrr:exitrstripshutilcopytreerr'rPrrrapathlibr\rglobis_filer~strr&namer%r(r$rLrNr;r)r )rUrVrorWrrrrXrrrrrrYrrZ all_files_setrkr\r Zheader_include_dir_pathrUrr|rr[rr?sn                )rrrFF)F)F)FrRrSrrrTrFrTFFFN)`rnrOryrr`r:rTrrZcuda_to_hip_mappingsrrtypingrrrr collections.abcr r enumr r rreZHipifyFinalResultr.r __annotations__ZPYTORCH_TEMPLATE_MAP__all__rr"r#r$r%boolr&r~r'r(r)r*compilerr+r,r-r.Z RE_ASSERTr/rr0rr1r2r3r4r5r6r7r8Z CAFFE2_TRIErZ PYTORCH_TRIErobjectrmappingr ritemssrcridstZ meta_dataZ API_CAFFE2rWZ API_SPECIALrZ API_PYTORCHrzr2r1r3r4r5r6r9r:r;r<Z RE_INCLUDEr=r>r?rrrrsZ         ! -   # &    H   =            9   .