Vh!UdZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddl m Z ddlmZdd lmZmZdd lmZddlZddlZGd d eZGd dZeeefZdZiZeed< dddZgdZGdde Z!dZ"GddZ#GddZ$dedefdZ%dedede&fd Z'd!Z( d]d"ed#ed$eded%e&d&e&deefd'Z)d(ed)ed*ed+ed,eee*fd-e&d&e&d.e$d/e&ddfd0Z+d1Z,d2Z-ej\d3Z/d4Z0d5Z1d6Z2d7Z3ej\d8Z4d9Z5ej\d:Z6d;Z7ej\d<Z8d=Z9d^d>Z:d?Z;d@ZdCZ?GdDdEZ@GdFdGZAeAZBiZCeAZDiZEeeeFfedH<iZGe D]ZHeIeHesJeHjD]\ZKZLeLdZMeLddZNejeNvr||_t|_g|_yrH)keep_intermediatessetfiles_to_clean dirs_to_clean)rr_s rr zGeneratedFileCleaner.__init__ps"4!erc|SrHrr"s r __enter__zGeneratedFileCleaner.__enter__us rctjj|s8|jj tjj |t |g|i|SrH)ospathexistsraaddabspathrQ)rfnargskwargss rrQzGeneratedFileCleaner.openxsHww~~b!    # #BGGOOB$7 8B((((rctjj|\}}|s"tjj|\}}|r4|r2tjj|s|j |dtjj |r|sNtj ||jjtjj|yy)NT)exist_ok) rfrgsplitrhmakedirsisdirmkdirrbappendrj)rdnroparentns rrqzGeneratedFileCleaner.makedirs}sGGMM"%  f-IFA av 6 MM&4M 0ww}}R  HHRL    % %bggoob&9 :)1rc|jsS|jD]}tj||jdddD]}tj |yy)N)r_rarfunlinkrbrmdir)rtypevalue tracebackfds r__exit__zGeneratedFileCleaner.__exit__sY&&((  !  ''"-   'rNF) rrr__doc__r rdrQrqrrrrr+r+ns5 ) ;rr+rgreturncB|jtjdS)N/)replacerfsep)rgs r _to_unix_pathrs << $$rrR extensionsc,tfd|DS)zz#match_extensions..s8x  #8any)rRrs` rr,r,s 8Z8 88rc,tfd|DS)Nc3JK|]}tj|ywrH)fnmatch)rpatternfilepaths rrz_fnmatch..sJgwx1Js #r)rpatternss` r_fnmatchrs JJ JJr root_pathincludesignoresout_of_place_onlyis_pytorch_extensionc#Kt|}tj|dD]+\}}} tjj ||} | dk(rPd|vr|j dd|vr|j dd|vr"|j d|j d| D]} ttjj|| } ttjj| | } t| |sbt| |rot| |s| |vs|s%t| s t| s|r t| s| .yw)NT)topdown.z.gitbuild third_partyzthird_party/nvfuser)r`rfwalkrgrelpathremovertrjoinrr,r:r=r9)rrrrrr exact_matches abs_dirpathdirs filenames rel_dirpathrRr rel_filepaths rr-r-s(MM+-'')T*J&dIggook9= # ~ F#$ G$$ M* 12! H$RWW\\+x%HIH(k8)LML8,!(G4%h ;x=?X+*<8AST`Aa (1N  sDE EE$1Eoutput_directoryr all_filesheader_include_dirsstatship_clang_launch clean_ctx show_progressc ftjjtjj||} t t j | } | t| <t||||||||| } |r2d| jvr$t| d| j| jd| t| <y)N)rrignoredz->T)flush) rfrgrjrrr rr$r?rprintr) rrrrrrrrrfin_path hipify_resultresults rr.r.swwrww||,TUV=W XY $))% &' 8UCT=U9VWX8Y Z[]s A&cd}d}|jddjdd}tdDcgc]}i}}d||d<t|D]b\}}|dkDrnX|dk(r|dz }n |d k(r|dz}|d k(s|t|dz k(s9|dk(s?||d k7z||d <|dz }|dksX|dz||d<d||dd|dd dz}||dd|dd } ||dd|dd jd dj d } ||dd|dd jd dj d } d| d } d| d } |j| | }| j| | }|j|| z||z}|Scc}w)zBadds dim3() to the second and third arguments in the kernel launchr<<>>rstartr(),end  zdim3()rrange enumeraterstrip) kernel_string cuda_kernelcountclosure_arg_locsindc first_arg_rawsecond_arg_rawfirst_arg_cleansecond_arg_cleanfirst_arg_dim3second_arg_dim3first_arg_raw_dim3second_arg_raw_dim3s rr0r0s EG!))%4<z'processKernelLaunches..s!1syy|nB.Orc|d|ddddddddd}ddi}d}d}d }d }|}t|d ddz ddD]} |} |||fvrJ| d k(r||k(r |}||d d<|dxxdz cc<| dk(r$|dxxdzcc<|ddk(r||k(r ||d d<|}||k7s^ |js |dvr-||k7r |}||dd<|dk(sd|dd<|d|d |d gcS||k(s||dd<|d|d |d gcSy)Nrrrrry) kernel_launch kernel_nametemplatez<>rrrr>r<>r#rr:r)risalnum) in_kernelposrSTART AT_TEMPLATEAFTER_TEMPLATEAT_KERNEL_NAMEricharstrings rgrab_method_and_templatez7processKernelLaunches..grab_method_and_templates(1'9)EBRS%'3"$R0 q  s?+G4q8"bA% aA!9D%--3;!,12J.$K1$K3;$K1$KT{a'V{-B34J0!/$!9$$&&)7P*P/!/45M*51Av67M*73"%]!3s:#oJ^``/67M*73"%]!3s:#oJ^``K% arcd}g}|jd|dk7re|jd|}|jd|dz}|dkr td|j|||||d|jd|dk7re|S)zKFinds the starting and ending points for all kernel launches in the string.rrryrrzno kernel end found)rrr)findr(rt)r kernel_endkernel_positions kernel_starts rfind_kernel_boundsz1processKernelLaunches..find_kernel_boundsTs kk%,2!;;uj9L UL9A=JQ !677  # #l:.4\:.N%P Qkk%,2 rcd}d}d}|D]z}|dk(r,|dk(r|dk(rd}nT|dk(r|dk(rd}nG|dk(rB|dk7r=|dk7r8d}n5|dk(r |d k(s|d k(r&d}n#|dk(r |dk(r|dk(rd}n|dk(r |dk(r|dk7rd}|}|dk(r||z }v|d z }||S) Nrrz//*z/*"\' rxr)r in_commentprev_c new_stringrs r mask_commentsz,processKernelLaunches..mask_commentsls   "AR8# !%J#X&C-!%J#X&D.Vs]!$Jt#9T !#Jt#8# !#Js"8$!#JFRa c! 3 "4rrrrrrryrrrrzhipLaunchKernelGGL(z, 0rr) RE_KERNEL_LAUNCHsublistrr0rrCrrt)rrrrrget_kernel_positions output_stringkernelparams parenthesisrrend_param_indexkernel_name_with_templatecuda_kernel_dim3num_klp hip_kernels` rr1r1s ! !"OQW XF9av 0B 2=3H IJM'4)&1kk#ve}5 VAYw/ a@ vgve}= %ay/25!1$*6!9W+=f_>UV[>\_`>`$a!#M;?'6'?+B+B5#+N+V+VW\^a+bcd*-=a-C-K-K 5AK(50.229'%2Fww 4H !:CB[<[^a>> find_closure_group("(hi)", 0, ["(", ")"]) (0, 3) Fr)ryryTr)NN)r) input_stringrrinside_parenthesisparensrp_startp_ends rr2r2s F CNGU L! !  a (!U*%)"!  # %( */A aKF{~% q L! ! rc"t||ddgS)z%Finds the first balanced parantheses.{}rr2rrs rr3r3 lE#s DDrc"t||ddgS)z!Finds the first balanced bracket.rrrr&r's rr4r4r(rz\bassert[ ]*\(c^|}tD]!}|j|dt|d}#|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)rrfuncs rr5r5sJ!M#\%--$qk>QRV>W=XXY;Z[ \ rz:?:?\b(__syncthreads)\b(\w*\()c|ddg}tfd|DrSdv}|dvz }|dvz }|tjduz }|rd|zS) 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.hc3XK|]!}tjd|d|d#yw)z #include ("z"|)N)research)rextrs rrz#hip_header_magic..s+ W299 C5C53] C Ws'*hipLaunchKernelGGL __global__ __shared__Nz#include "hip/hip_runtime.h" )rRE_SYNCTHREADSr/)rheadershasDeviceLogicrs @rr6r6s!M#$;\s]+)\s+(\w+)\s*\[\s*\]\s*;c8|}tjd|}|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)" c|d|jdxsdd|jdd|jddS) NzHIP_DYNAMIC_SHARED(rrrrrrrrrs rrz'replace_extern_shared..s<)#))A,*<")=Qsyy|nBsyyYZ|n\]^r)RE_EXTERN_SHAREDr)rrs rr7r7s(!M$((^`moM rctjj|rJ|s t|s|Stjj |\}}tjj |\}}|dk(rd}|}|}|j dd}|j dd}|j dd}|j dd}|j dd}|d k7r|j dd}|s%||k(r tjj|d}|r||k(r ||z|k(r|d z}tjj|||zS) z3 Returns the new name of the hipified file .cu.hipcudahipCUDAHIPTHCTHHz caffe2/core_hip)rfrgisabsr9rpsplitextrr)rrdirpathrRrootr0 orig_filename orig_dirpaths rr8r8"s/ ww}}\** *  (E l3GX  *ID#H e|MLoofe,Goofe,GooeU+G << &D << &D-||E5) G|$;'',,w.< 7TCZMt|rd|jvSy)NsparseFr:lowerrPs rr;r;s!|$<--/// rct|r9d|jvryd|jvrd|jvryyy)NrSTlinalgbatchlinearalgebralibblasFrTrPs rr<r<sI|$ |))+ + ++- -*l.@.@.BB rc tjj|rJ|jdrytjj |}tjj |\}}d|vxs|dvxrd|vS)Nzc10/cudaTgpur;.cuhcudnn)rfrgrDrObasenamerE)rrRrr0s rr=r=stww}}\** *z*ww -H WW  h 'FAs X  7!7 VgX>UVrceZdZdZdZy)TrieNodezA Trie node whose children are represented as a directory of char: TrieNode. A special char '' represents end of word ci|_yrH)childrenr"s rr zTrieNode.__init__s  rN)rrrrr rrrr`r`s rr`c^eZdZdZdZdZdZdZdZe jdZ dZ d Z y ) 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.ct|_tjd|_|jj |_y)z,Initialize the trie with an empty root node.F)usedforsecurityN)r`rGhashlibmd5_hashdigest_digestr"s rr z Trie.__init__s/J [[7 zz((* rc>|jj|j|jj|_|j }|D]5}|j j|t|j |}7d|j d<y)zAdd a word to the Trie. TrN) rhupdateencoderirjrGrb setdefaultr`rwordnoders rrizTrie.addsz $++-(zz((* yy 'D MM $ $T8: 6==&D '! brc|jS)zReturn the root node of Trie. )rGr"s rdumpz Trie.dumps yyrc,tj|S)z Escape a char for regex. )r.escape)rrs rquotez Trie.quotesyyrc|j}|D]!}||jvr|j|}!yd|jvS)zZSearch whether word is present in the Trie. Returns True if yes, else return FalseFr)rGrbros rr/z Trie.searchsJyy Dt}}$}}T*  T]]""rcN|}d|jvr't|jjdk(ryg}g}d}t|jjD]o}t |j|t rN |j |j||j}|j|j||znd}qt|dkD } t|dkDrIt|dk(r|j|dn&|jddj|zdzt|dk(r|d} nddj|zd z} |r| r| d z } | Sd| d } | S#t$r$|j|j|YBwxYw) 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)?) rbrkeyssorted isinstancer`_patternrjrtrv Exceptionr) rrGrirqaltccqrrecursecconlyrs rrz Trie._patterns  3t}}'9'9';#<#A  4==--/0 D$---x80"mmDMM$,?NGJJtzz$/'9: X\! r7Q;2w!| 2a5! 3,s23 s8q=VFSXXc]*S0F #  vhb) -!0IIdjj./0sA E77)F$#F$cN|j|j|jSz#Export the Trie to a regex pattern.rrGrjr"s rrz Trie.pattern}}TYY 55rcN|j|j|jSrrr"s rexport_to_regexzTrie.export_to_regexrrN)rrrrr rirsrvr/ functools lru_cacherrrrrrr>r>sFS+ ! #))V66rr> PYTORCH_MAPrz(?<=\W)(z)(?=\W)z#include "([^"]+)"z#include <([^>]+)>z"#define THC_GENERIC_FILE "([^"]+)"z\.cu\bc b tjjtjj|t} |vr%d| _d| _tj| _ | Sttjj|} td5} | jtk(r.d| _d| _tj| _ | cdddS| jd| j!} ddd } tjjtjjt#| }tjj%tjj'|s.j)tjj'|dfd}rt*j-| } n]t/| rt*j-|| } n;t1| rt*j-| } nd }t2j-|| } d#f d }t4j-|d d | } t6j-|d d| } t8j-|d| } |j;dr:| j=dd} | j=dd} t>j-d| } s tA| } |j;drd|vr tC| } tE| } ri| | k(rdtjj'tjj'|k(r%| _d| _tj| _ | S|k7rtGdr t| z} d }tjj%|r)t|d5}|j!| k7}ddd|rT j|dd5}|jI| ddd|| _d| _tj| _ | S|| _d"| _tj| _ | S#1swYmxYw#1swYxYw#1swYlxYw#tJ$r~}tMtNjPd|d|jRddtNjT tVjX | _d!| _tj| _ | cYd}~Sd}~wwxYw)$z< Executes the CUDA -> HIP conversion on the specified file. Nz[ignored, not to be hipified]zutf-8)encodingz#[ignored, input is hipified output]rc2t|jdSNr)rrms rpt_replzpreprocessor..pt_repl[s1771:&&rcZtj|jd|Sr)PYTORCH_SPECIAL_MAPgetr)rrs rpt_special_replz%preprocessor..pt_special_repl^s""&&qwwqz71:>>rc2t|jdSr) CAFFE2_MAPrrs rc2_replzpreprocessor..c2_replks!!''!*--rTc *  f d}|S)Nc N |jd}tjj| |j ds"|j dr;|j ds*j t |jdSrt fd Drnd}d}rtjj}tjjtjj||}tjj|r|}|}|D]}tjj|}tjjtjj||}tjj|s|}|}||jdS|tvrt|   n|tvrt|}|jtj k(rtjj#|}tjjtjjt |} | |_|t|<j tjj#| | |S||St|j$} j tjj#| | |S||S|jdS)Nr)z ATen/cudazATen/native/cudazATen/native/nested/cudazATen/native/quantized/cudazATen/native/sparse/cudazATen/native/transformers/cudazTHC/rATHCPc3@K|]}|jywrHr)rsrRs rrz>preprocessor..mk_repl..repl..s,UaQZZ-A,Urr)rrfrgr^rOformatr8rdirnamerjrrhr$r.rr rrr)rr header_dirheader_filepathheader_dir_to_checkheader_path_to_checkheader_include_dir header_resultheader_rel_pathheader_fout_pathhipified_header_filepathrRrrrrrinclude_current_dirrrrrtempls @rreplz+preprocessor..mk_repl..replqs Aww''*H &'e$Q\\&-A||$5aggajBV$WXX$,U9,U)U! "&&*,''//(*C'+-77??277<"*.AC*.0ggll;KM_.`+/1wwrww||L_ab?c/d,77>>*>?)+O+]+](||BGGOOPhPt4LJT%VWW9H*%VWW771: rr) rrrrrrrrrrrrs `` rmk_replzpreprocessor..mk_replps9 9 t rz#include "{0}"z#include <{0}>Fz#define THC_GENERIC_FILE "{0}"zCMakeLists.txtr?r@rArBr<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)-rfrgrjrr$rrr rrrrrQreadlineHIPIFY_C_BREADCRUMBseekreadr8rhrrqRE_PYTORCH_PREPROCESSORrr<r:RE_CAFFE2_PREPROCESSORRE_QUOTE_HEADERRE_ANGLE_HEADERRE_THC_GENERIC_FILErr RE_CU_SUFFIXr1r5r6r,writePermissionErrorrr*rXstrerrorrZsysstderr)rrrrrrrrrrrfin output_sourceorig_output_source fout_pathrrrdo_writefout_oldfoutrrrs` ``````` @@rr?r?5sdwwrww||, &2&7&7 # ;K!LML h )#S <<>0 0*.M '#HM *6*;*;M ' ##   #' -=?PQ]_s?t uvI 77>>"''//)4 5277??956'? /33G]K < (377WM \ *377OM .266w NM;;x$''0@$(GWM#''0@%(H-XM'++G4T,UWdeM)*%--fe< %--eU; $((?  -mUC  /* 80K.}= %]3M  - / GGOOH %)C C&. #6 &2&7&7 #9!1(M *6*;*;M '  !sU0.snrr) r)rr.rrurrrtruncate)r search_stringrstrictrcontentss ` rr@r@s x !668 vvRYY}%=$>cBD\^fgH'' ~FH q   s B B  B)ct|d5}|j}|ddk7r|ddk7rd|d}d|d|z}|jd|j||j dddy#1swYyxYw) Nrrrryrrz #include z )r)rrrr)rheaderrrs rrArAs x !668 !9 r c 1]Fxs+x7 q   s A"A88Bc*|jdd}|S)z||dk(r|dxxdz cc<n(||dk(r ||dz dk7r|ddkDr |dxxdzcc<|ddk(r|ddk(r|j||d  |S|ddk(r)|ddk(r!||d k(r|j||d |dz}|dz }|t|kr|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)rrt)rr argumentsclosurescurrent_positionargument_start_poss rrCrC sjI  H)A- S[ ( " #s * SMQ M $ % , SMQ M $ % , SMQ M $ % ,8H18L1MQT1TYabeYfijYj SMQ M C=A (3-1"4   '9BRS T   C=A (3-1"4@P9QUX9X   '9BRS T!1A!5 A+ S[ (. rcx|jdvry|jdvrytjd)zArgumentParser doesn't support type=bool. Thus, this helper method will convert from possible string types to True / False.)yestruety1T)nofalserrw0FzBoolean value expected.)rUargparseArgumentTypeError)vs rrDrD;s: wwy22 4 4(()BCCrproject_directory show_detailedheader_extensions extra_fileshipify_extra_files_onlyc d |dk(rtj}tjj|s t dt j d|s|jd|dz}||k7r@Dcgc]}|j||c} Dcgc]}|j||c} tjj|stj||tttttt  tt| ||| }t|}|D]W}tjj!|s tjj#||}||vsG|j%|Yddlm}|D]{}tjj!|r ||}n&|tjj#||}|j+ fd |j-d D}| t/d }ggd }| s|n|D]}t1|||||| | || t t2j4dzt2j6zt j8|r t;|t<Scc}wcc}w)Nrz,The project folder specified does not exist.rr_amd)rrrrrr)Pathc3K|]`}|jstt|r8tt|s#t|jr t|bywrH)is_filerstrr,name)rrgrrrs rrzhipify..sP  TH-c$i1 ,=> I s A)AA)rT)r_)rrz-Successfully preprocessed all matching files.r)rfgetcwdrgrhrrexitrstriprshutilcopytreermaprr-r`rDrrtpathlibrextendrglobr+r.r*rWrZrr/r$)rrrrrrrrrrrrrrrincluderOr all_files_setrrrheader_include_dir_pathrrs ` ` ` rrErEFsk"BIIK 77>>+ , <=     %,v5,,X`aWGOO$57GHaU\]66>>"35EF] 77>>* +)+;<C x01H3}g./G'(8807J:K=QSTI NM  ww}}Q -q1A M !   Q   1   77==+ ,&*+=&> #&*277<<8HJ\+]&^ # "9"?"?"D   (DA 46"ME&=Y;q'(8(IOb(-/?AUW`bo qq '//K Kgll Zadakakle gb]s =J(J-)rrrFFr)F)r;r\rrrrz.inr)r\rrrr)rrFrTFFFN)crrrr.rrrfrrcuda_to_hip_mappingsrrtypingrcollections.abcrr r enumr rrfr rdictrHipifyFinalResultrr$__annotations__PYTORCH_TEMPLATE_MAP__all__rr(r)r*r+rboolr,rr-rr.r/r0compilerr1r2r3r4 RE_ASSERTr5r4r6r9r7r8r9r:r;r<r=r`r> CAFFE2_TRIEr PYTORCH_TRIErobjectrmappingritemssrcr}dst meta_data API_CAFFE2ri API_SPECIALr API_PYTORCHrrrrrrrr?r@rArB RE_INCLUDErCrDrErrrrsz0 65$- 4 DDl*+T)+&+!+*= s..1 &D%%%9s99T9 K  !"'%* )))) )  ) # ) 08} )X++++& + CI +  +#+(++!%+2 \ F2::>?Wt#LE E BJJ( )   =>:2::hi E-P  W\6\6|f  v !# T#v+ ##"G gw '' 'mmo " UAh!"I   y 0   S !$$ 1kooc26N+.#C(#& C  1i6K6KS\6\ OOC !JsO "" $K$?$?$AB$"**1M1M1O0PPW%XY"**23"**23 bjj!FGrzz)$ vvvv& v CI v  v#v(vv!-vp  RZZ( ) +\D T"8$&#"!&$)04RRRR R  R " RRRRRRRR"R,-R !Rr