o
    ۾iJ                     @   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   X/home/ubuntu/.local/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   zif 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 sw   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 sw   Y  |d @ |j}| |ddd | |ddd |j}||j	d}| t
|d	 |j}||jd}| t
|d	 W d    d S 1 sw   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|'