a hη@sUdZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddl m Z ddlmZdd lmZmZdd lmZddlZddlZGd d d eZGd ddZeeefZdZiZeed<dddZgdZGddde Z!ddZ"GdddZ#GdddZ$eedddZ%eee&dd d!Z'd"d#Z(dneeeee&e&eed&d'd(Z)eeeeeee*fe&e&e$e&dd) d*d+Z+d,d-Z,d.d/Z-e.d0Z/d1d2Z0d3d4Z1d5d6Z2d7d8Z3e.d9Z4d:d;Z5e.d<Z6d=d>Z7e.d?Z8d@dAZ9dodBdCZ:dDdEZ;dFdGZdLdMZ?GdNdOdOZ@GdPdQdQZAeAZBiZCeAZDiZEeeeFfedR<iZGe D]ZHeIeHesJeHJD]\ZKZLeLdZMeLddZNejOeNvreDPeKejQeNvreEReKdSreMeGeK<neMeEeK<ejSeNvrejQeNvreBPeKeMeCeK<qqze.eBTZUe.dTeDTdUZVe.dVZWe.dWZXe.dXZYe.dYZZeeeeeee*fe&e&e$e&ed) dZd[Z[dpd\d]Z\d^d_Z]d`daZ^e.dbZ_dcddZ`dedfZadqee&eeeeeee&ee&e&e&e&e e$edkdldmZbdS)ra 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__ INITIALIZEDDONErr\/var/www/html/assistant/venv/lib/python3.9/site-packages/torch/utils/hipify/hipify_python.pyr -sr c@seZdZddZddZdS) HipifyResultcCs||_||_d|_dS)N current_state hipified_pathstatus)selfrrrrr__init__2szHipifyResult.__init__cCsd|jd|jd|jS)NzHipifyResult:: current_state: z, hipified_path : z , status: rrrrr__str__7szHipifyResult.__str__N)r rrrrrrrrr1srz;// !!! 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@rrrrrRszInputError.__str__)r rrrr __classcell__rrrArr Ks r cCst||ddS)Nignore)errors)open)filenamemoderrrr!Vsr!c@s,eZdZdZdZdZdZdZdZdZ dZ d S) r"zzzzzzzzN) r rrHEADERZOKBLUEOKGREENWARNINGZFAILENDCZBOLDZ 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)rrMrrrrpszGeneratedFileCleaner.__init__cCs|Sr>rrrrr __enter__uszGeneratedFileCleaner.__enter__cOs6tj|s |jtj|t|g|Ri|Sr>)ospathexistsrOaddabspathrF)rfnargskwargsrrrrFxs zGeneratedFileCleaner.opencCsxtj|\}}|s$tj|\}}|rF|rFtj|sF|j|ddtj|rV|stt||jtj |dS)NT)exist_ok) rRrSsplitrTmakedirsisdirmkdirrPappendrV)rdnrZparentnrrrr\}s zGeneratedFileCleaner.makedirscCs@|js<|jD]}t|q |jdddD]}t|q,dS)N)rMrOrRunlinkrPrmdir)rtypevalue tracebackfdrrr__exit__s   zGeneratedFileCleaner.__exit__N)F)F) r rr__doc__rrQrFr\rkrrrrr#ns   r#)rSreturncCs|tjdS)N/)replacerRsep)rSrrr _to_unix_pathsrq)rG extensionsrmcstfdd|DS)zendswith).0erGrr z#match_extensions..any)rGrrrrwrr$sr$cstfdd|DS)Nc3s|]}t|VqdSr>)fnmatch)rupatternfilepathrrrxryz_fnmatch..rz)rpatternsrr~r_fnmatchsrrF) root_pathincludesignoresrrout_of_place_onlyis_pytorch_extensionrmccst|}tj|ddD]\}}} tj||} | dkrvd|vrH|dd|vrZ|dd|vrv|d|d| D]x} ttj|| } ttj| | } t | |rzt | |szt | |s| |vrz|st | st | sqz|rt | sqz| VqzqdS)NT)topdown.z.gitbuildZ third_partyzthird_party/nvfuser)rNrRwalkrSrelpathremover_rqjoinrr$r2r5r1)rrrrrrr exact_matchesZ abs_dirpathdirs filenamesZ rel_dirpathrGr rel_filepathrrrr%s8     r%) output_directoryr all_filesheader_include_dirsstatship_clang_launchr clean_ctx show_progressrmc Csttjtj||} ttj| d} | t| <t||||||||| } |rhd| j vrht | d| j | j dd| t| <dS)N)rrZignoredz->T)flush) rRrSrVrrr rrr7rprintr) rrrrrrrrrfin_path hipify_resultresultrrrr&s   r&cCsPdd|dD}tdt|dtd|tdt|dddS) NcSsh|] \}}|qSrr)ruZ cuda_callZ _filepathrrr ryz compute_stats..unsupported_callsz1Total number of unsupported CUDA function calls: rj, z+ Total number of replaced kernel launches: kernel_launches)rlenr)rrrrrr'sr'cCsd}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)ru_rrr ryzadd_dim3..r startr(),end  zdim3()rorange 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    **    r(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)Nrr z::groupinprrrryz'processKernelLaunches..c s|d|ddddddddd}ddi}d}d}d }d }|}t|d ddddD]"}|} |||fvr| d kr||kr|}||d d<|dd7<| dkr|dd8<|ddkr||kr||d d<|}||krZ|s|dvrP||kr|}||dd<|dkr~d|dd<|d|d |d gSqZ||krZ||dd<|d|d |d gSqZdS)Nrrrrrc) kernel_launch kernel_nametemplatez<>rrr r>r<>rr#:rr)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.rrrcrrzno kernel end found)rrr)findr r_)rZ kernel_endZkernel_positionsZ kernel_startrrrfind_kernel_boundsSs  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) Nrrnz//*z/*"\' rxr)rZ in_commentZprev_cZ new_stringrrrr mask_commentsks2  z,processKernelLaunches..mask_commentsrrrrrrcrrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr(rr;ror_)rrrrrZget_kernel_positions output_stringZkernelparamsZ parenthesisrrZend_param_indexZkernel_name_with_templateZcuda_kernel_dim3Znum_klpZ hip_kernelrrrr)s6 ;!   "  r)cCsd}d}|}d\}}|t|kr|||dkrP|durFd}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)rcrcTr)NN)r) input_stringrrZinside_parenthesisZparensrZp_startZp_endrrrr*s$   r*cCst||ddgdS)z%Finds the first balanced parantheses.{}rr*rrrrrr+sr+cCst||ddgdS)z!Finds the first balanced bracket.rrrrrrrrr,sr,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)rro)rrfuncrrrr-s r-z:?:?\b(__syncthreads)\b(\w*\()csh|ddg}tfdd|Dr&Sdv}|dv7}|dv7}|tdu7}|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)ruextrrrrxryz#hip_header_magic..ZhipLaunchKernelGGL __global__Z __shared__Nz#include "hip/hip_runtime.h" )r{RE_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/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 rrrrrrrrrryz'replace_extern_shared..)RE_EXTERN_SHAREDr)rrrrrr/s r/cCstj|rJ|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)rRrSisabsr1r[splitextror)rrdirpathrGrootr orig_filenameZ orig_dirpathrrrr0!s* $       r0cCs>tj|rJ|drdS|dr,dS|dr:dSdS)Ntorch/Fthird_party/nvfuser/tools/autograd/templates/TrRrSr startswithrrrrr1is   r1cCsZtj|rJ|dr,|dr(dSdS|dr:dS|drHdS|drVdSdS)Nzaten/zaten/src/ATen/core/FTrrrrrrrrr2us     r2cCst|rd|vSdS)NsparseFr2lowerrrrrr3s r3cCs<t|r8d|vrdSd|vr8d|vr4dSdSdS)NrTZlinalgZbatchlinearalgebralibblasFrrrrrr4s   r4cCsRtj|rJ|drdStj|}tj|\}}d|vsJ|dvoPd|vS)Nzc10/cudaTZgpur.cuhZcudnn)rRrSrrbasenamer)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>)childrenrrrrrszTrieNode.__init__N)r rrrlrrrrrrsrc@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)rrhashlibmd5_hashdigest_digestrrrrrsz Trie.__init__cCsT|j||j|_|j}|D]}|j|t|j|}q&d|jd<dS)zAdd a word to the Trie. TrN) rupdateencoderrrr setdefaultrrwordnoderrrrrUs  zTrie.addcCs|jS)zReturn the root node of Trie. )rrrrrdumpsz Trie.dumpcCs t|S)z Escape a char for regex. )rescape)rrrrrquotesz Trie.quotecCs6|j}|D] }||jvr$|j|}q dSq d|jvS)zZSearch whether word is present in the Trie. Returns True if yes, else return FalseFr)rrr rrrrs   z Trie.searchc CsH|}d|jvr$t|jdkr$dSg}g}d}t|jD]j}t|j|trz,||j||j}|| ||Wqt y|| |Yq0q>d}q>t|dk } t|dkrt|dkr||dn|dd |dt|dkr|d} ndd |d } |rD| r8| d 7} n d| 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)?) rrkeyssorted isinstancer_patternrr_r Exceptionr) rrrrZaltccqrZrecurseZcconlyrrrrrs6      z Trie._patterncCs||j|jSz#Export the Trie to a regex pattern.rrrrrrrr}sz Trie.patterncCs||j|jSrrrrrrexport_to_regexszTrie.export_to_regexN)r rrrlrrUrrr functools lru_cacherr}rrrrrr6s   +r6 PYTORCH_MAPrz(?<=\W)(z)(?=\W)z#include "([^"]+)"z#include <([^>]+)>z"#define THC_GENERIC_FILE "([^"]+)"z\.cu\bc  stjtj|t|}t} |vrFd| _d| _tj | _ | Sttj |} t ddT} | tkrd| _d| _tj | _ | WdS| d| } Wdn1s0Y| } tjtjt| }tjtj|stj|ddfd d }r8t| } nDt| rPt|| } n,t| rht| } nd d }t|| } d' f dd }t|dd | } t|dd| } t|d| } |dr| dd} | dd} td| } st | } |dr,d|vr,t!| } t"| } rv| | krvtjtj|krv| _d| _tj | _ | S|krt#drt| } d }tj|rt |dd}|| k}Wdn1s0Y|rzVj |ddd}|$| Wdn1s0Y|| _d| _tj | _ | WSt%y}zTt&t'j(d |d!|j)d"d#t'j* t+j,d$| _d%| _tj | _ | WYd}~Sd}~00n|| _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)r"rmrrrpt_repl[szpreprocessor..pt_replcst|d|Sr$)PYTORCH_SPECIAL_MAPgetrr%)r'rrpt_special_repl^sz%preprocessor..pt_special_replcSst|dSr$) CAFFE2_MAPrr%rrrc2_replkszpreprocessor..c2_replTc s$ f dd}|S)Nc s|d}tj||ds4|drJ|dsJ t|dSrtfddDrd}d}rtj}tj tj ||}tj |r|}|}|durD]<}tj |}tj tj ||}tj |r|}|}q|dur|dS|t vr,t |  nz|t vrt |}|jtjkrtj|}tj tj t|} | |_|t |< tj| dur| n||St |j} ttj| dur| n||S|dS)Nr)z ATen/cudazATen/native/cudazATen/native/nested/cudazATen/native/quantized/cudazATen/native/sparse/cudazATen/native/transformers/cudazTHC/rZTHCPc3s|]}|VqdSr>rs)rusrwrrrxryz>preprocessor..mk_repl..repl..r)rrRrSrrformatr0r{dirnamerVrrTrr&rr rrrrq) r&riZ header_dirZheader_filepathZheader_dir_to_checkZheader_path_to_checkheader_include_dirZ header_resultZheader_rel_pathZheader_fout_pathZhipified_header_filepath) rrrrrinclude_current_dirrrrrtemplrwrreplqsd           z+preprocessor..mk_repl..replr)r2r1r3) rrrrrrrrr)r1r2rmk_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)-rRrSrVrrqrrrr rrrrFreadlineHIPIFY_C_BREADCRUMBseekreadr0rTr/r\RE_PYTORCH_PREPROCESSORrr4r2RE_CAFFE2_PREPROCESSORRE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILErtro RE_CU_SUFFIXr)r-r.r$writeOSErrorrr"rKstrerrorrLsysstderr)rrrrrrrrrrrZfinZ output_sourceZorig_output_sourceZ fout_pathr*r,r4Zdo_writeZfout_oldZfoutrvr) rrrrrrrr'rrrr74s   &    <     ,*&r7cst|dj}|}|r>tdt|dfdd|}n ||}|d|||Wdn1sz0YdS)Nr+z\b(z)\bcsSr>r)rreplace_stringrrrryz+file_specific_replacement..r) r!r@rrrror?rGtruncate)rZ search_stringrNstrictricontentsrrMrr8s &   r8cCst|dh}|}|ddkr8|ddkr8d|d}d|d|}|d|||Wdn1sx0YdS) NrLrrrcrrz #include z )r!r@r?rGrO)rheaderrirQrrrr9s    r9cCs|dd}|S)z.csg|]}|qSrrS)rurDrkrrrfry)rrrrrrr)Pathc3sF|]>}|rtt|rtt|st|jrt|VqdSr>)is_filerstrr$name)rurS)rgrrrrrxs   zhipify..rT)rM)rrz-Successfully preprocessed all matching files.r;)rRgetcwdrSrTrrJexitrstripshutilcopytreermaprqr%rNrrr_pathlibrlextendrglobr#r&r"rJrLrKr'r)rerfrrrgrrrrhrrrrrrirrZ all_files_setrirlr0Zheader_include_dir_pathrrr)rgrrrrerr=FsZ               r=)rrrFF)F)F)FrbrcrrrdrFrTFFFN)crlr_r|rrsrJrRrrZcuda_to_hip_mappingsrrtypingrcollections.abcrrr enumr r rr rdictrnZHipifyFinalResultr>r__annotations__ZPYTORCH_TEMPLATE_MAP__all__rr r!r"r#rqboolr$rr%rr&r'r(compilerr)r*r+r,Z RE_ASSERTr-rr.rr/r0r1r2r3r4r5rr6Z CAFFE2_TRIEr+Z PYTORCH_TRIEr"objectr(mappingritemssrcrgdstZ meta_dataZ API_CAFFE2rUZ API_SPECIALr)Z API_PYTORCHrrBrArCrDrErFr7r8r9r:Z RE_INCLUDEr;r<r=rrrrs&          ! -   # &    H  ^            :   .