o
    i'                     @   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re&  dS 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   const
array_likeCONST1Dgrid)ACi r    b/home/ubuntu/transcripts/venv/lib/python3.10/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 ]R}||7 }||7 }|	|7 }	|
|7 }
||7 }||9 }||9 }||9 }||9 }||9 }|| }|| }|| }|| }|| }||K }||K }||K }||K }||K }q,|| |	 |
 | | td< | td  || | | | 7  < | td  || | | | 7  < | td  || | | | 7  < d S )Nr   r   r   )ranger   r   )xabcdefa1a2a3a4a5b1b2b3b4b5c1c2c3c4c5d1d2d3d4d5r   r    r    r!   func_with_lots_of_registers   sZ   
&&*r?   c                 C   sN   t jd|}t d}|dkrtdD ]}|||< qt   || | |< d S )Nd   r   r   )r   sharedarrayr   r#   syncthreads)arydtysmr   jr    r    r!   simple_smemH   s   

rH   c                 C   sT   t d\}}t jdt}|d |d  |||f< t   |||f | ||f< d S )N   )r      r   )r   r   rA   rB   r   rC   )rD   r   rG   rF   r    r    r!   coop_smem2dR   s
   rK   c                 C   s   t d}|| |< d S Nr   )r   r   )rD   r   r    r    r!   simple_maxthreadsZ   s   
rM   i  c                 C   sR   t jt|}t|jd D ]}| | ||< qt|jd D ]}|| ||< qd S Nr   )r   localrB   	LMEM_SIZEr#   shape)r   BrE   r   r   r    r    r!   simple_lmemb   s   rS   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
TestLinkerNUMBA_CUDA_USE_NVIDIA_BINDING0c                 C   s   t jdd}~dS )z9Simply go through the constructor and destructor
        )      )ccN)r   new)self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 rL   )r   r   r^   )r$   yr   r    r    r!   foo   s   
z%TestLinker._test_linking.<locals>.foo{   r   iA  )r   r   r   i  )
r   declare_devicer^   strr   jitnprB   r   
assertTrue)r[   eagerr`   argsrb   r   rR   r    r    r!   _test_linkingu   s   
zTestLinker._test_linkingc                 C      | j dd d S )NFri   rk   r[   r    r    r!   test_linking_lazy_compile      z$TestLinker.test_linking_lazy_compilec                 C   rl   )NTrm   rn   ro   r    r    r!   test_linking_eager_compile   rq   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 )Nr^   r_   z
jitlink.cur`   c                    s.   t d}|t| k r || | |< d S d S rL   )r   r   len)rr$   r   r^   r    r!   kernel   s   
z*TestLinker.test_linking_cu.<locals>.kernelr   r   )r       rI   )r   rd   re   r   rf   rg   aranger   
zeros_liketestingassert_array_equal)r[   r`   rw   r$   ru   expectedr    rv   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 d    n1 s.w   Y  | t	|d
d | 
dt|d j | 
dt|d j d S )Nr^   r_   zwarn.cuT)recordvoid(int32)rs   c                        |  d S Nr    r$   rv   r    r!   rw         z6TestLinker.test_linking_cu_log_warning.<locals>.kernelr   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   rd   re   r   warningscatch_warningsr   rf   assertEqualrt   assertInmessage)r[   r`   wrw   r    rv   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 d    n1 s*w   Y  |jjd }| 	d	| | 	d
| | 	d| d S )Nr^   r_   zerror.cur   rs   c                    r   r   r    r   rv   r    r!   rw      r   z0TestLinker.test_linking_cu_error.<locals>.kernelr   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")
r   rd   re   r   assertRaisesr
   rf   	exceptionrj   r   )r[   r`   r)   rw   msgr    rv   r!   test_linking_cu_error   s   z TestLinker.test_linking_cu_errorc                 C   N   d}|  t| tjddgddd }W d    d S 1 s w   Y  d S )Nz/Don't know how to link file with extension .cuhvoid()z
header.cuhrs   c                   S      d S r   r    r    r    r    r!   rw         z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernelassertRaisesRegexRuntimeErrorr   rf   r[   expected_errrw   r    r    r!   #test_linking_unknown_filetype_error   
   "z.TestLinker.test_linking_unknown_filetype_errorc                 C   r   )Nz-Don't know how to link file with no extensionr   datars   c                   S   r   r   r    r    r    r    r!   rw      r   zDTestLinker.test_linking_file_with_no_extension_error.<locals>.kernelr   r   r    r    r!   )test_linking_file_with_no_extension_error   r   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   rs   c                   S   r   r   r    r    r    r    r!   rw      r   z7TestLinker.test_linking_cu_cuda_include.<locals>.kernel)re   r   r   rf   )r[   r`   rw   r    r    r!   test_linking_cu_cuda_include   s   z'TestLinker.test_linking_cu_cuda_includec                 C   sV   |  t}tjddgddd }W d    n1 sw   Y  | d|jj d S )Nvoid(int32[::1])znonexistent.ars   c                 S   s   d| d< d S rN   r    r   r    r    r!   r*      r   z2TestLinker.test_try_to_link_nonexistent.<locals>.fznonexistent.a not found)r   r	   r   rf   r   r   rj   )r[   r)   r*   r    r    r!   test_try_to_link_nonexistent   s
   z'TestLinker.test_try_to_link_nonexistentc                 C   s:   t t}|jtdgtdR  }| | 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.rx      9   N)	r   rf   r?   
specializerg   emptyr#   assertGreaterget_regs_per_threadr[   compiledr    r    r!   test_set_registers_no_max   s   
z$TestLinker.test_set_registers_no_maxc                 C   @   t jddt}|jtdgtdR  }| | d d S )Nr   max_registersrx   r   	r   rf   r?   r   rg   r   r#   assertLessEqualr   r   r    r    r!   test_set_registers_57      z TestLinker.test_set_registers_57c                 C   r   )N&   r   rx   r   r   r   r    r    r!   test_set_registers_38   r   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   rf   r?   r   r   )r[   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 rL   )	r   r   r   rf   r"   get_const_mem_sizeassertGreaterEqualr   nbytes)r[   r   r   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gtdR  }| }| |d d S )Nrx   r   r   )	r   rf   r?   r   rg   r   r#   get_shared_mem_per_blockr   )r[   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   rg   r   rf   rH   r   r   )r[   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 )Nr@   r   i   )
r   rf   rH   r   rg   zerosr   r   r   r   )r[   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   rf   rK   get_max_threads_per_blockr   )r[   r   max_threadsr    r    r!   test_get_max_threads_per_block  s   z)TestLinker.test_get_max_threads_per_blockc              
   C   sz   t dt}| }|d }tj|tjd}z|d|f | W d S  ty< } z| d|j	 W Y d }~d S d }~ww )Nr   r   r   cuLaunchKernel)
r   rf   rM   r   rg   r   r   r   r   r   )r[   r   r   nelemrD   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 rL   )r   r   r   rg   r   rf   rS   get_local_mem_per_threadr   itemsizerP   r   )r[   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   rf   rS   r   rg   r   rP   r   r   r   r   r   r   )r[   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___NUMBA_NVIDIA_BINDING_0_ENVr   r]   rk   rp   rr   r~   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r    r    r    r!   rT   j   s6    

		
rT   __main__)'numpyrg   r   numba.cuda.testingr   r   r   r   r   numba.cuda.cudadrv.driverr   r   r	   numba.cuda.cudadrv.errorr
   
numba.cudar   numba.tests.supportr   numbar   r   r   r   r   r   r   ry   r   r"   r?   rH   rK   rM   rP   rS   rT   r   mainr    r    r    r!   <module>   s0    $0
 R