
    
9j              
         % S SK Jr  S SKrS SKrS SKrS SKrS SKrS SKrS SKrS SK	r	S SK
r
S SKrS SKJr  S SKJr  S SKJrJrJr  S SKJr  S SKJr  S SKJr  S S	KJr  S S
KJr  \R:                  " 5       rSqSq S r!\	R
                  RE                  S5      q#Sr$Sq%\" 5       q&S\'S'   S@S jr( " S S\)5      r* " S S\)5      r+ " S S\)5      r,SAS jr-\R\                  " 5       S 5       r/SBS jr0SBS jr1S r2\R\                  " 5       S 5       r3Sr4\R\                  " 5       S 5       r5\R\                  " SS 9S! 5       r6\R\                  " SS 9SAS" j5       r7S# r8S$ r9S% r:S& r;\;" S'S(5      r<\;" S)S(5      r=S(q>S* r?S+ r@   SCS, jrA   SCS- jrB   SDS. jrCS/ rD0 qES0\'S1'    SESSS2S(SSS(S(S3.S4 jjrF   SFS5 jrG " S6 S7\)5      rH " S8 S95      rIS: rJSAS; jrKS< rLS= rMSqNS> rO   SGS? jrPg)H    )annotationsN)device)function)DiskKernelCacheBackendKernelCacheBackend_hash_hexdigest)driver)runtimenvrtc)_environment)_utilc                     [         b  [         $ [        R                  (       d!  [        R                  " 5       S:  a  Sq [         $ Sq [         $ )aA  Check if the CUDA driver supports function enumeration APIs.

cuModuleGetFunctionCount, cuModuleEnumerateFunctions, and cuFuncGetName
are required.  The first two were introduced in CUDA driver 12.4 (12040)
and cuFuncGetName in 12.3 (12030), so the effective minimum is 12040.

Returns False unconditionally on HIP/ROCm.
i/  TF)_func_enum_supportedr
   is_hipdriverGetVersion     L/home/wildlama/miniconda3/lib/python3.13/site-packages/cupy/cuda/compiler.py_is_function_enum_supportedr   #   sC     '##>>g668EA#    %r   win32)
--device-cz-dcz	-rdc=truez--relocatable-device-code=true_KernelCacheBackend_kernel_cache_backendc                    U q g)zSet the global kernel cache backend.

This is a private API to allow programmatically changing the cache backend.

Args:
    backend: The kernel cache backend instance to use.
N)r   )backends    r   _set_kernel_cache_backendr   @   s
     $r   c                      \ rS rSrSrg)NVCCExceptionL   r   N__name__
__module____qualname____firstlineno____static_attributes__r   r   r   r   r   L       r   r   c                      \ rS rSrSrg)HIPCCExceptionP   r   Nr!   r   r   r   r)   r)   P   r'   r   r)   c                      \ rS rSrSrg)JitifyExceptionT   r   Nr!   r   r   r   r,   r,   T   r'   r   r,   c           	         [         R                  n[        (       a[  [        5       nUbN  U[         R                  -   [         R                  R                  SS5      -   n[        R                  " U5      nXdS'   [        R                  " XU[        R                  S[        (       a  [        R                  OSS9nUb  UR                  U5        U$ ! [        R                   ad  nSR                  UUR                  UR                   UR"                  5      n	US:X  a  [%        U	5      eUS:X  a  ['        U	5      e[)        U	5      eS nAf[*         a-  nS	[-        U5      -   n	[+        U	R                  U5      5      eS nAff = f)
NPATH Tr   )cwdenvstderruniversal_newlinescreationflagszZ`{}` command returns non-zero exit status. 
command: {}
return-code: {}
stdout/stderr: 
{}nvcchipccz>Failed to run `{0}` command. Check PATH environment variable: )osenviron_win32_get_extra_path_for_msvcpathsepgetcopydeepcopy
subprocesscheck_outputSTDOUTCREATE_NO_WINDOWwriteCalledProcessErrorformatcmd
returncodeoutputr   r)   RuntimeErrorOSErrorstr)
rG   r1   r   
log_streamr2   
extra_pathpathlogemsgs
             r   _run_ccrS   X   sR   ,+ jj6 23J%!BJJ.1KKmmC("F%%c$$#:@&:66a	J
 !S!
(( $ F755<<88%	 	 f$$ %%s## +2A cjj)**	+s%   CC E<#AEE<(E77E<c                     [         R                  " S5      n U (       a  g [        5       nU(       a  U$ [        5       nU(       a  U$ g )Ncl.exe)shutilwhich_get_cl_exe_dir_get_cl_exe_dir_fallback)cl_exe
cl_exe_dirs     r   r;   r;      s;    \\(#F "J)+Jr   c                      SS K n U R                  R                  [        R
                  " 5       5      R                  nU HK  n[        R                  R                  US5      n[        R                  R                  U5      (       d  MI  Us  $    [        R                  " SU 35        g ! [         a     g f = f! [         a/  n[        R                  " S[        U5       SU 35         S nAg S nAff = f)Nr   rU   zcl.exe could not be found in z,Failed to find cl.exe with setuptools.msvc: : )setuptools.msvc	ExceptionmsvcEnvironmentInfoplatformmachineVCToolsr8   rO   joinexistswarningswarntype)
setuptoolsvctoolsrO   rZ   rQ   s        r   rX   rX      s    K	 # //11(2B2B2DEMMDWW\\$1Fww~~f%%  	5gY?@   		  K:47)2aSI	K 	KKsA   B* A?B: 
B: B: *
B74B: 6B77B: :
C3%C..C3c                 n    SSK Jn   SSKJn  U" U " SS05      5      nUR	                  5         UR
                  R                  5         [        R                  R                  UR
                  R                  5      $ ! [         a/  n[        R                  " S[        U5       SU 35         S nAg S nAff = f)Nr   )Distribution)	build_extnamecupy_cl_exe_discoverz'Failed to find cl.exe with setuptools: r]   )rj   rm   setuptools.command.build_extrn   setup_shlib_compilershlib_compiler
initializer8   rO   dirnameccr_   rg   rh   ri   )rm   rn   extrQ   s       r   rY   rY      s    
	F+:f.D%EFG  "%%'wws114455 F5d1gYbD	F 	FFs   A8A; ;
B4%B//B4c                 F    [         c  [        R                  " 5       q [         $ N)_nvrtc_versionr   
getVersionr   r   r   _get_nvrtc_versionr|      s    ))+r   c                 &    SSK Jn   U R                  $ )Nr   core)
cupy._corer   CUPY_CACHE_KEYr~   s    r   _get_cupy_cache_keyr      s    r   )3253627287c                 b    [        5       u  pU S:X  a
  US:  a  SnU$ U S:X  a
  US:X  a  SnU$ SnU$ )N      90120121)r|   )majorminornvrtc_max_compute_capabilitys      r   _get_max_compute_capabilityr      sR    %'LE	'+$ (' 2+%1*',$
 (' (-$''r   T)for_each_devicec                     [        5       n [        R                  " 5       R                  nU[        ;   a  U$ [        X[        S9$ )N)key)r   r   Devicecompute_capability_tegra_archsminint)r   archs     r   	_get_archr      s8     $?#@ ==?--D|43??r   c                    U c
  [        5       n [        (       d'  [        U 5      [        [        5       5      ::  a  SU  3S4$ SU  3S4$ )Nz	-arch=sm_cubin-arch=compute_ptx)r   _use_ptxr   r   )r   s    r   _get_arch_for_options_for_nvrtcr      sN     |{HI8:;;4&!7**D6"E))r   c                &    [        S U  5       5      $ )Nc              3  <   #    U  H  o[         ;   d  M  Uv   M     g 7fry   
_rdc_flags.0os     r   	<genexpr>'_is_cudadevrt_needed.<locals>.<genexpr>  s     5'Q*_qq'   	)anyoptionss    r   _is_cudadevrt_neededr     s    5'555r   c                 8   [         b  [         $ SSKJn   U " 5       nUc  [        S5      e[        (       a  US-  nO1US-   n[
        R                  R                  U5      (       d  US-  nOUn[
        R                  R                  U5      (       d  [        S5      eU$ )Nr   )get_cuda_pathzCUDA is not found.z/lib/x64/cudadevrt.libz/lib64/libcudadevrt.az/lib/libcudadevrt.az>Relocatable PTX code is requested, but cudadevrt is not found.)
_cudadevrt	cupy.cudar   rJ   r:   r8   rO   isfile)r   	cudadevrtcudadevrt64s      r   _get_cudadevrt_pathr     s     ( I/00v--	"99ww~~k**..I#I77>>)$$ 	 r   c                &    [        S U  5       5      $ )Nc              3  <   #    U  H  o[         ;  d  M  Uv   M     g 7fry   r   r   s     r   r   %_remove_rdc_option.<locals>.<genexpr>$  s     ;Gq
':Gr   )tupler   s    r   _remove_rdc_optionr   #  s    ;G;;;r   c                    [         R                  R                  U 5      nUb  [        U5      S:X  a  U$  [	        U5      S:H  $ ! [
         a     gf = f)Nr      F)r8   r9   r=   lenr   
ValueError)ro   defaultvals      r   _get_bool_env_variabler   '  sM    
**..
C
{c#h!m3x1} s   A 
AACUPY_COMPILE_WITH_PTXFCUPY_NVRTC_USE_PCHc                   SSK Jn  [        (       d7  SSKJn  UR                  5         UR                  UR                  5       5        SqU nUS-   U -   n  UR                  X5      u  papxXb:X  d   eXU4$ ! [         ac  n	[        [        U	5      XRUS5      n
[        SS5      nU(       a  U
R                  [        R                  5        [        [        U
5      5      U	eS n	A	ff = f)	Nr   jitifyr~   T
r   CUPY_DUMP_CUDA_SOURCE_ON_ERRORF)r   r   #_jitify_header_source_map_populatedr   r   _init_module_add_sources_get_header_source_mapr_   CompileExceptionrL   r   dumpsysr3   r,   )sourcer   cu_pathr   r   
old_sourcero   headersinclude_namesrQ   cexr   s               r   _jitify_prepr   6  s      /.#D779:.2+ Jt^f$F/06f0N-w ??]**  /s1vzGXN%,e5HHSZZ c#h'Q./s   A3 3
C =ACC c                |    U (       a  [         R                  " S[        SS9  g [         R                  " S[        SS9  g )Nzjitify=True is deprecated and its support is staged for removal in CuPy v15.0.
Please try compiling without jitify using the CCCL headers as needed.
Also see https://nvidia.github.io/cccl/python/ for e.g. Thrust/CUB algorithm exposure to Python.   )
stacklevelz}The jitify argument is deprecated and staged for removal in CuPy v15.0. Avoid passing `jitify=False` to silence this warning.)rg   rh   DeprecationWarningr   s    r   _jitify_deprecation_warningr   [  s8    7 1	. 	D 1		.r   c	           	     X  ^ U4S jn	U(       dd  [         R                  " 5        n
[        R                  R	                  X5      n[        US5       nUR                  U 5        S S S 5        S S S 5        OU(       d  SOUnU	" XWUXWU5      $ ! , (       d  f       N.= f! , (       d  f       N+= f)Nc           	       > Ub  US:X  d   eUST 34-  nO*[         R                  (       d  [        T5      u  pvX4-  nOSnU(       a  [        XU5      u  pn	OK[         R                  (       d  US-  nS=p[	        5       u  pU
S:  d  U
S:X  a  US:  a  [
        (       a  US	-  n[        XXX6S
9n UR                  X5      u  pX4$ ! [         a8  n[        SS5      nU(       a  UR                  [        R                  5        e S nAff = f)Nltor   r   )z#--device-as-default-execution-spacer      r   r   )z--pch)name_expressionsmethodr   F)r
   r   r   r   r|   _use_pch_NVRTCProgramcompiler   r   r   r   r3   )r   r   r   r   rM   r   r   arch_optr   r   major_versionminor_versionprogcompiled_objmappingrQ   r   r   s                    r   _compile1_compile_using_nvrtc_no_warning.<locals>._compiles  s    U?"?./11G>tDH{"GF.:/*+Gm >>CC&((G+=+?(M"$"b(]a-?H:%Vg.>O	$(LL$E!L $$   	)0%9Dszz"	s   /C 
D3DDwr0   )tempfileTemporaryDirectoryr8   rO   re   openrD   )r   r   r   filenamer   rM   cache_in_memoryr   r   r   root_dirr   cu_files     `          r   _compile_using_nvrtc_no_warningr   m  s    '%R ((*hggll86Ggs#wf% $ +* #"FW.>0 0 $# +*s#   ,BB
"B

B	B
B)c	                <    Ub  [        U5        [        XX#XEXgU5	      $ ry   )r   r   )	r   r   r   r   r   rM   r   r   r   s	            r   compile_using_nvrtcr     s,    
 #F+*)9) )r   c           	        SSK Jn  U(       d
  [        5       nUS;  a  [        S5      eUS:X  a	  U(       a   eSR	                  US9nU" 5       n	U	R                  5       n
U
R                  U5        [        R                  " 5        nUR                  S5      S   n[        R                  R                  X5      nS	U-  nU< SU< 3n[        US
5       nUR                  U 5        S S S 5        U(       dB  U
R                  SU-  5        U
[        U5      -  n
U
R                  U5         [        XSU5        OU
R/                  5       nUR                  S5        US-   nU
[        USU4-   5      -  n
U
R                  U5         [        XSU5        [1        U5      nUSUSUS-   4-  nU[        U5      -   n
 [        XSU5        US:X  a/  [        US5       nUR3                  5       sS S S 5        sS S S 5        $ US:X  a/  [        US5       nUR3                  5       sS S S 5        sS S S 5        $  U5       e! , (       d  f       GNT= f! [          aP  n[#        [%        U5      XUS5      n['        SS5      nU(       a  UR)                  [*        R,                  5        UeS nAff = f! [          aP  n[#        [%        U5      XUS5      n['        SS5      nU(       a  UR)                  [*        R,                  5        UeS nAff = f! [          a  n[#        [%        U5      SSUS5      nUeS nAff = f! , (       d  f       O= f! , (       d  f       O= fS S S 5        g ! , (       d  f       g = f)Nr   )get_nvcc_path)r   r   z,Invalid code_type %s. Should be cubin or ptxr   z'-gencode=arch=compute_{cc},code=sm_{cc})rv   .z%s.cur   z--%sr6   r   Fz--cubinz.o-oz--device-link.cubinr0   rbr   )r   r   r   r   rF   splitappendr   r   r8   rO   re   r   rD   listrS   r   r   rL   r   r   r   r3   r>   r   read)r   r   r   r   	code_typeseparate_compilationrM   r   arch_str_nvccrG   r   
first_partrO   r   result_pathr   rQ   r   r   cmd_partialobjptx_filebin_files                           r   compile_using_nvccr    s    ({((GHHE'''8??4?HHOE
++-CJJx		$	$	&(^^C(+
ww||H1D.!%y1'37MM&!   $JJv	)*4= CJJwvz: ((*Ky)+C44+-..CJJwvz: )1GdD8ODDGW-Cvz:
 k4(H}} )(w 
'	&z '!k4(H}} )(} 
'	&B $)#5u   ! 	&s1vv'-/ .4e=HHSZZ(		( ! 	&s1vv'-/ .4e=HHSZZ(		" ! &s1vr2wG	
 )( )(} 
'	&	&s   AMH1$AM'I4AM J %M3K= ML)"	M5ML:	M*M1
I 	;M
JAJJM 
K:*AK55K::M=
L&L!!L&&M)
L7	3M:
M	M
M"c                p   US:X  a5  USR                  U5      4-  n[        U 5      n UR                  U5      u  pVO'US:X  a   US-   n[        XUSSS	9nO[        S
U-  5      e[        U[        5      (       d   eSR                  S UR                  5       R                  5        5       5      $ ! [         a8  n[	        SS5      nU(       a  UR                  [        R                  5        e S nAff = f! [         a8  n[	        SS5      nU(       a  UR                  [        R                  5        e S nAff = f)Nr   z-arch=compute_{}r   Fr6   )r   zpreprocess.ptxzpreprocess.cur   )r  Invalid backend %sr   c              3  T   #    U  H  oR                  S 5      (       d  M  Uv   M      g7f)z//N)
startswith)r   xs     r   r   _preprocess.<locals>.<genexpr>-  s       G/a<<3E/s   (	()rF   r   r   r   r   r   r   r3   r  r   
isinstancebytesre   decode
splitlines)	r   r   r   r   r   result_rQ   r   s	            r   _preprocessr    s9   ' 	&--d355V$	W-IFA 
F			 88G'279F -788fe$$$$ 99 G==?--/G G G/   	)0%9Dszz"	   	)0%9Dszz"	s/   B. C3 .
C083C++C03
D5=3D00D5dict_empty_file_preprocess_cacher   )r   extra_sourcer   enable_cooperative_groupsr   rM   r   to_ltoirc                  U(       a   [         R                  (       a  [        S5      eUb  US:w  a  [        e[	        SS5      =(       a    US:H  n
[         R                  (       a  US:X  a  SOSn[        XX#UXgU
5      $ [        XX#UXVUXU	5      $ )Nz+Cooperative groups is not supported in HIP.r   CUPY_CACHE_IN_MEMORYFhiprtcr7   )r
   r   r   NotImplementedErrorr   _compile_with_cache_hip_compile_with_cache_cuda)r   r   r   r  r   r  r   rM   r   r   r   s              r   _compile_module_with_cacher'  4  s    
 !>>=? ? #7(:!!
 	5u= 	w  ~~%0(g&T/; 	; (T%X/ 	/r   c                   Uc
  [        5       nUS-  nU
(       a  US-  nU(       a  US-  n[        SS5      (       a  US-  nSU;   nU	(       a  U(       d  US-  nOU(       a	  U	(       d  S	n	U	(       a  US
:w  a  [        S5      eX![        5       U4[	        U5      -   n[
        R                  US 5      nUc  [        SXU5      nU[
        U'   [        5       nU< SU< SU < SU< S[        5       < 3	nU(       a'  U(       a   USSR                  [        U5      5      -   -  nUR                  S5      n[        U5      U
(       a  SOS-   nU
(       d  [        R                  " 5       nU(       d  U(       + =(       d    UnU(       a  [         R#                  U5      nUb  U
(       a  U$ WR#                  U5        U(       af  UR%                  ['        U5      5        UR(                  (       d9  UR*                  (       a'  [,        R.                  " UR*                  5        SUl        OU$ U$ O US
:X  a  U(       a  SOUS-   n[1        XUUUXxX(       a  SOS 5	      u  nn[3        U5      (       aZ  U
(       dS  [        R4                  " 5       nUR7                  US5        [9        5       nUR;                  U5        UR=                  5       nOUnU
(       d  WR?                  U5        O>US:X  a*  U
(       a  [@        e[3        U5      n[C        XUUS-   SUUS9nO[        SU-  5      eU(       d  [         RE                  UUU 5        O U
(       a  U$ WR#                  U5        U$ )N)z	-ftz=true)z-dlto)r   CUPY_CUDA_COMPILE_WITH_DEBUGF)z--device-debugz--generate-line-info-DCUPY_USE_JITIFY)r*  Tr   zjitify only works with NVRTCr0    ,utf-8z.ltoirr   r   .cur   zcupy.ptxr6   r   )r  r  rM   r  )#r   r   r   r|   r   r  r=   r  r   r   re   sortedencoder   r   Moduler   load_enumerate_and_build_mappingr   r   ptrr	   moduleUnloadr   r   	LinkStateadd_ptr_datar   add_ptr_filecomplete_set_mappingr$  r  save)r   r   r   r  r   r  r   rM   r   r   r   is_jitify_requestedr2   basecan_enumkey_srcro   mod	use_cacher   cu_namer   r   lsr   rdcs                             r   r&  r&  S  s    |{ ~G: ?" <eDD==.'9) 	))	V'W$788 -/9,T23C'++C6D|2wg6,0$S)*,H 	T6<)<)>@GH
 	3&)9":;;;nnW%G 7#8xJDoo )(4H	)..t4E  LHHUO'88!"235"{{"ww & 3 3CGG <*+ $'J"
 	''"TE\6T7,<(NW  ((##%BOOC,,.JOOJ'KKMEEW%	F	%%"7+"6D#'%<78;.8:
 -788""47 	
r   c                  J   ^  \ rS rSrS	U 4S jjrS rS rS rS rS r	Sr
U =r$ )
r   i  c                ^   > Xl         X l        X0l        X@l        XPl        [
        TU ]  5         g ry   )_msgr   ro   r   r   super__init__)selfrR   r   ro   r   r   	__class__s         r   rI  CompileException.__init__  s'    		r   c                    [        U 5      U R                  U R                  U R                  U R                  U R
                  44$ ry   )ri   rG  r   ro   r   r   rJ  s    r   
__reduce__CompileException.__reduce__  s7    T
TYYTYY!\\4<<9 : 	:r   c                    [        U 5      $ ry   )rL   rN  s    r   __repr__CompileException.__repr__  s    4yr   c                "    U R                  5       $ ry   )get_messagerN  s    r   __str__CompileException.__str__  s    !!r   c                    U R                   $ ry   )rG  rN  s    r   rU  CompileException.get_message  s    yyr   c           	     l   U R                   R                  S5      n[        [        R                  " [        R
                  " [        U5      5      5      5      S-   nSR                  U5      nUR                  SR                  U R                  R                  5       5      5        UR                  SR                  U 5      5        UR                  S5        UR                  SR                  U R                  5      5        UR                  SR                  S	R                  U R                  5      5      5        UR                  S
5        [        U5       H<  u  pVUR                  UR                  US-   5      UR                  5       -   S-   5        M>     UR                  S5        UR!                  5         g )Nr   r   z
{{:0{}d}} z{} zcompilation error: {}
z-----
z	Name: {}
zOptions: {}
r+  zCUDA source:
)r   r   r   mathfloorlog10r   rF   rD   r   upperro   re   r   	enumeraterstripflush)rJ  flinesdigits	linum_fmtilines          r   r   CompileException.dump  s.   !!$'TZZ

3u: 6781< ''/		T\\//123	)0067			##DII./	&&sxx'=>?	 ! 'GAGGI$$QU+dkkm;dBC (				r   )rG  r   ro   r   r   r   )r"   r#   r$   r%   rI  rO  rR  rV  rU  r   r&   __classcell__)rK  s   @r   r   r     s&    :" r   r   c                  L    \ rS rSr  SS jr\R                  4S jrSS jrSr	g)	r   i  Nc                   S U l         [        U[        5      (       a  UR                  S5      n[        U[        5      (       a  UR                  S5      nXl        X l        [        R                  " XX45      U l         XPl        X`l	        g )NzUTF-8)
r4  r  r  r  srcro   r   createProgramr   r   )rJ  rl  ro   r   r   r   r   s          r   rI  _NVRTCProgram.__init__  sh    c5!!**W%CdE"";;w'D	&&s'I 0r   c                    U" 5       (       a  g U R                   (       a!  [        R                  " U R                   5        g g ry   )r4  r   destroyProgram)rJ  is_shutting_downs     r   __del___NVRTCProgram.__del__  s,    88  * r   c                    U R                   (       a4  U R                    H$  n[        R                  " U R                  U5        M&     [        R                  " U R                  U5        S nU R                   (       a8  0 nU R                    H&  n[        R
                  " U R                  U5      XC'   M(     Ub/  UR                  [        R                  " U R                  5      5        U R                  S:X  a"  [        R                  " U R                  5      U4$ U R                  S:X  a"  [        R                  " U R                  5      U4$ U R                  S:X  a"  [        R                  " U R                  5      U4$ [        S5      e! [        R                   a^    [        R                  " U R                  5      n[        XPR                  U R                   U["        R$                  (       d  S5      eS5      ef = f)Nr   r   r   zUnknown NVRTC compile methodr   r#  )r   r   addNameExpressionr4  compileProgramgetLoweredNamerD   getProgramLogr   getCUBINgetPTXgetLTOIRrJ   
NVRTCErrorr   rl  ro   r
   r   )rJ  r   rM   kerr   rP   s         r   r   _NVRTCProgram.compile  s   	P$$00C++DHHc: 1  73G$$00C#(#7#7##FGL 1%  !4!4TXX!>?{{g%~~dhh/88%||DHH-w66%~~dhh/88"#ABB 	P%%dhh/C"3$))W29..7P PFNP P	Ps   DF 1F 	1F ;F A2G8)r   ro   r   r4  rl  )default_programr   r   Nr   )r   N)
r"   r#   r$   r%   rI  r   rq  rr  r   r&   r   r   r   r   r     s#    <>AF (-'='= +Pr   r   c                4    [         R                  " SU 5      S L$ )Nz^[a-zA-Z_][a-zA-Z_0-9]*$)rematch)ro   s    r   is_valid_kernel_namer  5  s    88.5TAAr   c           	     &   SS/[        U5      -   n[        R                  " 5        n[        R                  R                  US5      nUS-   nUS-   n[        US5       n	U	R                  U 5        S S S 5        XGSU/-  n [        XESU5      n
[        R                  R!                  U5      (       d  [        S
R#                  XJ5      5      e[        US5       n	U	R%                  5       sS S S 5        sS S S 5        $ ! , (       d  f       N= f! [         aP  n[        [        U5      XUS5      n[        SS	5      nU(       a  UR                  [        R                  5        UeS nAff = f! , (       d  f       O= f S S S 5        g ! , (       d  f       g = f)Nr7   z--gencokernz.cpp.hsacor   r   r   FzN`hipcc` command does not generate output file. 
command: {}
stdout/stderr: 
{}r   )r  r   r   r8   rO   re   r   rD   rS   r)   r   rL   r   r   r   r3   r   rF   r  )r   r   r   rM   rG   r   rO   in_pathout_pathrb  rI   rQ   r   r   s                 r   compile_using_hipccr  9  sW    I
g
.C		$	$	&(ww||Hf--(?'31GGFO   	x((	SGZ@F ww~~h''  VC(	* *
 (D!Q668 "!9 
'	&
    		"3q66G#*,C *0%9D$I		  "!!9 
'	&	&sa   7FC9/F?D
A
FE'&	F9
D	F

E$AEE$$F'
E5	1F
Fc                   SS/[        U5      -   n[        R                  " 5        n[        R                  R                  US5      nSU-  n[        US5       nUR                  U 5        S S S 5        UR                  U5        [        X#S5      n[        U[        5      (       d   e[        R                  " SSU5      sS S S 5        $ ! , (       d  f       Nc= f! , (       d  f       g = f)Nr7   z--preprocessr  z%s.cppr   z	(?m)^#.*$r0   )r  r   r   r8   rO   re   r   rD   r  rS   r  rL   r  sub)r   r   rG   r   rO   r   r   pp_srcs           r   _preprocess_hipccr  a  s    N
#d7m
3C		$	$	&(ww||Hf-T/'37MM&!   	

70&#&&&&vvk2v. 
'	&  	 
'	&s$   2CC*AC
C	C
C&c                   [         S:  a  SnOSn[        U5      n UR                  U5      u  pE[        U[        5      (       d   eU$ ! [         a8  n[	        SS5      nU(       a  UR                  [        R                  5        e S nAff = f)Nthz}
        // hiprtc segfaults if the input code is empty
        __global__ void _cupy_preprocess_dummy_kernel_() { }
        z
        // hiprtc segfaults if the input code is empty
        #include <hip/hip_runtime.h>
        __global__ void _cupy_preprocess_dummy_kernel_() { }
        r   F)
_cuda_hip_versionr   r   r   r   r   r   r3   r  r  )r   r   coder   r  r  rQ   r   s           r   _preprocess_hiprtcr  p  s    H$
 DLL)	 fe$$$$M  %,e5FF3::s   A 
B
3BB
c                   U(       d  SU -   $ [         S:  a  U $ [         S:  a  SU -   $ [        ci  Ubf  UR                  S5      nU Vs/ s H5  nUR                  S5      (       a  M  UR                  S5      (       a  M3  UPM7     nnSR	                  U5      =qnU R                  S5      n U  Vs/ s H  o3R                  S5      (       a  M  UPM     n nS[        -   SR	                  U 5      -   n U $ s  snf s  snf )Nz#include <hip/hip_runtime.h>
r  i  r   z#includez#pragma oncez7#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
)r  _hip_extra_sourcer   r  re   )r   r  	is_hiprtcrg  s       r   _convert_to_hip_sourcer    s    /&88H$C/&88
  #'--d3L-9 :\TOOJ/ !7 !\L : 04yy/FF\\$F%Ivt__Z-HdvFII!"$(IIf$56F M: Js   C/%C/=C/-C4
C4c	           	     l   [        U5      (       a  [        S5      eUS-  nU[        S [        R                  " 5        5       5      -  nUc  [
        R                  " 5       R                  nU(       a  [        XUS:H  S9n X![        5       U4n	[        R                  U	S 5      n
U
c(  US:X  a  [        SU5      n
O[        SU5      n
U
[        U	'   U	< SU
< SU < SU< 3nUR                  S5      n[        U5      S	-   n[         R"                  " 5       nU(       d3  U(       d+  [$        R'                  U5      nUb  UR'                  U5        U$ O US:X  a%  [)        XX,S
-   UXg5      u  pUR+                  U5        O[-        XX&5      nU(       d  [$        R/                  XU 5        O UR'                  U5        U$ )Nz,separate compilation is not supported in HIP)z-fcuda-flush-denormals-to-zeroc              3  ,   #    U  H
  nS U 3v   M     g7f)z-INr   )r   include_dirs     r   r   *_compile_with_cache_hip.<locals>.<genexpr>  s      AK [MAs   r#  )r  r0   r+  r-  r  r.  )r   r   r   r   _get_hipcc_include_dirsr   r   r   r  r|   r  r=   r  r  r0  r   r   r1  r   r2  r   r:  r  r;  )r   r   r   r  r   r   rM   r   use_converterr2   r=  r?  ro   r@  binaryr   s                   r   r%  r%    s    G$$GHH 22G u '??A  G | }}11'3:h3FI ,.
8C'++C6D|h%b'2D$R1D,0$S)"D&,?GnnW%G7#h.D
//
C
  *//5F! 
 	(-T%<1A) 	!$VdG""48 	HHVJr   )r   r   returnNonery   )r  z
str | None)r   Nkern.cuNNFNN)r   Nr  r   FN)r   )Nr   FNNFFF)r#  NNFT)Q
__future__r   r>   r[  r8   rb   r  rV   r@   r   r   rg   r   r   r   cupy.cuda._compiler_cacher   _DiskKernelCacheBackendr   r   r   cupy_backends.cuda.apir	   r
   cupy_backends.cuda.libsr   cupyr   r   get_build_versionr  rz   r   r   r  r:   r   r   r   __annotations__r   r_   r   r)   r,   rS   memoizer;   rX   rY   r|   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r  r  r'  r&  r   r   r  r  r  r  r  r  r%  r   r   r   <module>r     s{   "   	  	   
     
 * * )   ,,.    & 
	 	 	)0

 .E-F * F	$	I 		Y 		i 	.+b  "($   . ( ( t$	@ %	@ t$* %*"6:< ""95A!"6>&+ #"+J.& -6&*/390z -6&*/3
) 155<>BU$p GF &( d ' /$(t5$uu/@ ;B:>GLEP$y $N0P 0PfB#P/8  > @D=B*.Vr   