o
    ºiŠJ  ã                   @   s   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	m
Z
 d dlmZmZmZmZmZmZ d dlmZ d dlmZmZ edƒG dd	„ d	eeƒƒZedƒG d
d„ deeƒƒZdd„ ZedƒG dd„ deeƒƒZdd„ ZedƒG dd„ deeƒƒZedƒG dd„ deƒƒZdS )é    N)Úcuda)ÚNumbaWarning)ÚCUDATestCaseÚskip_on_cudasimÚskip_unless_cc_60Úskip_if_cudadevrt_missingÚskip_if_mvc_enabledÚtest_data_dir)ÚSerialMixin)ÚDispatcherCacheUsecasesTestÚskip_bad_accessz$Simulator does not implement cachingc                   @   sö   e Zd Zej e¡Zej ed¡Z	dZ
dd„ Zdd„ Zdd„ Zd	d
„ Zdd„ Zdd„ Zdd„ Zdd„ Zdd„ Zdd„ Zeeedƒdd„ ƒƒƒZeeedƒdd„ ƒƒƒZdd„ Zee ejdkd¡d d!„ ƒƒZee ejdkd¡d"d#„ ƒƒZ d$d%„ Z!d&S )'ÚCUDACachingTestúcache_usecases.pyÚcuda_caching_test_fodderc                 C   ó   t  | ¡ t | ¡ d S ©N©r   ÚsetUpr   ©Úself© r   úa/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_caching.pyr      ó   
zCUDACachingTest.setUpc                 C   r   r   ©r   ÚtearDownr   r   r   r   r   r      r   zCUDACachingTest.tearDownc                 C   sÔ   |   d¡ |  ¡ }|   d¡ |j}|  |ddƒd¡ |   d¡ |  |ddƒd¡ |   d¡ |  |jdd¡ |j}||jdƒ}|  t|ƒd¡ |j	}||j
dƒ}|  t|ƒd¡ |   d¡ |  |jdd¡ |  ¡  d S )	Nr   é   é   é   ç      @ç      @é   ©r   g     ÀE@)Úcheck_pycacheÚimport_moduleÚadd_usecaseÚassertPreciseEqualÚ
check_hitsÚfuncÚrecord_return_alignedÚaligned_arrÚtupleÚrecord_return_packedÚ
packed_arrÚrun_in_separate_process)r   ÚmodÚfÚrecr   r   r   Útest_caching!   s$   




zCUDACachingTest.test_cachingc                 C   s.   |   ¡ }|j}|  |ddƒd¡ |  d¡ d S )Nr   r   r   r   )r#   Úadd_nocache_usecaser%   r"   ©r   r.   r/   r   r   r   Útest_no_caching:   s   zCUDACachingTest.test_no_cachingc                 C   s0   |   d¡ |  ¡ }|j}|d ƒ  |   d¡ d S )Nr   )r    r    r   )r"   r#   Úmany_localsr3   r   r   r   Útest_many_localsA   s
   

z CUDACachingTest.test_many_localsc                 C   s¤   |   ¡ }t ¡ @ t dt¡ |j}|  |dƒd¡ |j}|  |dƒd¡ |j}|  |dƒd¡ |j	}|  |dƒd¡ |  
d¡ W d   ƒ d S 1 sKw   Y  d S )NÚerrorr   r   é   é
   é   é   )r#   ÚwarningsÚcatch_warningsÚsimplefilterr   Úclosure1r%   Úclosure2Úclosure3Úclosure4r"   r3   r   r   r   Útest_closureM   s   
"õzCUDACachingTest.test_closurec                 C   sò   |   ¡ }| dd¡ | dd¡ | dd¡ | dd¡ | |jd¡ | |jd¡ | d¡ |  	¡ }|  
|jjdd¡ |   ¡ }|  ||¡ |j}|ddƒ |  
|jdd¡ |ddƒ |  
|jdd¡ |  |  	¡ |¡ |  ¡  |  |  	¡ |¡ d S )Nr   r   r   g      @r   r    )r#   r$   Úouter_uncachedÚouterr+   r,   r(   r)   Úsimple_usecase_callerÚget_cache_mtimesr&   r'   ÚassertIsNotÚassertEqualr-   )r   r.   ÚmtimesÚmod2r/   r   r   r   Útest_cache_reuse]   s(   


z CUDACachingTest.test_cache_reusec                 C   sz   |   ¡ }|j}|  |ddƒd¡ t| jdƒ}| d¡ W d   ƒ n1 s&w   Y  |   ¡ }|j}|  |ddƒd¡ d S )Nr   r   r   Úaz
Z = 10
é   )r#   r$   r%   ÚopenÚmodfileÚwriter3   r   r   r   Útest_cache_invalidatex   s   ÿz%CUDACachingTest.test_cache_invalidatec                 C   s†   |   ¡ }|j}|  |ddƒd¡ |   ¡ }|j}d|_|  |ddƒd¡ |j ¡  |  |ddƒd¡ |   ¡ }|j}|  |ddƒd¡ d S )Nr   r   r   r9   rN   )r#   r$   r%   ÚZr'   Ú	recompiler3   r   r   r   Útest_recompile…   s   
zCUDACachingTest.test_recompilec                 C   s8   |   ¡ }|j}|  |dƒd¡ |j}|  |dƒd¡ d S )Nr   é   r8   )r#   Úrenamed_function1r%   Úrenamed_function2r3   r   r   r   Útest_same_names—   s
   zCUDACachingTest.test_same_nameszCG not supported with MVCc                 C   s<   |   d¡ |  ¡ }|   d¡ | d¡ |   d¡ |  ¡  d S )Nr   r   )r"   r#   Ú
cg_usecaser-   )r   r.   r   r   r   Útest_cache_cgŸ   s   



zCUDACachingTest.test_cache_cgc                 C   st   |   d¡ dt| j| jd }tjtjd|gtjtjd}|j	dd\}}|j
dkr8td|j
| ¡ | ¡ f ƒ‚d S )	Nr   zžif 1:
            import sys

            sys.path.insert(0, %(tempdir)r)
            mod = __import__(%(modname)r)
            mod.cg_usecase(0)
            )ÚtempdirÚmodnamez-c)ÚstdoutÚstderré<   )ÚtimeoutzBprocess failed with code %s: 
stdout follows
%s
stderr follows
%s
)r"   Údictr\   r]   Ú
subprocessÚPopenÚsysÚ
executableÚPIPEÚcommunicateÚ
returncodeÚAssertionErrorÚdecode)r   ÚcodeÚpopenÚoutÚerrr   r   r   Útest_cache_cg_clean_run¯   s    
úþ
ýÿÿz'CUDACachingTest.test_cache_cg_clean_runc                 C   s†   |   ¡ }|j}| jtj|jjjdd |  |ddƒd¡ |  	|jdd¡ |   ¡ }|j}|  |ddƒd¡ |  	|jdd¡ |  
d¡ dS )	zy
        With a disabled __pycache__, test there is a working fallback
        (e.g. on the user-wide cache dir)
        T)Úignore_errorsr   r   r   r   r    N)r#   r$   Ú
addCleanupÚshutilÚrmtreer'   ÚstatsÚ
cache_pathr%   r&   r"   )r   r.   r/   rK   r   r   r   Ú_test_pycache_fallbackÎ   s   ÿz&CUDACachingTest._test_pycache_fallbackÚntz3cannot easily make a directory read-only on Windowsc                 C   s:   t  | j¡j}t  | jd¡ |  t j| j|¡ |  ¡  d S )Né@  )ÚosÚstatr\   Úst_modeÚchmodrr   rw   )r   Ú	old_permsr   r   r   Útest_non_creatable_pycacheè   s   z*CUDACachingTest.test_non_creatable_pycachec                 C   sN   t j | jd¡}t  |¡ t  |¡j}t  |d¡ |  t j||¡ |  	¡  d S )NÚ__pycache__ry   )
rz   ÚpathÚjoinr\   Úmkdirr{   r|   r}   rr   rw   )r   Úpycacher~   r   r   r   Útest_non_writable_pycacheó   s   
z)CUDACachingTest.test_non_writable_pycachec                 C   s\   t td ƒ}d}|  t|¡ tjdd|gddd„ ƒ}W d   ƒ d S 1 s'w   Y  d S )Nzjitlink.ptxz0Cannot pickle CUDACodeLibrary with linking fileszvoid()T)ÚcacheÚlinkc                   S   s   d S r   r   r   r   r   r   r/     s   z>CUDACachingTest.test_cannot_cache_linking_libraries.<locals>.f)Ústrr	   ÚassertRaisesRegexÚRuntimeErrorr   Újit)r   r‡   Úmsgr/   r   r   r   Ú#test_cannot_cache_linking_libraries   s   "þz3CUDACachingTest.test_cannot_cache_linking_librariesN)"Ú__name__Ú
__module__Ú__qualname__rz   r   ÚdirnameÚ__file__Úherer‚   Úusecases_filer]   r   r   r1   r4   r6   rC   rL   rR   rU   rY   r   r   r   r[   rp   rw   r   ÚunittestÚskipIfÚnamer   r…   r   r   r   r   r   r      sD    ÿÿ
r   c                   @   sJ   e Zd Zej e¡Zej ed¡Z	dZ
dd„ Zdd„ Zdd„ Zd	d
„ ZdS )ÚCUDAAndCPUCachingTestzcache_with_cpu_usecases.pyÚ cuda_and_cpu_caching_test_fodderc                 C   r   r   r   r   r   r   r   r     r   zCUDAAndCPUCachingTest.setUpc                 C   r   r   r   r   r   r   r   r     r   zCUDAAndCPUCachingTest.tearDownc                 C   sÔ   |   d¡ |  ¡ }|   d¡ |j}|j}|  |dƒd¡ |   d¡ |  |dƒd¡ |   d¡ |  |jdd¡ |  |jdd¡ |  |dƒd¡ |   d¡ |  |dƒd¡ |   d¡ |  |jdd¡ |  |jdd¡ d S )Nr   r;   r   r   r    ç      @rV   )r"   r#   Ú
assign_cpuÚassign_cudar%   r&   r'   )r   r.   Úf_cpuÚf_cudar   r   r   Útest_cpu_and_cuda_targets  s"   





z/CUDAAndCPUCachingTest.test_cpu_and_cuda_targetsc                 C   s  |   ¡ }| d¡ | d¡ | d¡ | d¡ |  ¡ }|  |jjdd¡ |  |jjdd¡ |   ¡ }|  ||¡ |j}|j}|dƒ |  |jdd¡ |dƒ |  |jdd¡ |dƒ |  |jdd¡ |dƒ |  |jdd¡ |  |  ¡ |¡ |  ¡  |  |  ¡ |¡ d S )Nr;   rš   r   r   r    r   )	r#   r›   rœ   rG   r&   r'   rH   rI   r-   )r   r.   rJ   rK   r   rž   r   r   r   Útest_cpu_and_cuda_reuse0  s.   



z-CUDAAndCPUCachingTest.test_cpu_and_cuda_reuseN)rŽ   r   r   rz   r   r‘   r’   r“   r‚   r”   r]   r   r   rŸ   r    r   r   r   r   r˜   	  s    r˜   c               	   C   s¤   t jd } |  t  ¡ jj}W d   ƒ n1 sw   Y  t jdd … D ]+}| t  ¡ jj}||kr@| |fW  d   ƒ   S W d   ƒ n1 sJw   Y  q$d S )Nr   r    )r   ÚgpusÚcurrent_contextÚdeviceÚcompute_capability)Ú	first_gpuÚfirst_ccÚgpuÚccr   r   r   Úget_different_cc_gpusS  s   
ÿýþ€r©   c                   @   sB   e Zd Zej e¡Zej ed¡Z	dZ
dd„ Zdd„ Zdd„ Zd	S )
ÚTestMultiCCCachingr   Ú!cuda_multi_cc_caching_test_fodderc                 C   r   r   r   r   r   r   r   r   j  r   zTestMultiCCCaching.setUpc                 C   r   r   r   r   r   r   r   r   n  r   zTestMultiCCCaching.tearDownc                 C   s  t ƒ }|s
|  d¡ |  d¡ |  ¡ }|  d¡ |d ^ |j}|  |ddƒd¡ |  d¡ |  |ddƒd¡ |  d¡ |  |jdd¡ |j}||j	dƒ}|  t
|ƒd	¡ |j}||jdƒ}|  t
|ƒd	¡ |  d¡ |  |jdd¡ W d   ƒ n1 s}w   Y  |d ^ |j}|  |ddƒd¡ |  d¡ |  |ddƒd¡ |  d¡ |  |jdd¡ |j}||j	dƒ}|  t
|ƒd	¡ |j}||jdƒ}|  t
|ƒd	¡ |  d¡ |  |jdd¡ W d   ƒ n1 sçw   Y  |  ¡ }|  ||¡ |d ^ |j}|  |ddƒd¡ |  d
¡ |  |ddƒd¡ |  d¡ |  |jdd¡ |j}||j	dƒ}|  t
|ƒd	¡ |j}||jdƒ}|  t
|ƒd	¡ |  d¡ |  |jdd¡ W d   ƒ n	1 s\w   Y  |  ¡ }|  ||¡ |d ? |j}|  |ddƒd¡ |  |ddƒd¡ |j}||j	dƒ}|  t
|ƒd	¡ |j}||jdƒ}|  t
|ƒd	¡ W d   ƒ n	1 s²w   Y  |d @ |j}|  |ddƒd¡ |  |ddƒd¡ |j}||j	dƒ}|  t
|ƒd	¡ |j}||jdƒ}|  t
|ƒd	¡ W d   ƒ d S 1 sÿw   Y  d S )Nz.Need two different CCs for multi-CC cache testr   r   r   r   r   r   r    r!   é   r8   r9   )r©   ÚskipTestr"   r#   r$   r%   r&   r'   r(   r)   r*   r+   r,   rH   )r   r¡   r.   r/   r0   rK   Úmod3r   r   r   Ú
test_cacher  s    






ð



ð



ð
õ
$õzTestMultiCCCaching.test_cacheN)rŽ   r   r   rz   r   r‘   r’   r“   r‚   r”   r]   r   r   r¯   r   r   r   r   rª   d  s    rª   c                  C   s   ddl m}  d| _d| _d S )Nr   ©Úconfig)Ú
numba.corer±   ÚCUDA_LOW_OCCUPANCY_WARNINGSÚCUDA_WARN_ON_IMPLICIT_COPYr°   r   r   r   Úchild_initializerá  s   
rµ   c                   @   sF   e Zd ZdZej e¡Zej 	ed¡Z
dZdd„ Zdd„ Zdd	„ Zd
S )ÚTestMultiprocessCacheFr   Úcuda_mp_caching_test_fodderc                 C   r   r   r   r   r   r   r   r   ô  r   zTestMultiprocessCache.setUpc                 C   r   r   r   r   r   r   r   r   ø  r   zTestMultiprocessCache.tearDownc                 C   sŒ   |   ¡ }|j}d}zt d¡}W n ty   t}Y nw | |t¡}zt| |t	|ƒ¡ƒ}W | 
¡  n| 
¡  w |  |||d  d ¡ d S )Nr   Úspawnr    r   )r#   rF   ÚmultiprocessingÚget_contextÚAttributeErrorÚPoolrµ   ÚsumÚimapÚrangeÚcloserI   )r   r.   r/   ÚnÚctxÚpoolÚresr   r   r   Útest_multiprocessingü  s   ÿz*TestMultiprocessCache.test_multiprocessingN)rŽ   r   r   Ú_numba_parallel_test_rz   r   r‘   r’   r“   r‚   r”   r]   r   r   rÅ   r   r   r   r   r¶   é  s    r¶   z0Simulator does not implement the CUDACodeLibraryc                   @   s   e Zd Zdd„ ZdS )ÚTestCUDACodeLibraryc                 C   sZ   ddl m} tƒ }d}|||ƒ}|  td¡ | ¡  W d   ƒ d S 1 s&w   Y  d S )Nr   )ÚCUDACodeLibraryÚlibraryzCannot pickle unfinalized)Únumba.cuda.codegenrÈ   Úobjectr‰   rŠ   Ú_reduce_states)r   rÈ   Úcodegenr—   Úclr   r   r   Ú!test_cannot_serialize_unfinalized  s   

"ÿz5TestCUDACodeLibrary.test_cannot_serialize_unfinalizedN)rŽ   r   r   rÏ   r   r   r   r   rÇ     s    rÇ   )r¹   rz   rs   rc   re   r•   r<   Únumbar   Únumba.core.errorsr   Únumba.cuda.testingr   r   r   r   r   r	   Únumba.tests.supportr
   Únumba.tests.test_cachingr   r   r   r˜   r©   rª   rµ   r¶   rÇ   r   r   r   r   Ú<module>   s2      vI|'