U
    hd.                     @   s  d dl Zd dlZd dlmZmZmZ d dlmZmZm	Z	 d dlm
Z
mZ d dlmZmZmZmZ d dlmZ d dlmZmZ d dlmZmZmZmZmZmZmZ ejd	ejd
Zdd Z dd Z!dd Z"dd Z#dd Z$dZ%dd Z&edG dd de
Z'e(dkre)  dS )    N)skip_if_mvc_enabledskip_unless_cc_53unittest)skip_on_cudasimskip_unless_cuda_pythonskip_if_cuda_includes_missing)CUDATestCasetest_data_dir)CudaAPIErrorLinkerLinkerError
NvrtcError)require_context)TestCaseignore_internal_warnings)cudavoidfloat64int64int32typeoffloat32
   dtypec                 C   s*   t jt}t d}|| d | |< d S )N         ?)r   constZ
array_likeCONST1Dgrid)ACi r#   ]/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudadrv/test_linker.pysimple_const_mem   s    
r%   c                 C   s  d}d}d}	d}
d}d}d}d}d}d}d}d}d}d}d}d}d}d}d}d}t |D ]}||7 }||7 }|	|7 }	|
|7 }
||7 }||9 }||9 }||9 }||9 }||9 }|| }|| }|| }|| }|| }||K }||K }||K }||K }||K }qX|| |	 |
 | | td< | td  || | | | 7  < | td  || | | | 7  < | td  || | | | 7  < d S )Nr   r   r   )ranger   r   )xabcdefZa1Za2a3Za4Za5b1b2Zb3Zb4Zb5c1c2c3Zc4Zc5Zd1Zd2Zd3Zd4Zd5r"   r#   r#   r$   func_with_lots_of_registers   sZ    
&&r4   c                 C   sN   t jd|}t d}|dkr6tdD ]}|||< q(t   || | |< d S )Nd   r   r   )r   sharedarrayr   r&   syncthreads)arydtysmr"   jr#   r#   r$   simple_smemH   s    

r=   c                 C   sT   t d\}}t jdt}|d |d  |||f< t   |||f | ||f< d S )N   )r      r   )r   r   r6   r7   r   r8   )r9   r"   r<   r;   r#   r#   r$   coop_smem2dR   s
    r@   c                 C   s   t d}|| |< d S Nr   )r   r   )r9   r"   r#   r#   r$   simple_maxthreadsZ   s    
rB   i  c                 C   sR   t jt|}t|jd D ]}| | ||< qt|jd D ]}|| ||< q<d S Nr   )r   localr7   	LMEM_SIZEr&   shape)r    Br:   r!   r"   r#   r#   r$   simple_lmemb   s
    rH   z$Linking unsupported in the simulatorc                   @   s(  e Zd ZddiZedd Zdd Zdd Zd	d
 Ze	ddd Z
e	ddd Ze	ddd Zedejeddd Zeejeddd Zdd Zdd Zee	d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d0d1 Z d2d3 Z!d4d5 Z"d6d7 Z#d8S )9
TestLinkerZNUMBA_CUDA_USE_NVIDIA_BINDING0c                 C   s   t jdd}~dS )z9Simply go through the constructor and destructor
        )      )ccN)r   new)selfZlinkerr#   r#   r$   test_linker_basicn   s    zTestLinker.test_linker_basicc                 C   s   t ddattd }|r$dg}ng }t j|d|gidd }tjdgtjd	}tjd
gtjd	}|d || | 	|d dk d S )Nbarint32(int32)zjitlink.ptxzvoid(int32[:], int32[:])linkc                 S   s&   t d}| |  t|| 7  < d S rA   )r   r   rQ   )r'   yr"   r#   r#   r$   foo   s    
z%TestLinker._test_linking.<locals>.foo{   r   iA  )r   r   r   i  )
r   declare_devicerQ   strr	   jitnpr7   r   
assertTrue)rO   eagerrS   argsrU   r    rG   r#   r#   r$   _test_linkingu   s    
zTestLinker._test_linkingc                 C   s   | j dd d S )NFr\   r^   rO   r#   r#   r$   test_linking_lazy_compile   s    z$TestLinker.test_linking_lazy_compilec                 C   s   | j dd d S )NTr_   r`   ra   r#   r#   r$   test_linking_eager_compile   s    z%TestLinker.test_linking_eager_compilezNVIDIA Binding needed for NVRTCc                    st   t dd ttd }t j|gd fdd}tjdtjd}t|}|d	 || |d
 }tj	
|| d S )NrQ   rR   
jitlink.curS   c                    s*   t d}|t| k r& || | |< d S rA   )r   r   len)rr'   r"   rQ   r#   r$   kernel   s    
z*TestLinker.test_linking_cu.<locals>.kernelr   r   )r       r>   )r   rW   rX   r	   rY   rZ   aranger   Z
zeros_liketestingZassert_array_equal)rO   rS   ri   r'   rg   expectedr#   rh   r$   test_linking_cu   s    
zTestLinker.test_linking_cuc              	      s   t dd ttd }tjdd(}t  t jd|gd fdd	}W 5 Q R X | t	|d
d | 
dt|d j | 
dt|d j d S )NrQ   rR   zwarn.cuT)recordvoid(int32)re   c                    s    |  d S Nr#   r'   rh   r#   r$   ri      s    z6TestLinker.test_linking_cu_log_warning.<locals>.kernelr   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   rW   rX   r	   warningscatch_warningsr   rY   assertEqualrf   assertInmessage)rO   rS   wri   r#   rh   r$   test_linking_cu_log_warning   s    z&TestLinker.test_linking_cu_log_warningc              	      s~   t dd ttd }| t"}t jd|gd fdd}W 5 Q R X |jjd }| 	d	| | 	d
| | 	d| d S )NrQ   rR   zerror.curp   re   c                    s    |  d S rq   r#   rr   rh   r#   r$   ri      s    z0TestLinker.test_linking_cu_error.<locals>.kernelr   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")
r   rW   rX   r	   assertRaisesr   rY   	exceptionr]   rv   )rO   rS   r,   ri   msgr#   rh   r$   test_linking_cu_error   s    z TestLinker.test_linking_cu_errorz0NVRTC not available when ctypes binding is used.)envvarsc              	   C   sD   t td }d}| t| tjd|gddd }W 5 Q R X d S )Nrd   zBLinking CUDA source files is not supported with the ctypes bindingvoid()re   c                   S   s   d S rq   r#   r#   r#   r#   r$   r-      s    z8TestLinker.test_linking_cu_ctypes_unsupported.<locals>.f)rX   r	   assertRaisesRegexNotImplementedErrorr   rY   )rO   rS   r|   r-   r#   r#   r$   "test_linking_cu_ctypes_unsupported   s
    z-TestLinker.test_linking_cu_ctypes_unsupportedc              	   C   s2   d}|  t| tddd }W 5 Q R X d S )NzLUse of float16 requires the use of the NVIDIA CUDA bindings and setting the zvoid(float16[::1])c                 S   s   t j| d  d S rC   )r   Zfp16Zhexp10rr   r#   r#   r$   hexp10_vectors   s    zCTestLinker.test_linking_float16_unsupported.<locals>.hexp10_vectors)r   r   r   rY   )rO   r|   r   r#   r#   r$    test_linking_float16_unsupported   s    z+TestLinker.test_linking_float16_unsupportedc              	   C   s8   d}|  t| tjddgddd }W 5 Q R X d S )Nz/Don't know how to link file with extension .cuhr   z
header.cuhre   c                   S   s   d S rq   r#   r#   r#   r#   r$   ri      s    z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernelr   RuntimeErrorr   rY   rO   Zexpected_errri   r#   r#   r$   #test_linking_unknown_filetype_error   s    z.TestLinker.test_linking_unknown_filetype_errorc              	   C   s8   d}|  t| tjddgddd }W 5 Q R X d S )Nz-Don't know how to link file with no extensionr   datare   c                   S   s   d S rq   r#   r#   r#   r#   r$   ri      s    zDTestLinker.test_linking_file_with_no_extension_error.<locals>.kernelr   r   r#   r#   r$   )test_linking_file_with_no_extension_error   s    z4TestLinker.test_linking_file_with_no_extension_errorc                 C   s(   t td }tjd|gddd }d S )Nzcuda_include.cur   re   c                   S   s   d S rq   r#   r#   r#   r#   r$   ri     s    z7TestLinker.test_linking_cu_cuda_include.<locals>.kernel)rX   r	   r   rY   )rO   rS   ri   r#   r#   r$   test_linking_cu_cuda_include   s    z'TestLinker.test_linking_cu_cuda_includec              	   C   sB   |  t}tjddgddd }W 5 Q R X | d|jj d S )Nvoid(int32[::1])znonexistent.are   c                 S   s   d| d< d S rC   r#   rr   r#   r#   r$   r-   	  s    z2TestLinker.test_try_to_link_nonexistent.<locals>.fznonexistent.a not found)rz   r   r   rY   rv   r{   r]   )rO   r,   r-   r#   r#   r$   test_try_to_link_nonexistent  s    z'TestLinker.test_try_to_link_nonexistentc                 C   s8   t t}|jtdftd }| | d dS )a  Ensure that the jitted kernel used in the test_set_registers_* tests
        uses more than 57 registers - this ensures that test_set_registers_*
        are really checking that they reduced the number of registers used from
        something greater than the maximum.rj      9   N)	r   rY   r4   
specializerZ   emptyr&   assertGreaterget_regs_per_threadrO   compiledr#   r#   r$   test_set_registers_no_max  s    
z$TestLinker.test_set_registers_no_maxc                 C   s>   t jddt}|jtdftd }| | d d S )Nr   Zmax_registersrj   r   	r   rY   r4   r   rZ   r   r&   assertLessEqualr   r   r#   r#   r$   test_set_registers_57  s    z TestLinker.test_set_registers_57c                 C   s>   t jddt}|jtdftd }| | d d S )N&   r   rj   r   r   r   r#   r#   r$   test_set_registers_38  s    z TestLinker.test_set_registers_38c                 C   sD   t td d d tttttt}tj|ddt}| | d d S )Nr   r   r   )r   r   r   r   rY   r4   r   r   )rO   sigr   r#   r#   r$   test_set_registers_eager!  s    z#TestLinker.test_set_registers_eagerc                 C   s:   t td d d }t|t}| }| |tj d S rA   )	r   r   r   rY   r%   Zget_const_mem_sizeassertGreaterEqualr   nbytes)rO   r   r   Zconst_mem_sizer#   r#   r$   test_get_const_mem_size&  s    z"TestLinker.test_get_const_mem_sizec                 C   s<   t t}|jtdftd }| }| |d d S )Nrj   r   r   )	r   rY   r4   r   rZ   r   r&   get_shared_mem_per_blockru   )rO   r   shared_mem_sizer#   r#   r$   test_get_no_shared_memory,  s    
z$TestLinker.test_get_no_shared_memoryc                 C   s@   t td d d ttj}t|t}| }| |d d S )Nr   i  )	r   r   r   rZ   r   rY   r=   r   ru   )rO   r   r   r   r#   r#   r$   test_get_shared_mem_per_block2  s    z(TestLinker.test_get_shared_mem_per_blockc                 C   s<   t t}|tjdtjdtj}| }| 	|d d S )Nr5   r   i   )
r   rY   r=   r   rZ   zerosr   r   r   ru   )rO   r   compiled_specializedr   r#   r#   r$   #test_get_shared_mem_per_specialized8  s    
 z.TestLinker.test_get_shared_mem_per_specializedc                 C   s&   t dt}| }| |d d S )Nzvoid(float32[:,::1])r   )r   rY   r@   get_max_threads_per_blockr   )rO   r   max_threadsr#   r#   r$   test_get_max_threads_per_block?  s    z)TestLinker.test_get_max_threads_per_blockc              
   C   sx   t dt}| }|d }tj|tjd}z|d|f | W n0 tk
rr } z| d|j	 W 5 d }~X Y nX d S )Nr   r   r   ZcuLaunchKernel)
r   rY   rB   r   rZ   r   r   r
   rv   r|   )rO   r   r   Znelemr9   r,   r#   r#   r$   test_max_threads_exceededD  s    z$TestLinker.test_max_threads_exceededc                 C   s^   t td d d td d d ttj}t|t}| }ttjj	t
 }| || d S rA   )r   r   r   rZ   r   rY   rH   get_local_mem_per_threadr   itemsizerE   r   )rO   r   r   local_mem_size	calc_sizer#   r#   r$   test_get_local_mem_per_threadN  s
    &z(TestLinker.test_get_local_mem_per_threadc                 C   s\   t t}|tjttjdtjttjdtj}|	 }t
tjjt }| || d S )Nr   )r   rY   rH   r   rZ   r   rE   r   r   r   r   r   r   )rO   r   r   r   r   r#   r#   r$   "test_get_local_mem_per_specializedU  s    
z-TestLinker.test_get_local_mem_per_specializedN)$__name__
__module____qualname__Z_NUMBA_NVIDIA_BINDING_0_ENVr   rP   r^   rb   rc   r   rn   ry   r}   r   r   Zrun_test_in_subprocessr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r#   r#   r#   r$   rI   j   sH   





		
rI   __main__)*numpyrZ   rs   Znumba.cuda.testingr   r   r   r   r   r   r   r	   Znumba.cuda.cudadrv.driverr
   r   r   r   Z
numba.cudar   Znumba.tests.supportr   r   Znumbar   r   r   r   r   r   r   rk   r   r%   r4   r=   r@   rB   rE   rH   rI   r   mainr#   r#   r#   r$   <module>   s*   $0
 v
