a ¢º”hηã@sUdZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddl m Z ddlmZdd lmZmZdd lmZddlZddlZGd d „d eƒZGd d„dƒZeeefZdZiZeed<dddœZgd¢ZGdd„de ƒZ!dd„Z"Gdd„dƒZ#Gdd„dƒZ$eedœdd„Z%eee&dœd 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 .d0¡Z/d1d2„Z0d3d4„Z1d5d6„Z2d7d8„Z3e .d9¡Z4d:d;„Z5e .d<¡Z6d=d>„Z7e .d?¡Z8d@dA„Z9dodBdC„Z:dDdE„Z;dFdG„ZdLdM„Z?GdNdO„dOƒZ@GdPdQ„dQƒZAeAƒZBiZCeAƒZDiZEeeeFfedR<iZGe D]ªZHeIeHeƒsŽJ‚eH J¡D]Š\ZKZLeLdZMeLdd…ZNejOeNvrôeD PeK¡ejQeNvrìeE ReKdS¡rìeMeGeK<neMeEeK<ejSeNvr–ejQeNvr–eB PeK¡eMeCeK<q–qze .eB T¡¡ZUe .dTeD T¡›dU¡ZVe .dV¡ZWe .dW¡ZXe .dX¡ZYe .dY¡ZZeeeeeee*fe&e&e$e&ed)œ dZd[„Z[dpd\d]„Z\d^d_„Z]d`da„Z^e .db¡Z_dcdd„Z`dedf„Zadqee&eeeeeee&ee&e&e&e&e e$edkœdldm„ZbdS)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)ÚMappingÚIterable)ÚEnumc@seZdZdZdZdS)Ú CurrentStateréN)Ú__name__Ú __module__Ú __qualname__Ú INITIALIZEDÚDONE©rrúN/var/www/auris/lib/python3.9/site-packages/torch/utils/hipify/hipify_python.pyr -sr c@seZdZdd„Zdd„ZdS)Ú HipifyResultcCs||_||_d|_dS)NÚ©Ú current_stateÚ hipified_pathÚstatus)ÚselfrrrrrÚ__init__2szHipifyResult.__init__cCsd|j›d|j›d|j›S)NzHipifyResult:: current_state: z, hipified_path : z , status: r©rrrrÚ__str__7szHipifyResult.__str__N)r rrrrrrrrr1srz;// !!! This is a file automatically generated by hipify!!! ÚHIPIFY_FINAL_RESULTZscalar_t)ZDtypeÚT)!Ú InputErrorÚopenfÚbcolorsÚGeneratedFileCleanerÚmatch_extensionsÚmatched_files_iterÚpreprocess_file_and_save_resultÚ compute_statsÚadd_dim3ÚprocessKernelLaunchesÚfind_closure_groupÚfind_bracket_groupÚfind_parentheses_groupÚreplace_math_functionsÚhip_header_magicÚreplace_extern_sharedÚget_hip_file_pathÚis_out_of_placeÚis_pytorch_fileÚis_cusparse_fileÚis_special_fileÚis_caffe2_gpu_filer5ÚTrieÚ preprocessorÚfile_specific_replacementÚfile_add_headerÚfix_static_global_kernelsÚextract_argumentsÚstr2boolr rÚhipifycs$eZdZ‡fdd„Zdd„Z‡ZS)r cstƒ |¡||_dS©N)ÚsuperrÚmessage)rr@©Ú __class__rrrNs zInputError.__init__cCs d|j›S)Nz Input error: )r@rrrrrRszInputError.__str__)r rrrrÚ __classcell__rrrArr Ks r cCst||ddS)NÚignore)Úerrors)Úopen)ÚfilenameÚmoderrrr!Vsr!c@s,eZdZdZdZdZdZdZdZdZ dZ d S) r"zzzzzzzzN) r rrÚHEADERZOKBLUEÚOKGREENÚWARNINGÚFAILÚENDCZBOLDZ UNDERLINErrrrr"[sr"c@s<eZdZdZddd„Zdd„Zdd„Zdd d „Zd d „Zd S)r#z+Context Manager to clean up generated filesFcCs||_tƒ|_g|_dSr>)Úkeep_intermediatesÚsetÚfiles_to_cleanÚ dirs_to_clean)rrNrrrrpszGeneratedFileCleaner.__init__cCs|Sr>rrrrrÚ __enter__uszGeneratedFileCleaner.__enter__cOs6tj |¡s |j tj |¡¡t|g|¢Ri|¤ŽSr>)ÚosÚpathÚexistsrPÚaddÚabspathrF)rÚfnÚargsÚkwargsrrrrFxs zGeneratedFileCleaner.opencCsxtj |¡\}}|s$tj |¡\}}|rF|rFtj |¡sF|j|ddtj |¡rV|stt |¡|j tj  |¡¡dS)NT)Úexist_ok) rSrTÚsplitrUÚmakedirsÚisdirÚmkdirrQÚappendrW)rÚdnr[ÚparentÚnrrrr]}s zGeneratedFileCleaner.makedirscCs@|js<|jD]}t |¡q |jddd…D]}t |¡q,dS)Néÿÿÿÿ)rNrPrSÚunlinkrQÚrmdir)rÚtypeÚvalueÚ tracebackÚfÚdrrrÚ__exit__‡s   zGeneratedFileCleaner.__exit__N)F)F) r rrÚ__doc__rrRrFr]rlrrrrr#ns   r#)rTÚreturncCs| tjd¡S)Nú/)ÚreplacerSÚsep)rTrrrÚ _to_unix_pathsrr)rGÚ extensionsrncst‡fdd„|DƒƒS)z©Úendswith)Ú.0Úe©rGrrÚ ”óz#match_extensions..©Úany)rGrsrrxrr$’sr$cst‡fdd„|DƒƒS)Nc3s|]}t ˆ|¡VqdSr>)Úfnmatch)rvÚpattern©Úfilepathrrry˜rzz_fnmatch..r{)r€ÚpatternsrrrÚ_fnmatch—sr‚rF)Ú root_pathÚincludesÚignoresrsÚout_of_place_onlyÚis_pytorch_extensionrnccsút|ƒ}tj|ddD]Þ\}}} tj ||¡} | dkrvd|vrH| d¡d|vrZ| d¡d|vrv| d¡| d¡| D]x} ttj || ¡ƒ} ttj | | ¡ƒ} t | |ƒrzt | |ƒszt | |ƒsÈ| |vrz|sìt | ƒsÞt | ƒsÞqz|rìt | ƒsìqz| VqzqdS)NT)ÚtopdownÚ.z.gitÚbuildZ third_partyzthird_party/nvfuser)rOrSÚwalkrTÚrelpathÚremover`rrÚjoinr‚r$r2r5r1)rƒr„r…rsr†r‡Z exact_matchesZ abs_dirpathÚdirsÚ filenamesZ rel_dirpathrGr€Ú rel_filepathrrrr%›s8    ÿþýý r%) Úoutput_directoryr€Ú all_filesÚheader_include_dirsÚstatsÚhip_clang_launchr‡Ú clean_ctxÚ show_progressrnc Csttj tj ||¡¡} ttj| d} | t| <t|||||||||ƒ } |rhd| j vrht | d| j | j dd| t| <dS)N)rrZignoredz->T)Úflush) rSrTrWrŽrr rrr7rÚprintr) r’r€r“r”r•r–r‡r—r˜Úfin_pathÚ hipify_resultÚresultrrrr&Çs  ÿ þr&cCsPdd„|dDƒ}tdt|ƒd›ƒtd |¡ƒtdt|dƒd›ƒdS) NcSsh|] \}}|’qSrr)rvZ cuda_callZ _filepathrrrÚ árzz compute_stats..Úunsupported_callsz1Total number of unsupported CUDA function calls: rkú, z+ Total number of replaced kernel launches: Úkernel_launches)ršÚlenrŽ)r•rŸrrrr'àsr'cCs¦d}d}| dd¡ dd¡}dd„tdƒDƒ}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 … dd¡ d¡} ||d d|d d … dd¡ d¡} d| ›d } d| ›d } | | | ¡} | | | ¡}| ||| |¡}|S)zBadds dim3() to the second and third arguments in the kernel launchrú<<>>cSsg|]}i‘qSrr)rvÚ_rrrÚ òrzzadd_dim3..r Ústartrú(ú)ú,ÚendÚ ú zdim3()rpÚrangeÚ enumerater¢Ústrip)Ú kernel_stringÚ cuda_kernelÚcountÚclosureZarg_locsÚindÚcZ first_arg_rawZsecond_arg_rawZfirst_arg_cleanZsecond_arg_cleanZfirst_arg_dim3Zsecond_arg_dim3Zfirst_arg_raw_dim3Zsecond_arg_raw_dim3rrrr(ís6    **    r(z([ ]+)(detail?)::[ ]+\\\n[ ]+csVt dd„ˆ¡‰‡fdd„}dd„}dd„}t||ˆƒƒƒ}ˆ}|D] }||ƒ}ˆ d |d ¡} ˆ|d d | d …} ˆ|d |d …} |d d dkr¢d nd } ˆ|d d || d d …} t| | ƒ}ttd |d dd ¡ dd¡ƒƒ}d|d d… ddd|d¡ dd¡ dd¡ | d | d¡}| | |¡}|d |¡qD|S)zK Replace the CUDA style Kernel launches with the HIP style kernel launches.cSs| d¡›| d¡›dS)Nrr z::©Úgroup©ÚinprrrÚrzz'processKernelLaunches..c s„|d|ddœdddœdddœdœ}ddi}d}d}d }d }|}t|d ddddƒD]"}ˆ|} |||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)Nr§r«©r§r«rd)Ú kernel_launchÚ kernel_nameÚtemplatez<>rrr ér½ú>r¿ú<>ú#ú:r©r¨r¥r¾)r®Úisalnum) Z in_kernelÚposr³ÚSTARTZ AT_TEMPLATEZAFTER_TEMPLATEZAT_KERNEL_NAMErÚiÚchar©ÚstringrrÚgrab_method_and_templatesDý        z7processKernelLaunches..grab_method_and_templatecSsdd}g}| d|¡dkr`| d|¡}| d|¡d}|dkrDtdƒ‚| |||||…dœ¡q|S)zKFinds the starting and ending points for all kernel launches in the string.rr£rdr¤rÀzno kernel end found)r§r«r¸)Úfindr r`)rËZ kernel_endZkernel_positionsZ kernel_startrrrÚfind_kernel_boundsSs  ÿ z1processKernelLaunches..find_kernel_boundscSsâd}d}d}|D]Ì}|dkrf|dkr2|dkr2d}q¾|dkrH|dkrHd}q¾|dkr¾|dkr¾|dkr¾d}nX|dkr„|d ks~|d kr¾d}n:|dkr¢|dkr¾|dkr¾d}n|dkr¾|dkr¾|dkr¾d}|}|dkrÔ||7}q|d 7}q|S) Nrroz//Ú*z/*ú"ú\ú'ú r¬Úxr)rËZ in_commentZprev_cZ new_stringr¶rrrÚ mask_commentsks2  z,processKernelLaunches..mask_commentsr¨r«rr§rrdr¸r£r¤r©zhipLaunchKernelGGL(z, 0ér r¡) ÚRE_KERNEL_LAUNCHÚsubÚlistrÍr(r¢r;rpr`)rËr•rÌrÎrÕZget_kernel_positionsÚ output_stringÚkernelÚparamsZ parenthesisr²r±Zend_param_indexZkernel_name_with_templateZcuda_kernel_dim3Znum_klpZ hip_kernelrrÊrr)s6 ;!   "ÿÿþ þ r)cCsŽd}d}|}d\}}|t|ƒkrŠ|||dkrP|durFd}d}|}q€|d7}n0|||dkr€|r€|d8}|dkr€|}||fS|d7}qdS)aÊGeneralization 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)rdrdTr)NN)r¢)Ú input_stringr§r¸Zinside_parenthesisÚparensrÆZp_startZp_endrrrr*­s$   r*cCst||ddgdS)z%Finds the first balanced parantheses.Ú{Ú}r·©r*©rÝr§rrrr+Ósr+cCst||ddgdS)z!Finds the first balanced bracket.r¨r©r·rárârrrr,Øsr,z\bassert[ ]*\(cCs.|}tD] }| |›dt|›d¡}q|S)a‹FIXME: 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¨)rrp)rÝrÚÚfuncrrrr-às r-z:?:?\b(__syncthreads)\b(\w*\()csh|‰ddg}t‡fdd„|Dƒƒr&ˆSdˆv}|dˆv7}|dˆv7}|t ˆ¡du7}|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(|] }t d|›d|›dˆ¡VqdS)z #include ("z"|)N)ÚreÚsearch)rvÚext©rÚrrryÿrzz#hip_header_magic..ZhipLaunchKernelGGLÚ __global__Z __shared__Nz#include "hip/hip_runtime.h" )r|ÚRE_SYNCTHREADSrå)rÝÚheadersZhasDeviceLogicrrçrr.ós  r.zGextern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;cCs|}t dd„|¡}|S)a€Match 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| d¡pd›d| d¡›d| d¡›dS) NzHIP_DYNAMIC_SHARED(rrr­r r rÀr©r·r¹rrrr»rzz'replace_extern_shared..)ÚRE_EXTERN_SHAREDrØ)rÝrÚrrrr/s ÿr/cCsðtj |¡rJ‚|s t|ƒs |Stj |¡\}}tj |¡\}}|dkrLd}|}|}| dd¡}| dd¡}| dd¡}| dd¡}| dd¡}|d kr¤| dd¡}|s¾||kr¾tj |d¡}|rÞ||krÞ|||krÞ|d }tj |||¡S) z3 Returns the new name of the hipified file ú.cuú.hipZcudaZhipÚCUDAÚHIPÚTHCÚTHHz caffe2/coreZ_hip)rSrTÚisabsr1r\ÚsplitextrprŽ)r‘r‡ÚdirpathrGÚrootræÚ orig_filenameZ orig_dirpathrrrr0!s* $       r0cCs>tj |¡rJ‚| d¡rdS| d¡r,dS| d¡r:dSdS)Nútorch/Fúthird_party/nvfuser/útools/autograd/templates/T©rSrTròÚ startswith©r‘rrrr1is   r1cCsZtj |¡rJ‚| d¡r,| d¡r(dSdS| d¡r:dS| d¡rHdS| d¡rVdSdS)Nzaten/zaten/src/ATen/core/FTr÷rørùrúrürrrr2us     r2cCst|ƒrd| ¡vSdS)NÚsparseF©r2Úlowerrürrrr3„s r3cCs<t|ƒr8d| ¡vrdSd| ¡vr8d| ¡vr4dSdSdS)NrýTZlinalgZbatchlinearalgebralibblasFrþrürrrr4Šs   r4cCsRtj |¡rJ‚| d¡rdStj |¡}tj |¡\}}d|vsJ|dvoPd|vS)Nzc10/cudaTZgpu©rìú.cuhZcudnn)rSrTròrûÚbasenameró)r‘rGr¥rærrrr5”s   r5c@seZdZdZdd„ZdS)ÚTrieNodezA Trie node whose children are represented as a directory of char: TrieNode. A special char '' represents end of word cCs i|_dSr>)Úchildrenrrrrr¡szTrieNode.__init__N)r rrrmrrrrrrœsrc@sVeZdZdZdd„Zdd„Zdd„Zdd „Zd d „Ze j d d „ƒZ dd„Z dd„Z dS)r6z£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&tƒ|_tjdd|_|j ¡|_dS)z,Initialize the trie with an empty root node.F)ÚusedforsecurityN)rrõÚhashlibÚmd5Ú_hashÚdigestÚ_digestrrrrr¨sz Trie.__init__cCsT|j | ¡¡|j ¡|_|j}|D]}|j |tƒ¡|j|}q&d|jd<dS)zAdd a word to the Trie. TrN) rÚupdateÚencoder r rõrÚ setdefaultr©rÚwordÚnoderÉrrrrV®s  zTrie.addcCs|jS)zReturn the root node of Trie. )rõrrrrÚdump¹sz Trie.dumpcCs t |¡S)z Escape a char for regex. )räÚescape)rrÉrrrÚquote½sz 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)rõrrrrrråÁs   z Trie.searchc CsH|}d|jvr$t|j ¡ƒdkr$dSg}g}d}t|j ¡ƒD]j}t|j|tƒr¤z,| |j||j¡}| |  |¡|¡Wq¨t y | |  |¡¡Yq¨0q>d}q>t|ƒdk } t|ƒdkrøt|ƒdkrà| |d¡n| dd  |¡d¡t|ƒdkr|d} ndd  |¡d } |rD| r8| d 7} n d| ›d } | S) zžConvert 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)?) rr¢ÚkeysÚsortedÚ isinstancerÚ_patternr r`rÚ ExceptionrŽ) rrõr rZaltÚccÚqrÉZrecurseZcconlyrrrrrÎs6      z Trie._patterncCs| |j|j¡S©z#Export the Trie to a regex pattern.©rrõr rrrrr~úsz Trie.patterncCs| |j|j¡Srr rrrrÚexport_to_regexþszTrie.export_to_regexN)r rrrmrrVrrråÚ functoolsÚ lru_cacherr~r!rrrrr6¤s   +r6Ú PYTORCH_MAPrz(?<=\W)(z)(?=\W)z#include "([^"]+)"z#include <([^>]+)>z"#define THC_GENERIC_FILE "([^"]+)"z\.cu\bc  sÊtj tj ˆ|¡¡‰t|ƒ}tˆ} |ˆvrFd| _d| _tj | _ | Sttj  |ˆ¡ƒ} t ˆddT} |   ¡tkr˜d| _d| _tj | _ | WdƒS|  d¡|  ¡} Wdƒn1s¾0Y| } tj tj ˆt| ˆƒ¡¡}tj tj |¡¡sˆ tj |¡¡dd„‰‡fd d „}ˆr8t ˆ| ¡} nDt| ƒrPt || ¡} n,t| ƒrht ˆ| ¡} nd d „}t || ¡} d'‡‡‡‡‡‡‡‡‡ f dd„ }t |dd ƒ| ¡} t |ddƒ| ¡} t |dƒ| ¡} | d¡rþ|  dd¡} |  dd¡} t d| ¡} ˆst | ˆ ƒ} | d¡r,d|vr,t!| ƒ} t"| ƒ} ˆrv| | krvtj ˆ¡tj |¡krvˆ| _d| _tj | _ | Sˆ|kr”t#ˆdƒr”t| } d }tj |¡ràt |dd}| ¡| k}Wdƒn1sÖ0Y|r®zVˆj |ddd}| $| ¡Wdƒn1s0Y|| _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| d¡S©Nr)r$r¸©ÚmrrrÚpt_repl[szpreprocessor..pt_replcst | d¡ˆ|ƒ¡Sr&)ÚPYTORCH_SPECIAL_MAPÚgetr¸r')r)rrÚpt_special_repl^sz%preprocessor..pt_special_replcSst| d¡Sr&)Ú CAFFE2_MAPr¸r'rrrÚc2_replkszpreprocessor..c2_replTc s$‡‡‡‡‡‡‡‡‡ ‡ ‡f dd„}|S)Nc sà| d¡}tj |¡‰| d¡s4| d¡rJ| d¡sJˆ  t| d¡ˆƒ¡SˆrÖt‡fdd„ˆDƒƒrÖd}d}ˆrªtj ˆ¡}tj  tj  ||¡¡}tj  |¡rª|}|}|durôˆD]<}tj  ˆ|¡}tj  tj  ||¡¡}tj  |¡r¶|}|}q¶|dur| d¡S|t vr,t ˆ|ˆˆˆ ˆˆˆˆ ƒ nz|t vr¦t |}|jtjkr¦tj |ˆ¡}tj  tj  ˆt|ˆƒ¡¡} | |_|t |<ˆ  tj | durœ| n||¡¡St |j} ˆ  ttj | durÊ| n||¡ƒ¡S| d¡S)Nr)z ATen/cudazATen/native/cudazATen/native/nested/cudazATen/native/quantized/cudazATen/native/sparse/cudazATen/native/transformers/cudazTHC/rðZTHCPc3s|]}| ˆ¡VqdSr>rt)rvÚsrxrrry€rzz>preprocessor..mk_repl..repl..r)r¸rSrTrrûÚformatr0r|ÚdirnamerWrŽrUrr&rr rrŒrrr) r(rjÚ header_dirZheader_filepathZheader_dir_to_checkZheader_path_to_checkÚheader_include_dirZ header_resultZheader_rel_pathZheader_fout_pathZhipified_header_filepath) r“r—r›r”r–Úinclude_current_dirr‡r’r˜r•ÚtemplrxrÚreplqsd  ÿøø       ý ÿÿ ÿz+preprocessor..mk_repl..replr)r5r4r6) r“r—r›r”r–r‡r’r˜r•)r4r5rÚmk_replps :zpreprocessor..mk_replz#include "{0}"z#include <{0}>Fz#define THC_GENERIC_FILE "{0}"zCMakeLists.txtrîrïrðrñrírZ PowKernelz[skipped, no changes])rìrú.cú.ccú.cppú.hú.hppÚwz[ok]zFailed to save z with "z ", leaving z unchanged.©Úfilez[skipped, no permissions]zkipped, already hipified])T)-rSrTrWrŽrrrrrr rrrŒrFÚreadlineÚHIPIFY_C_BREADCRUMBÚseekÚreadr0rUr1r]ÚRE_PYTORCH_PREPROCESSORrØr4r2ÚRE_CAFFE2_PREPROCESSORÚRE_QUOTE_HEADERÚRE_ANGLE_HEADERÚRE_THC_GENERIC_FILErurpÚ RE_CU_SUFFIXr)r-r.r$ÚwriteÚOSErrorršr"rKÚstrerrorrMÚsysÚstderr)r’r€r“r”r•r–r‡r—r˜rœr‘ZfinZ output_sourceZorig_output_sourceZ fout_pathr,r.r7Zdo_writeZfout_oldZfoutrwr) r“r—r›r”r–r‡r’r)r˜r•rr74s    &    <     ÿþý,*&ÿr7csˆt|dƒj}| ¡}|r>t dt |¡›d‡fdd„|¡}n | |ˆ¡}| d¡| |¡| ¡Wdƒn1sz0YdS)Núr+z\b(z)\bcsˆSr>r)rÔ©Úreplace_stringrrr»ñrzz+file_specific_replacement..r) r!rCrärØrrprBrJÚtruncate)r€Z search_stringrQÚstrictrjÚcontentsrrPrr8ís &   r8cCs†t|dƒh}| ¡}|ddkr8|ddkr8d|›d}d|›d|}| d¡| |¡| ¡Wdƒn1sx0YdS) NrOrrÂrdrÁrÐz #include z )r!rCrBrJrR)r€ÚheaderrjrTrrrr9ùs    r9cCs| dd¡}|S)z.csg|]}| ˆˆ¡‘qSrrV)rvrDrnrrr¦frz)r„r…rsr†r‡r)ÚPathc3sF|]>}| ¡rtt|ƒˆƒrtt|ƒˆƒst|jˆƒrt|ƒVqdSr>)Úis_filer‚Ústrr$Úname)rvrT)rjr…r„rrrys   ýzhipify..rÏT)rN)rŸr¡z-Successfully preprocessed all matching files.r>)rSÚgetcwdrTrUršrMÚexitÚrstripÚshutilÚcopytreerÙÚmaprrr%rOròrŽr`ÚpathlibroÚextendÚrglobr#r&r"rJrMrNr'r)rhrirsrjr’r”r„rkr†r…r˜r–r‡rlr—r“Z all_files_setrjror3Zheader_include_dir_pathr•r€r)rjr…r„r’rhrr=FsZ     ý    ÿ      ÿr=)rrrFF)F)F)FrerfrrrgrFrTFFFN)crmrbr}rärvrMrSrrZcuda_to_hip_mappingsrrÚtypingrÚcollections.abcrrr Úenumr r"rr rÚdictrqZHipifyFinalResultrArÚ__annotations__ZPYTORCH_TEMPLATE_MAPÚ__all__rr r!r"r#rrÚboolr$r‚r%rÙr&r'r(Úcompiler×r)r*r+r,Z RE_ASSERTr-rér.rër/r0r1r2r3r4r5rr6Z CAFFE2_TRIEr-Z PYTORCH_TRIEr$Úobjectr*ÚmappingrÚitemsÚsrcrhÚdstZ meta_dataZ API_CAFFE2rVZ API_SPECIALr+Z API_PYTORCHr!rErDrFrGrHrIr7r8r9r:Z RE_INCLUDEr;r<r=rrrrÚs&          !úú - ÷  # &    H  ^           ÷ :   . ñð