U
    ,‰dÚB  ã                   @   s  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
 d dlmZmZ d dlmZ d dlmZmZ d dl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)ÚCUDACodeLibrary)ÚCUDATestCaseÚskip_on_cudasim)ÚSerialMixin)ÚDispatcherCacheUsecasesTestÚskip_bad_access)ÚPathz$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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.pyZcuda_caching_test_fodderc                 C   s   t  | ¡ t | ¡ d S ©N©r   ÚsetUpr   ©Úself© r   úH/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudapy/test_caching.pyr      s    
zCUDACachingTest.setUpc                 C   s   t  | ¡ t | ¡ d S r   ©r   ÚtearDownr   r   r   r   r   r      s    
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   Zadd_nocache_usecaser    r   ©r   r)   r*   r   r   r   Útest_no_caching8   s    zCUDACachingTest.test_no_cachingc                 C   s0   |   d¡ |  ¡ }|j}|d ƒ  |   d¡ d S )Nr   )r   r   r   )r   r   Zmany_localsr-   r   r   r   Útest_many_locals?   s
    

z CUDACachingTest.test_many_localsc              	   C   sŽ   |   ¡ }t ¡ t t dt¡ |j}|  |dƒd¡ |j}|  |dƒd¡ |j}|  |dƒd¡ |j	}|  |dƒd¡ |  
d¡ W 5 Q R X d S )NÚerrorr   r   é   é
   é   é   )r   ÚwarningsÚcatch_warningsÚsimplefilterr   Zclosure1r    Zclosure2Zclosure3Zclosure4r   r-   r   r   r   Útest_closureK   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   Z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   sf   |   ¡ }|j}|  |ddƒd¡ t| jdƒ}| d¡ W 5 Q R X |   ¡ }|j}|  |ddƒd¡ d S )Nr   r   r   Úaz
Z = 10
é   )r   r   r    ÚopenZmodfileÚwriter-   r   r   r   Útest_cache_invalidatev   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   r2   rB   )r   r   r    ÚZr"   Z	recompiler-   r   r   r   Útest_recompileƒ   s    
zCUDACachingTest.test_recompilec                 C   s8   |   ¡ }|j}|  |dƒd¡ |j}|  |dƒd¡ d S )Nr   é   r1   )r   Zrenamed_function1r    Zrenamed_function2r-   r   r   r   Útest_same_names•   s
    zCUDACachingTest.test_same_namesc                 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*   r?   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ÚstatÚtempdirÚst_modeÚchmodrK   rP   )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__rR   )
rS   ÚpathÚjoinrU   ÚmkdirrT   rV   rW   rK   rP   )r   ÚpycacherX   r   r   r   Útest_non_writable_pycacheÂ   s    
z)CUDACachingTest.test_non_writable_pycachec              	   C   sT   t j ttƒjjddd¡}d}|  t|¡  tj	dd|gddd	„ ƒ}W 5 Q R X d S )
NZcudadrvÚdataz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)
rS   r[   r\   r
   Ú__file__ÚparentÚassertRaisesRegexÚRuntimeErrorr   Zjit)r   rb   Úmsgr*   r   r   r   Ú#test_cannot_cache_linking_librariesÏ   s      ÿz3CUDACachingTest.test_cannot_cache_linking_librariesN)Ú__name__Ú
__module__Ú__qualname__rS   r[   Údirnamerc   Úherer\   Úusecases_fileÚmodnamer   r   r,   r.   r/   r8   r@   rE   rG   rI   rP   r	   ÚunittestZskipIfÚnamerY   r_   rh   r   r   r   r   r      s2   ÿÿ
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.pyZ cuda_and_cpu_caching_test_fodderc                 C   s   t  | ¡ t | ¡ d S r   r   r   r   r   r   r   ß   s    
zCUDAAndCPUCachingTest.setUpc                 C   s   t  | ¡ t | ¡ d S r   r   r   r   r   r   r   ã   s    
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   r4   r   r   r   ç      @rH   )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 )Nr4   rs   r   r   r   r   )	r   rt   ru   r;   r!   r"   r<   r=   r(   )r   r)   r>   r?   rv   rw   r   r   r   Útest_cpu_and_cuda_reuse   s.    



z-CUDAAndCPUCachingTest.test_cpu_and_cuda_reuseN)ri   rj   rk   rS   r[   rl   rc   rm   r\   rn   ro   r   r   rx   ry   r   r   r   r   rr   Ù   s   rr   c                  C   sz   t jd } |  t  ¡ jj}W 5 Q R X t jdd … D ]@}|2 t  ¡ jj}||krj| |fW  5 Q R £   S W 5 Q R X q4d S )Nr   r   )r   ÚgpusZcurrent_contextZdeviceZcompute_capability)Z	first_gpuZfirst_ccZgpuÚccr   r   r   Úget_different_cc_gpus#  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   Z!cuda_multi_cc_caching_test_fodderc                 C   s   t  | ¡ t | ¡ d S r   r   r   r   r   r   r   :  s    
zTestMultiCCCaching.setUpc                 C   s   t  | ¡ t | ¡ d S r   r   r   r   r   r   r   >  s    
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 5 Q R X |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 5 Q R X |  ¡ }|  ||¡ |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 5 Q R X |  ¡ }|  ||¡ |d t |j}|  |ddƒd¡ |  |ddƒd¡ |j}||j	dƒ}|  t
|ƒd	¡ |j}||jdƒ}|  t
|ƒd	¡ W 5 Q R X |d t |j}|  |ddƒd¡ |  |ddƒd¡ |j}||j	dƒ}|  t
|ƒd	¡ |j}||jdƒ}|  t
|ƒd	¡ W 5 Q R X d S )Nz.Need two different CCs for multi-CC cache testr   r   r   r   r   r   r   r   é   r1   r2   )r|   ZskipTestr   r   r   r    r!   r"   r#   r$   r%   r&   r'   r<   )r   rz   r)   r*   r+   r?   Zmod3r   r   r   Ú
test_cacheB  s–    
















zTestMultiCCCaching.test_cacheN)ri   rj   rk   rS   r[   rl   rc   rm   r\   rn   ro   r   r   r   r   r   r   r   r}   4  s   r}   c                  C   s   ddl m}  d| _d| _d S )Nr   ©Úconfig)Z
numba.corer   ZCUDA_LOW_OCCUPANCY_WARNINGSZ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   Zcuda_mp_caching_test_fodderc                 C   s   t  | ¡ t | ¡ d S r   r   r   r   r   r   r   Ä  s    
zTestMultiprocessCache.setUpc                 C   s   t  | ¡ t | ¡ d S r   r   r   r   r   r   r   È  s    
zTestMultiprocessCache.tearDownc                 C   s†   |   ¡ }|j}d}zt d¡}W n tk
r8   t}Y nX | |t¡}zt| 	|t
|ƒ¡ƒ}W 5 | ¡  X |  |||d  d ¡ d S )Nr   Úspawnr   r   )r   r:   ÚmultiprocessingZget_contextÚAttributeErrorZPoolr‚   ÚcloseÚsumÚimapÚranger=   )r   r)   r*   ÚnÚctxÚpoolÚresr   r   r   Útest_multiprocessingÌ  s    

z*TestMultiprocessCache.test_multiprocessingN)ri   rj   rk   Z_numba_parallel_test_rS   r[   rl   rc   rm   r\   rn   ro   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   s8   t ƒ }d}t||ƒ}|  td¡ | ¡  W 5 Q R X d S )NZlibraryzCannot pickle unfinalized)Úobjectr   re   rf   Z_reduce_states)r   Zcodegenrq   Zclr   r   r   Ú!test_cannot_serialize_unfinalizedæ  s
    
z5TestCUDACodeLibrary.test_cannot_serialize_unfinalizedN)ri   rj   rk   r’   r   r   r   r   r   á  s   r   )r…   rS   rL   rp   r5   Znumbar   Znumba.core.errorsr   Znumba.cuda.codegenr   Znumba.cuda.testingr   r   Znumba.tests.supportr   Znumba.tests.test_cachingr   r	   Úpathlibr
   r   rr   r|   r}   r‚   rƒ   r   r   r   r   r   Ú<module>   s0    HI|'