U
    9%e'                     @   s
  d dl Zd dlZd dlmZ d dlmZmZ d dlmZmZ d dl	m
Z
mZmZ d dlm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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)unittest)skip_on_cudasimskip_if_cuda_includes_missing)CUDATestCasetest_data_dir)CudaAPIErrorLinkerLinkerError)
NvrtcError)require_context)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   c/var/www/html/Darija-Ai-API/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    
&&r0   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    

r9   c                 C   sT   t d\}}t jdt}|d |d  |||f< t   |||f | ||f< d S )N   )r      r   )r   r   r2   r3   r   r4   )r5   r   r8   r7   r   r   r    coop_smem2dR   s
    r<   c                 C   s   t d}|| |< d S Nr   )r   r   )r5   r   r   r   r    simple_maxthreadsZ   s    
r>   i  c                 C   sR   t jt|}t|jd D ]}| | ||< qt|jd D ]}|| ||< q<d S Nr   )r   localr3   	LMEM_SIZEr"   shape)r   Br6   r   r   r   r   r    simple_lmemb   s
    rD   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dd Z	dd Z
dd Zdd Zdd Ze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d+d, Zd-d. Zd/d0 Zd1S )2
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 r=   )r   r   rM   )r#   yr   r   r   r    foo   s    
z%TestLinker._test_linking.<locals>.foo{   r   iA  )r   r   r   i  )
r   declare_devicerM   strr   jitnpr3   r   
assertTrue)rK   eagerrO   argsrQ   r   rC   r   r   r    _test_linkingu   s    
zTestLinker._test_linkingc                 C   s   | j dd d S )NFrX   rZ   rK   r   r   r    test_linking_lazy_compile   s    z$TestLinker.test_linking_lazy_compilec                 C   s   | j dd d S )NTr[   r\   r]   r   r   r    test_linking_eager_compile   s    z%TestLinker.test_linking_eager_compilec                    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 )NrM   rN   z
jitlink.curO   c                    s*   t d}|t| k r& || | |< d S r=   )r   r   len)rr#   r   rM   r   r    kernel   s    
z*TestLinker.test_linking_cu.<locals>.kernelr   r   )r       r:   )r   rS   rT   r   rU   rV   aranger   Z
zeros_liketestingZassert_array_equal)rK   rO   rd   r#   rb   expectedr   rc   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 )NrM   rN   zwarn.cuT)recordvoid(int32)r`   c                    s    |  d S Nr   r#   rc   r   r    rd      s    z6TestLinker.test_linking_cu_log_warning.<locals>.kernelr   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   rS   rT   r   warningscatch_warningsr   rU   assertEqualra   assertInmessage)rK   rO   wrd   r   rc   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 )NrM   rN   zerror.curk   r`   c                    s    |  d S rl   r   rm   rc   r   r    rd      s    z0TestLinker.test_linking_cu_error.<locals>.kernelr   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")
r   rS   rT   r   assertRaisesr
   rU   	exceptionrY   rq   )rK   rO   r(   rd   msgr   rc   r    test_linking_cu_error   s    z TestLinker.test_linking_cu_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 extension .cuhvoid()z
header.cuhr`   c                   S   s   d S rl   r   r   r   r   r    rd      s    z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernelassertRaisesRegexRuntimeErrorr   rU   rK   Zexpected_errrd   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 extensionry   datar`   c                   S   s   d S rl   r   r   r   r   r    rd      s    zDTestLinker.test_linking_file_with_no_extension_error.<locals>.kernelrz   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.cury   r`   c                   S   s   d S rl   r   r   r   r   r    rd      s    z7TestLinker.test_linking_cu_cuda_include.<locals>.kernel)rT   r   r   rU   )rK   rO   rd   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.ar`   c                 S   s   d| d< d S r?   r   rm   r   r   r    r)      s    z2TestLinker.test_try_to_link_nonexistent.<locals>.fznonexistent.a not found)ru   r	   r   rU   rq   rv   rY   )rK   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.re      9   N)	r   rU   r0   
specializerV   emptyr"   assertGreaterget_regs_per_threadrK   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_registersre   r   	r   rU   r0   r   rV   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   re   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   rU   r0   r   r   )rK   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 r=   )	r   r   r   rU   r!   Zget_const_mem_sizeassertGreaterEqualr   nbytes)rK   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 )Nre   r   r   )	r   rU   r0   r   rV   r   r"   get_shared_mem_per_blockrp   )rK   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   rV   r   rU   r9   r   rp   )rK   r   r   r   r   r   r    test_get_shared_mem_per_block  s    z(TestLinker.test_get_shared_mem_per_blockc                 C   s<   t t}|tjdtjdtj}| }| 	|d d S )Nr1   r   i   )
r   rU   r9   r   rV   zerosr   r   r   rp   )rK   r   compiled_specializedr   r   r   r    #test_get_shared_mem_per_specialized  s    
 z.TestLinker.test_get_shared_mem_per_specializedc                 C   s&   t dt}| }| |d d S )Nzvoid(float32[:,::1])r   )r   rU   r<   get_max_threads_per_blockr   )rK   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   rU   r>   r   rV   r   r   r   rq   rw   )rK   r   r   Znelemr5   r(   r   r   r    test_max_threads_exceeded   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 r=   )r   r   r   rV   r   rU   rD   get_local_mem_per_threadr   itemsizerA   r   )rK   r   r   local_mem_size	calc_sizer   r   r    test_get_local_mem_per_thread*  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   rU   rD   r   rV   r   rA   r   r   r   r   r   r   )rK   r   r   r   r   r   r   r    "test_get_local_mem_per_specialized1  s    
z-TestLinker.test_get_local_mem_per_specializedN)__name__
__module____qualname__Z_NUMBA_NVIDIA_BINDING_0_ENVr   rL   rZ   r^   r_   ri   rt   rx   r~   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r    rE   j   s4   

		
rE   __main__)'numpyrV   rn   Znumba.cuda.testingr   r   r   r   r   Znumba.cuda.cudadrv.driverr   r   r	   Znumba.cuda.cudadrv.errorr
   Z
numba.cudar   Znumba.tests.supportr   Znumbar   r   r   r   r   r   r   rf   r   r!   r0   r9   r<   r>   rA   rD   rE   r   mainr   r   r   r    <module>   s,   $0
 R
