o
    8ήcB                     @   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.pycuda_caching_test_fodderc                 C      t |  t|  d S Nr   setUpr   self r   K/tmp/pip-target-vg8gfxp4/lib/python/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_caching8   s   zCUDACachingTest.test_no_cachingc                 C   s0   |  d |  }|j}|d   |  d d S )Nr   )r   r   r   )r    r!   many_localsr1   r   r   r   test_many_locals?   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    r1   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"   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writer1   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   r7   rL   )r!   r"   r#   Zr%   	recompiler1   r   r   r   test_recompile   s   
zCUDACachingTest.test_recompilec                 C   s8   |   }|j}| |dd |j}| |dd d S )Nr      r6   )r!   renamed_function1r#   renamed_function2r1   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-   rI   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chmodrY   r^   )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__r`   )
ra   pathjoinrc   mkdirrb   rd   re   rY   r^   )r   pycacherf   r   r   r   test_non_writable_pycache   s   
z)CUDACachingTest.test_non_writable_pycachec                 C   sj   t jttjjddd}d}| t| tj	dd|gddd	 }W d    d S 1 s.w   Y  d S )
N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)
ra   ri   rj   r
   __file__parentassertRaisesRegexRuntimeErrorr   jit)r   rq   msgr-   r   r   r   #test_cannot_cache_linking_libraries   s   "z3CUDACachingTest.test_cannot_cache_linking_librariesN)__name__
__module____qualname__ra   ri   dirnamerr   hererj   usecases_filemodnamer   r   r/   r2   r4   rA   rJ   rP   rS   rW   r^   r	   unittestskipIfnamerg   rm   rx   r   r   r   r   r      s4    
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   r9   r   r   r         @rT   )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 )Nr9   r   r   r   r   r   )	r!   r   r   rE   r$   r%   rF   rG   r+   )r   r,   rH   rI   r   r   r   r   r   test_cpu_and_cuda_reuse   s.   



z-CUDAAndCPUCachingTest.test_cpu_and_cuda_reuseN)ry   rz   r{   ra   ri   r|   rr   r}   rj   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_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   !cuda_multi_cc_caching_test_fodderc                 C   r   r   r   r   r   r   r   r   :  r   zTestMultiCCCaching.setUpc                 C   r   r   r   r   r   r   r   r   >  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      r6   r7   )r   skipTestr    r!   r"   r#   r$   r%   r&   r'   r(   r)   r*   rF   )r   r   r,   r-   r.   rI   mod3r   r   r   
test_cacheB  s   
















$zTestMultiCCCaching.test_cacheN)ry   rz   r{   ra   ri   r|   rr   r}   rj   r~   r   r   r   r   r   r   r   r   r   4  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!   rD   multiprocessingget_contextAttributeErrorPoolr   sumimaprangecloserG   )r   r,   r-   nctxpoolresr   r   r   test_multiprocessing  s   z*TestMultiprocessCache.test_multiprocessingN)ry   rz   r{   _numba_parallel_test_ra   ri   r|   rr   r}   rj   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   sN   t  }d}t||}| td |  W d    d S 1 s w   Y  d S )NlibraryzCannot pickle unfinalized)objectr   rt   ru   _reduce_states)r   codegenr   clr   r   r   !test_cannot_serialize_unfinalized  s   

"z5TestCUDACodeLibrary.test_cannot_serialize_unfinalizedN)ry   rz   r{   r   r   r   r   r   r     s    r   )r   ra   rZ   r   r:   numbar   numba.core.errorsr   numba.cuda.codegenr   numba.cuda.testingr   r   numba.tests.supportr   numba.tests.test_cachingr   r	   pathlibr
   r   r   r   r   r   r   r   r   r   r   r   <module>   s2     HI|'