
    `i                     L   U d dl Z d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
mZ d dlZd dlmZ d dlmZ d dlmZ d dlmZ d dlmZ d dlmZ d d	lmZ d d
lmZ  ej                    Zdaej                            d          adZda G d de           Z! G d de           Z" G d de           Z#dBdZ$ ej%                    d             Z&dee'         fdZ(dee'         fdZ)d Z* ej%                    d             Z+dZ, ej%                    d             Z- ej%                    d             Z. ej%        d          d             Z/ ej%        d          dBd             Z0d! Z1d" Z2d# Z3d$ Z4 e4d%d&          Z5d&a6d' Z7d( Z8 e9 e8d)                    Z:	 	 	 dCd,Z;	 	 	 dDd.Z<d/ Z=ej>        ?                    d0          Z@d1 ZAi aBeCeDd2<   	 	 dEd&ddd&d4d5ZE	 	 	 dFd6ZF G d7 d8e           ZG G d9 d:eH          ZId; ZJdBd<ZKd= ZLd> ZMdaNd? ZO	 	 	 dGdAZPdS )H    N)Optional)device)function)get_rocm_path)driver)runtimenvrtc)_environment)_utilwin32)
--device-cz-dcz	-rdc=truez--relocatable-device-code=truec                       e Zd ZdS )NVCCExceptionN__name__
__module____qualname__     f/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/cupy/cuda/compiler.pyr   r   !           Dr   r   c                       e Zd ZdS )HIPCCExceptionNr   r   r   r   r   r   %   r   r   r   c                       e Zd ZdS )JitifyExceptionNr   r   r   r   r   r   )   r   r   r   c                    	 t           j        }t          rYt                      }|I|t           j        z   t           j                            dd          z   }t          j        |          }||d<   t          j	        | ||t          j
        dt          rt          j        nd          }||                    |           |S # t          j        $ re}d                    ||j        |j        |j                  }	|dk    rt%          |	          |dk    rt'          |	          t)          |	          d }~wt*          $ r9}d	t-          |          z   }	t+          |	                    |                    d }~ww xY w)
NPATH Tr   )cwdenvstderruniversal_newlinescreationflagsz^`{0}` command returns non-zero exit status. 
command: {1}
return-code: {2}
stdout/stderr: 
{3}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)
r6   r    backend
log_streamr!   
extra_pathpathlogemsgs
             r   _run_ccrC   -   s   ,+ j 	# 233J%!BJ.1K1KKmC(("F%Sc$#:@G:66a	J J J
 !S!!!
( $ $ $ VGELH& &	 	 f$$$ %%%s### + + +2A cjj))***	+s%   B;B> >E3A D--E3:4E..E3c                  ~    t          j        d          } | rd S t                      }|r|S t                      }|r|S d S )Ncl.exe)shutilwhich_get_cl_exe_dir_get_cl_exe_dir_fallback)cl_exe
cl_exe_dirs     r   r*   r*   ^   sU    \(##F t ""J )++J 4r   returnc                     	 	 dd l } n# t          $ r Y d S w xY w| j                            t	          j                              j        }|D ]E}t          j        	                    |d          }t          j        
                    |          r|c S Ft          j        d|            n># t          $ r1}t          j        dt          |           d|            Y d }~nd }~ww xY wd S )Nr   rE   zcl.exe could not be found in z,Failed to find cl.exe with setuptools.msvc: : )setuptools.msvc	ExceptionmsvcEnvironmentInfoplatformmachineVCToolsr'   r?   joinexistswarningswarntype)
setuptoolsvctoolsr?   rJ   rA   s        r   rH   rH   p   s=   K	 #"""" 	 	 	44	/11(2B2D2DEEM 	 	DW\\$11Fw~~f%% ?g??@@@@ K K KI477IIaII	K 	K 	K 	K 	K 	K 	K 	KK 4s4    B) 
B) A9B) B) )
C$3'CC$c                  t   	 ddl m}  ddlm}  | | ddi                    }|                                 |j                                         t          j        	                    |j        j
                  S # t          $ r1}t          j        dt          |           d|            Y d }~nd }~ww xY wd S )Nr   )Distribution)	build_extnamecupy_cl_exe_discoverz'Failed to find cl.exe with setuptools: rN   )r[   r^   setuptools.command.build_extr_   setup_shlib_compilershlib_compiler
initializer'   r?   dirnameccrP   rX   rY   rZ   )r^   r_   extrA   s       r   rI   rI      s   
	F++++++::::::if.D%EFFGG  """%%'''ws14555 F F FDd1ggDDDD	F 	F 	F 	F 	F 	F 	F 	FF 4s   A7A: :
B5'B00B5c                  D    t           t          j                    a t           S N)_nvrtc_versionr
   
getVersionr   r   r   _get_nvrtc_versionrm      s    )++r   c                      ddl m}  | j        S )Nr   core)
cupy._corerp   CUPY_CACHE_KEYro   s    r   _get_cupy_cache_keyrs      s    r   )3253627287c                      t                      \  } }| dk     rd}n;| dk    r	|dk    rd}n,| dk    r	|dk     rd}n| dk    r|dk    s| dk    r	|dk     rd}nd	}|S )
N   75r   80   86   90120)rm   )majorminornvrtc_max_compute_capabilitys      r   _get_max_compute_capabilityr      s    %''LE5rzz'+$$	"!'+$$	" (,$$
2++%1**%2++%!))'+$$ (-$''r   c                  |    t                      \  } }t          d t          j        | |          D                       S )Nc              3       K   | ]	}d | V  
dS )-INr   ).0ds     r   	<genexpr>z._get_extra_include_dir_opts.<locals>.<genexpr>   s<         	Q     r   )rm   tupler   $_get_include_dir_from_conda_or_wheel)r   r   s     r   _get_extra_include_dir_optsr      sQ    %''LE5  B5
 
     r   T)for_each_devicec                      t                      } t          j                    j        }|t          v r|S t          || t                    S )N)key)r   r   Devicecompute_capability_tegra_archsminint)r   archs     r   	_get_archr      sD     $?#@#@ =??-D|453????r   c                     | t                      } t          s3t          |           t          t                                k    rd|  dfS d|  dfS )Nz	-arch=sm_cubinz-arch=compute_ptx)r   _use_ptxr   r   )r   s    r   _get_arch_for_options_for_nvrtcr      s`     |{{+II8::;;;;!4!!7**"D""E))r   c                 4    t          d | D                       S )Nc              3   ,   K   | ]}|t           v |V  d S rj   
_rdc_flagsr   os     r   r   z'_is_cudadevrt_needed.<locals>.<genexpr>   s&      55QQ*__q____55r   )anyoptionss    r   _is_cudadevrt_neededr      s    55'555555r   c                  2   t           t           S ddlm}   |             }|t          d          t          r|dz  }n,|dz   }t
          j                            |          s|dz  }n|}t
          j                            |          st          d          |S )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   r9   r)   r'   r?   isfile)r   	cudadevrtcudadevrt64s      r   _get_cudadevrt_pathr      s     (''''' I/000 $--		"99w~~k** 	$..II#I7>>)$$   	 r   c                 4    t          d | D                       S )Nc              3   ,   K   | ]}|t           v|V  d S rj   r   r   s     r   r   z%_remove_rdc_option.<locals>.<genexpr>  s,      ;;qq
':':':':':':;;r   )r   r   s    r   _remove_rdc_optionr   
  s    ;;G;;;;;;r   c                     t           j                            |           }|t          |          dk    r|S 	 t	          |          dk    S # t
          $ r Y dS w xY w)Nr      F)r'   r(   r,   lenr   
ValueError)r`   defaultvals      r   _get_bool_env_variabler     sd    
*..

C
{c#hh!mm3xx1}   uus   A 
AACUPY_COMPILE_WITH_PTXFc                 
   ddl m} t          sCddlm} |                                 |                    |                                           da| }|dz   | z   } 	 |                    | |          \  }}}}n# t          $ rs}	t          t          |	          |||d          }
t          dd          }|r|
                    t          j                   t          t          |
                    |	d }	~	ww xY w||k    sJ |||fS )	Nr   )jitifyro   T
r   CUPY_DUMP_CUDA_SOURCE_ON_ERRORF)r   r   #_jitify_header_source_map_populatedrq   rp   _init_module_add_sources_get_header_source_maprP   CompileExceptionr;   r   dumpsysr"   r   )sourcer   cu_pathr   rp   
old_sourcer`   headersinclude_namesrA   cexr   s               r   _jitify_prepr     s<          / 3######D7799:::.2+ Jt^f$F/06fg0N0N-gw / / /s1vvz7GXNN%,e5 5 	!HHSZ   c#hh''Q./ 7????G]**s   A8 8
C5A.C00C5c                 R    t          j        | d                                          S )NF)usedforsecurity)hashlibsha1	hexdigest)values    r   _hash_hexdigestr   A  s#    <u555??AAAr   r   r   kern.cuc           	      r   fd}|st          j                    5 }	t          j                            |	|          }
t          |
d          5 }|                    |            d d d            n# 1 swxY w Y    || ||
|||          cd d d            S # 1 swxY w Y   d S |sdn|}
 || ||
|||          S )Nc                    t           j        st                    \  }}||fz  }nd}|rt          | ||          \  }}}	n dx}}	t	                      \  }
}|
dk    r|dz  }t          | |||	||          }	 |                    ||          \  }}nD# t          $ r7}t          dd          }|r|	                    t          j                    d }~ww xY w||fS )Nr   r   r   )z#--device-as-default-execution-space)name_expressionsmethodr   F)r   is_hipr   r   rm   _NVRTCProgramcompiler   r   r   r   r"   )r   r   r   r   r=   r   arch_optr   r   r   major_versionminor_versionprogcompiled_objmappingrA   r   r   s                    r   _compilez%compile_using_nvrtc.<locals>._compileK  s2    ~ 	>tDDHf{"GGF 		D.:/* /*+GWmm ')(Gm+=+?+?(M="" CCVWg}.>vO O O	$(LL*$E$E!L'' 	 	 	)0%9 9D #sz"""	 W$$s   7B 
C2CCwr   )tempfileTemporaryDirectoryr'   r?   rV   openr3   )r   r   r   filenamer   r=   cache_in_memoryr   r   root_dirr   cu_files     `         r   compile_using_nvrtcr   H  s   % % % % %@  ,(** 	Bhgll8X66Ggs## &wf%%%& & & & & & & & & & & & & & & 8FGW,j&B B	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B #0""x2B"F, , 	,s5   1BA/#B/A3	3B6A3	7BBBr   c           	         ddl m} |st                      }|dvrt          d          |dk    r|rJ d                    |          } |            }	|	                                }
|
                    |           t          j                    5 }|                    d          d         }t          j
                            ||          }d	|z  }|d|}t          |d
          5 }|                    |            d d d            n# 1 swxY w Y   |s|
                    d|z             |
t          |          z  }
|
                    |           	 t          |
|d|           n# t           $ rX}t#          t%          |          | ||d          }t'          dd          }|r|                    t*          j                   |d }~ww xY w|
                                }|                    d           |dz   }|
t          |d|fz             z  }
|
                    |           	 t          |
|d|           ne# t           $ rX}t#          t%          |          | ||d          }t'          dd          }|r|                    t*          j                   |d }~ww xY wt1          |          }|d|d|dz   fz  }|t          |          z   }
	 t          |
|d|           n4# t           $ r'}t#          t%          |          dd|d          }|d }~ww xY w|dk    rIt          |d          5 }|                                cd d d            cd d d            S # 1 swxY w Y   nY|dk    rIt          |d          5 }|                                cd d d            cd d d            S # 1 swxY w Y   n
J |            d d d            d S # 1 swxY w Y   d S )Nr   )get_nvcc_path)r   r   z,Invalid code_type %s. Should be cubin or ptxr   z'-gencode=arch=compute_{cc},code=sm_{cc})rg   .z%s.cur   z--%sr%   r   Fz--cubinz.o-oz--device-link.cubinr   rbr   )r   r   r   r   r5   splitappendr   r   r'   r?   rV   r   r3   listrC   r   r   r;   r   r   r   r"   r-   r   read)r   r   r   r   	code_typeseparate_compilationr=   r   arch_str_nvccr6   r   
first_partr?   r   result_pathr   rA   r   r   cmd_partialobjptx_filebin_files                           r   compile_using_nvccr   z  s    (''''' {{(((GHHHE''''8??4?HHHMOOE
++--CJJx		$	&	& A$(^^C((+
w||Hj11D.!%yy1'3 	"7MM&!!!	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" $ .	JJv	)***4== CJJwXvz::::  	 	 	&s1vvvw'-/ / .4e= = )HHSZ(((		 ((**Ky)))+C44+-...CJJwXvz::::  	 	 	&s1vvvw'-/ / .4e= = )HHSZ(((		 )11GdD8ODDGW-CXvz::::    &s1vvr2wGG	 k4(( 'H}}' ' ' ' ' ' 'wA$ A$ A$ A$ A$ A$ A$ A$v' ' ' ' ' ' ' ' ''!!k4(( 'H}}' ' ' ' ' ' '}A$ A$ A$ A$ A$ A$ A$ A$|' ' ' ' ' ' ' ' ' $)##5CA$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$ A$s   AO(D
>O
D	OD	AOE+)O+
G5AGGAO+H>=O>
J AJJ  0OK$#O$
L."LLO.MOM	O"M	#O=N*O*N.	.O1N.	2OOOc                    |dk    r|d                     |          fz  }t          |           }	 |                    |          \  }}n# t          $ r7}t	          dd          }|r|                    t          j                    d }~ww xY w|dk    r_	 |dz   }t          | ||dd	          }nV# t          $ r7}t	          dd          }|r|                    t          j                    d }~ww xY wt          d
|z            t          |t                    sJ d                    d |                                                                D                       S )Nr
   z-arch=compute_{}r   Fr%   )r   zpreprocess.ptxzpreprocess.cur   )r   Invalid backend %sr   c              3   D   K   | ]}|                     d           |V  dS )z//N
startswith)r   xs     r   r   z_preprocess.<locals>.<genexpr>  sP       G G1<<3E3EG	G G G G G Gr   )r5   r   r   r   r   r   r   r"   r   r   
isinstancebytesrV   decode
splitlines)	r   r   r   r<   r   result_rA   r   s	            r   _preprocessr    s   ' 	&--d3355V$$	W--IFAA 	 	 	)0%9 9D #sz"""	 
F				 88G'279 9 9FF 	 	 	)0%9 9D #sz"""	 -7888fe$$$$$ 99 G G==??--//G G G G G Gs/   A	 	
B
2BB
B. .
C/82C**C/z~/.cupy/kernel_cachec                  L    t           j                            dt                    S )NCUPY_CACHE_DIR)r'   r(   r,   _default_cache_dirr   r   r   get_cache_dirr    s    :>>*,>???r   _empty_file_preprocess_cacher
   )enable_cooperative_groupsr   r=   r   c                   |rt           j        rt          d          ||dk    rt          t	          dd          o|dk    }
t           j        r!|dk    rdnd}t          | ||||||||
	  	        S t          | |||||||||
|	          S )Nz+Cooperative groups is not supported in HIP.r
   CUPY_CACHE_IN_MEMORYFhiprtcr&   )r   r   r   NotImplementedErrorr   _compile_with_cache_hip_compile_with_cache_cuda)r   r   r   	cache_dirextra_sourcer<   r  r   r=   r   r   s              r   _compile_module_with_cacher    s    
 ! ?> 	?=? ? ? #7(:(:!!
 	5u== 	w  ~ 	%%00((g&GT9lGj/; ; 	; (GT9lG%'7V% % 	%r   c           
         |t                      }|t                      }|dz  }|r|dz  }t          dd          r|dz  }d|v }|
r|s|dz  }n|r|
sd}
|
r|d	k    rt          d
          |t	                      z  }||t                      |ft          |          z   }t                              |d           }|t          d|||          }|t          |<   |d|d| d|dt                      	}|                    d          }t          |          dz   }t          j                    }|	s(t          j                            |          st          j        |d           t          j                            ||          }t          j                            |          r|st+          |d          5 }|                                }d d d            n# 1 swxY w Y   t/          |          t0          k    r]|d t0                   }|t0          d          }t          |                              d          }||k    r|                    |           |S n	 |d	k    r|	rdn|dz   }t5          | ||||||	|
          \  }}t7          |          rat          j                    }|                    |d           t=                      }|                    |           |                                 }n|}|!                    |           nA|dk    r)t7          |          }tE          | |||dz   d||          }nt          d|z            |	st          |                              d          }tG          j$        |d          5 }|%                    |           |%                    |           |j&        }d d d            n# 1 swxY w Y   tO          j(        ||           t          dd          r@t+          |dz   d          5 }|%                    |            d d d            n# 1 swxY w Y   n	 |                    |           |S )N)z	-ftz=true)r   CUPY_CUDA_COMPILE_WITH_DEBUGF)z--device-debugz--generate-line-info-DCUPY_USE_JITIFY)r  Tr
   zjitify only works with NVRTCr    utf-8r   exist_okr   ascii.cuzcupy.ptxr%   r   )r   r   r=   r  dirdeleteCUPY_CACHE_SAVE_CUDA_SOURCEr   ))r  r   r   r   r   rm   r   r  r,   r  rs   encoder   r   Moduler'   r?   isdirmakedirsrV   rW   r   r   r   _hash_lengthloadr   r   	LinkStateadd_ptr_datar   add_ptr_filecomplete_set_mappingr   r   NamedTemporaryFiler3   r`   rF   move) r   r   r   r  r  r<   r  r   r=   r   r   is_jitify_requestedr!   basekey_srcr`   modr?   filedatahashr   
cubin_hashcu_namer   r   lsr   rdctf	temp_pathfs                                    r   r  r    sP    !OO	|{{~G  #?"<eDD >==.'9 )  	))	 V  9'W$$7888*,,,G'-//9,T223C'++C66D|2wg66,0$S) 	TTT666<<<)<)>)>)>@GnnW%%G7##h.D
/

C w}}Y'' 	2K	D1111
 w||It,,7>>$ 		(8 		dD!! #Tyy{{# # # # # # # # # # # # # # #4yyL((M\M*\]]+,U33::7CC
:%%HHUOOOJ 	''9""TE\*GT7,<1 1W  (( 	#%%BOOC,,,,..JOOJ'''KKMMEEE!!!!	F		"7++"67D#'%<78;.8: : :
 -7888 $U++227;;
 (YuEEE 	 HHZ   HHUOOOI	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	It$$$ ""?GG 	 dUlC((  A                              	HHUOOOJs6   8GG G2OOOP))P-0P-c                   >     e Zd Zd fd	Zd Zd Zd Zd Zd Z xZ	S )	r   r
   c                     || _         || _        || _        || _        || _        t          t          |                                            d S rj   )_msgr   r`   r   r<   superr   __init__)selfrB   r   r`   r   r<   	__class__s         r   rI  zCompileException.__init__  sI    		%%..00000r   c                 `    t          |           | j        | j        | j        | j        | j        ffS rj   )rZ   rG  r   r`   r   r<   rJ  s    r   
__reduce__zCompileException.__reduce__  s/    T

TYTY!\4<9 : 	:r   c                      t          |           S rj   )r;   rM  s    r   __repr__zCompileException.__repr__  s    4yyr   c                 *    |                                  S rj   )get_messagerM  s    r   __str__zCompileException.__str__  s    !!!r   c                     | j         S rj   )rG  rM  s    r   rR  zCompileException.get_message  s
    yr   c           	         | j                             d          }t          t          j        t          j        t          |                                        dz   }d                    |          }|                    d                    | j	        
                                                     |                    d                    |                      |                    d           |                    d                    | j                             |                    d                    d	                    | j                                       |                    d
           t          |          D ]H\  }}|                    |                    |dz             |                                z   dz              I|                    d           |                                 d S )Nr   r   z
{{:0{}d}} z{} zcompilation error: {}
z-----
z	Name: {}
zOptions: {}
r   zCUDA source:
)r   r   r   mathfloorlog10r   r5   r3   r<   upperr`   rV   r   	enumeraterstripflush)rJ  rD  linesdigits	linum_fmtilines          r   r   zCompileException.dump  s   !!$''TZ
3u:: 6 677881< ''//		T\//1122333	)0066777			##DI..///	&&sxx'='=>>???	 !!! '' 	D 	DGAtGGI$$QU++dkkmm;dBCCCC								r   r	   )
r   r   r   rI  rN  rP  rS  rR  r   __classcell__)rK  s   @r   r   r     s        1 1 1 1 1 1: : :  " " "        r   r   c                   6    e Zd Z	 	 ddZej        fdZd	dZdS )
r   default_programr   Nr   c                 ,   d | _         t          |t                    r|                    d          }t          |t                    r|                    d          }|| _        || _        t          j        ||||          | _         || _        || _	        d S )NzUTF-8)
ptrr  r  r	  srcr`   r
   createProgramr   r   )rJ  rg  r`   r   r   r   r   s          r   rI  z_NVRTCProgram.__init__  s    c5!! 	&**W%%CdE"" 	(;;w''D	&sD'=II 0r   c                 b     |            rd S | j         rt          j        | j                    d S d S rj   )rf  r
   destroyProgram)rJ  is_shutting_downs     r   __del__z_NVRTCProgram.__del__  sE     	F8 	+ *****	+ 	+r   c                    	 | j         r$| j         D ]}t          j        | j        |           t          j        | j        |           d }| j         r)i }| j         D ]}t          j        | j        |          ||<    |,|                    t          j        | j                             | j        dk    rt          j	        | j                  |fS | j        dk    rt          j
        | j                  |fS t          d          # t          j        $ rE t          j        | j                  }t          || j        | j        |t           j        sdnd          w xY w)Nr   r   zUnknown NVRTC compile methodr
   r  )r   r
   addNameExpressionrf  compileProgramgetLoweredNamer3   getProgramLogr   getCUBINgetPTXr9   
NVRTCErrorr   rg  r`   r   r   )rJ  r   r=   kerr   r@   s         r   r   z_NVRTCProgram.compile  sy   	P$ ;0 ; ;C+DHc:::: 7333G$ G0 G GC#(#7##F#FGCLL%  !4TX!>!>???{g%%~dh//88%%|DH--w66 ##ABBB 	P 	P 	P%dh//C"3$)W29.#N77hP P P	Ps   C
D %D 3D AE)rd  r   r   Nr   )r   N)r   r   r   rI  r   rk  rl  r   r   r   r   r   r     sd        <>AF    (-'= + + + +P P P P P Pr   r   c                 0    t          j        d|           d uS )Nz^[a-zA-Z_][a-zA-Z_0-9]*$)rematch)r`   s    r   is_valid_kernel_namery    s    8.55TAAr   c           	      l   ddgt          |          z   }t          j                    5 }t          j                            |d          }|dz   }|dz   }t          |d          5 }	|	                    |            d d d            n# 1 swxY w Y   ||d|gz  }	 t          ||d|          }
ne# t          $ rX}t          t          |          | ||d          }t          dd	          }|r|                    t          j                   |d }~ww xY wt          j                            |          s#t          d
                    ||
                    t          |d          5 }	|	                                cd d d            cd d d            S # 1 swxY w Y   	 d d d            d S # 1 swxY w Y   d S )Nr&   z--gencokern.cpp.hsacor   r   r   FzP`hipcc` command does not generate output file. 
command: {0}
stdout/stderr: 
{1}r   )r   r   r   r'   r?   rV   r   r3   rC   r   r   r;   r   r   r   r"   r   r5   r   )r   r   r   r=   r6   r   r?   in_pathout_pathrD  r8   rA   r   r   s                 r   compile_using_hipccr    s    I
g
.C		$	&	& (w||Hf---(?'3 	1GGFOOO	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	x((	S(GZ@@FF 		 		 		"3q6667G#*, ,C *0%9 9D %$$$I		 w~~h'' 	+  fS&))	+ + +
 (D!! 	Q6688	 	 	 	 	 	 	9       8	 	 	 	 	 	 	 	 	9                 s   ;F)#B9F)B			F)B		F)B,+F),
D6AD		DAF)#F7F)F	F)F	F))F-0F-c                    ddgt          |          z   }t          j                    5 }t          j                            |d          }d|z  }t          |d          5 }|                    |            d d d            n# 1 swxY w Y   |                    |           t          ||d          }t          |t                    sJ t          j        dd|          cd d d            S # 1 swxY w Y   d S )Nr&   z--preprocessr{  z%s.cppr   z	(?m)^#.*$r   )r   r   r   r'   r?   rV   r   r3   r   rC   r  r;   rw  sub)r   r   r6   r   r?   r   r   pp_srcs           r   _preprocess_hipccr    se   N
#d7mm
3C		$	&	& 
/(w||Hf--T/'3 	"7MM&!!!	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	

7h00&#&&&&&vk2v..
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/ 
/s6   6C*B 4C* B	C*B	AC**C.1C.c                 .   t           dk    rd}nd}t          |          }	 |                    |          \  }}nD# t          $ r7}t	          dd          }|r|                    t          j                    d }~ww xY wt          |t                    sJ |S )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   r"   r  r  )r   r   coder   r  r  rA   r   s           r   _preprocess_hiprtcr  -  s    H$$
 DLL))	   %,e5 5 	FF3: fe$$$$$Ms   : 
A;2A66A;c                 R   |sd| z   S t           dk    r| S t           dk    rd| z   S t          :|8|                    d          }d |D             }d                    |          xa}|                     d          } d | D             } dt          z   d                    |           z   } | S )Nz#include <hip/hip_runtime.h>
r  i  r   c                 f    g | ].}|                     d           |                     d          ,|/S )#includez#pragma oncer  r   ra  s     r   
<listcomp>z*_convert_to_hip_source.<locals>.<listcomp>\  sL     : : :TOOJ//:77:D : : :r   c                 <    g | ]}|                     d           |S )r  r  r  s     r   r  z*_convert_to_hip_source.<locals>.<listcomp>b  s)    IIItT__Z-H-HIdIIIr   z7#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
)r  _hip_extra_sourcer   rV   )r   r  	is_hiprtcs      r   _convert_to_hip_sourcer  L  s     9/&88H$$C/&88
  #'--d33L: :\ : : :L 04yy/F/FF\\$FIIvIIIFI!"$(IIf$5$56F Mr   r  c
           	      t   t          |          rt          d          |dz  }t          j                    }
|
dk    r|
dk     r|dt	                      z   dz   fz  }|t                      }|t          j                    j        }|	rt          | ||dk              } ||t                      |f}t                              |d           }|1|dk    rt          d	|          }nt          d	|          }|t          |<   |d
|d
| d
|}|                    d          }t!          |          dz   }t#          j                    }|s(t&          j                            |          st'          j        |d           t&          j                            ||          }t&          j                            |          r|st3          |d          5 }|                                }d d d            n# 1 swxY w Y   t7          |          t8          k    r]|d t8                   }|t8          d          }t!          |                              d          }||k    r|                    |           |S n	 |dk    r1t=          | |||dz   |||          \  }}|                    |           ntA          | |||          }|st!          |                              d          }tC          j"        |d          5 }|#                    |           |#                    |           |j$        }d d d            n# 1 swxY w Y   tK          j&        ||           tO          dd          r@t3          |dz   d          5 }|#                    |            d d d            n# 1 swxY w Y   n	 |                    |           |S )Nz,separate compilation is not supported in HIP)z-fcuda-flush-denormals-to-zeroifi ir   z/llvm/lib/clang/13.0.0/include/r  )r  r   r   r!  r}  Tr"  r   r$  r%  Fr&  r)  r|  r   )(r   r   r   get_build_versionr   r  r   r   r   r  rm   r  r,   r  r  r*  r   r   r+  r'   r?   r,  r-  rV   rW   r   r   r   r.  r/  r   r4  r  r   r5  r3   r`   rF   r6  r   )r   r   r   r  r  r<   r   r=   r   use_converterrocm_build_versionr!   r8  r9  r`   r:  r?   rD  r<  
hash_valuebinarybinary_hashr   rB  rC  s                            r   r  r  j  s    G$$ IGHHH 22G  133X%%*<x*G*G=??"%FFI 	I !OO	
 | }1 I'3:h3FI I I ,..
8C'++C66D|h%b'22DD$R11D,0$S)"ssDDD&&&,,?GnnW%%G7##h.D
/

C w}}Y'' 	2K	D1111
 w||It,,7>>$ 		(8 		dD!!  Qvvxx                             4yyL((!-<-0
lmm,-f55<<WEE,,HHV$$$J 	(-GT4%<1A) ) 	!!!!$VWdJGG %f--44W== (YuEEE 	 HH[!!!HHVI	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	It$$$ ""?GG 	 dVmS))  Q                              	HHVJs6   G..G25G232L11L58L54NNNrj   )r   Nr   NNFF)r   Nr   r   FN)r   NNNr
   )Nr
   FNNFF)r  NNFT)Qr-   r   rV  r'   rS   rw  rF   r/   r   r   typingr   rX   r   r   r   r   cupy_backends.cuda.apir   r   cupy_backends.cuda.libsr
   cupyr   r   r  r  rk   r  r)   r   r   rP   r   r   r   rC   memoizer*   r;   rH   rI   rm   rs   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r.  r   r   r  r?   
expanduserr  r  r  dict__annotations__r  r  r   objectr   ry  r  r  r  r  r  r  r   r   r   <module>r     s      				  				      



                     # # # # # # ) ) ) ) ) ) * * * * * * ) ) ) ) ) )            ,F,..  		 	 	)	)0

	 	 	 	 	I 	 	 		 	 	 	 	Y 	 	 		 	 	 	 	i 	 	 	.+ .+ .+ .+b   "#    ((3-    $      . ( ( (,    t$$$	@ 	@ %$	@ t$$$* * * %$*"6 6 6  :< < <   "!"95AA&+ #"+ "+ "+JB B B s??3''(( AJ:>6;/, /, /, /,d 155<>BU$ U$ U$ U$p G  G  GF W''(>?? @ @ @ &( d ' ' ' EI%6;$u% % % % %@ FM:>7<w w w wt$ $ $ $ $y $ $ $N0P 0P 0P 0P 0PF 0P 0P 0PfB B B# # # #P/ / /  8    > @D=B*.o o o o o or   