U
    hâËd€.  ã                   @   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 )Néd   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
