o
    8ήc                     @   s   d dl Zd dlZd dlZd dl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mZ d dlmZ d dlmZ d dlmZmZmZmZ d	d
 ZedG dd deZedkr`e  dS dS )    N)unittest)skip_on_cudasimskip_unless_cuda_pythonskip_if_cuda_includes_missingskip_with_cuda_python)CUDATestCase)LinkerLinkerError
NvrtcError)require_context)ignore_internal_warnings)cudavoidfloat64int64c                 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 )Ng      ?
      )ranger   grid)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d5i r1   K/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudadrv/test_linker.pyfunc_with_lots_of_registers   sZ   
&&*r3   z$Linking unsupported in the simulatorc                   @   s   e Zd 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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"S )#
TestLinkerc                 C   s   t  }~dS )z9Simply go through the constructor and destructor
        N)r   new)selflinkerr1   r1   r2   test_linker_basicB   s   zTestLinker.test_linker_basicc                 C   s   t ddatjtjtd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)datazjitlink.ptxzvoid(int32[:], int32[:])linkc                 S   s&   t d}| |  t|| 7  < d S Nr   )r   r   r9   )r   yr0   r1   r1   r2   fooT   s   
z%TestLinker._test_linking.<locals>.foo{   dtypeiA  )r   r   r   i  )r   declare_devicer9   ospathjoindirname__file__jitnparrayint32
assertTrue)r6   eagerr<   argsr?   ABr1   r1   r2   _test_linkingI   s   
zTestLinker._test_linkingc                 C      | j dd d S )NFrN   rR   r6   r1   r1   r2   test_linking_lazy_compile`      z$TestLinker.test_linking_lazy_compilec                 C   rS   )NTrT   rU   rV   r1   r1   r2   test_linking_eager_compilec   rX   z%TestLinker.test_linking_eager_compilezNVIDIA Binding needed for NVRTCc                    s   t dd tjtjtdd}t j|gd fdd}tj	dtj
d	}t|}|d
 || |d }tj|| d S )Nr9   r:   r;   
jitlink.cur<   c                    s.   t d}|t| k r || | |< d S d S r=   )r   r   len)rr   r0   r9   r1   r2   kernell   s   
z*TestLinker.test_linking_cu.<locals>.kernelr   rA   )r          )r   rC   rD   rE   rF   rG   rH   rI   rJ   arangerL   
zeros_liketestingassert_array_equal)r6   r<   r_   r   r]   expectedr1   r^   r2   test_linking_cuf   s   
zTestLinker.test_linking_cuc                    s   t dd tjtjtdd}tjdd}t	  t j
d|gd fd	d
}W d    n1 s4w   Y  | t|dd | dt|d j | dt|d j d S )Nr9   r:   r;   zwarn.cuT)recordvoid(int32)r[   c                        |  d S Nr1   r   r^   r1   r2   r_         z6TestLinker.test_linking_cu_log_warning.<locals>.kernelr   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   rC   rD   rE   rF   rG   rH   warningscatch_warningsr   rI   assertEqualr\   assertInstrmessage)r6   r<   wr_   r1   r^   r2   test_linking_cu_log_warning|   s   z&TestLinker.test_linking_cu_log_warningc                    s   t dd tjtjtdd}| t}t j	d|gd fdd}W d    n1 s0w   Y  |j
jd	 }| d
| | d| | d| d S )Nr9   r:   r;   zerror.curi   r[   c                    rj   rk   r1   rl   r^   r1   r2   r_      rm   z0TestLinker.test_linking_cu_error.<locals>.kernelr   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")r   rC   rD   rE   rF   rG   rH   assertRaisesr
   rI   	exceptionrO   rq   )r6   r<   r   r_   msgr1   r^   r2   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 )NzBLinking CUDA source files is not supported with the ctypes bindingvoid()rZ   r[   c                   S      d S rk   r1   r1   r1   r1   r2   r         z8TestLinker.test_linking_cu_ctypes_unsupported.<locals>.f)assertRaisesRegexNotImplementedErrorr   rI   )r6   rx   r   r1   r1   r2   "test_linking_cu_ctypes_unsupported   s
   "z-TestLinker.test_linking_cu_ctypes_unsupportedc                 C   rz   )Nz/Don't know how to link file with extension .cuhr{   z
header.cuhr[   c                   S   r|   rk   r1   r1   r1   r1   r2   r_      r}   z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernelr~   RuntimeErrorr   rI   r6   expected_errr_   r1   r1   r2   #test_linking_unknown_filetype_error   
   "z.TestLinker.test_linking_unknown_filetype_errorc                 C   rz   )Nz-Don't know how to link file with no extensionr{   r;   r[   c                   S   r|   rk   r1   r1   r1   r1   r2   r_      r}   zDTestLinker.test_linking_file_with_no_extension_error.<locals>.kernelr   r   r1   r1   r2   )test_linking_file_with_no_extension_error   r   z4TestLinker.test_linking_file_with_no_extension_errorc                 C   s4   t jt jtdd}tjd|gddd }d S )Nr;   zcuda_include.cur{   r[   c                   S   r|   rk   r1   r1   r1   r1   r2   r_      r}   z7TestLinker.test_linking_cu_cuda_include.<locals>.kernel)rD   rE   rF   rG   rH   r   rI   )r6   r<   r_   r1   r1   r2   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 )Nzvoid(int32[::1])znonexistent.ar[   c                 S   s   d| d< d S )Nr   r1   rl   r1   r1   r2   r      rm   z2TestLinker.test_try_to_link_nonexistent.<locals>.fznonexistent.a not found)rv   r	   r   rI   rq   rw   rO   )r6   r   r   r1   r1   r2   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.r`      9   N)	r   rI   r3   
specializerJ   emptyr   assertGreaterget_regs_per_threadr6   compiledr1   r1   r2   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_registersr`   r   	r   rI   r3   r   rJ   r   r   assertLessEqualr   r   r1   r1   r2   test_set_registers_57      z TestLinker.test_set_registers_57c                 C   r   )N&   r   r`   r   r   r   r1   r1   r2   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   rI   r3   r   r   )r6   sigr   r1   r1   r2   test_set_registers_eager   s   z#TestLinker.test_set_registers_eagerN)__name__
__module____qualname__r   r8   rR   rW   rY   r   rg   ru   ry   r   r   r   r   r   r   r   r   r   r   r   r1   r1   r1   r2   r4   ?   s0    





	r4   __main__)os.pathrD   numpyrJ   rn   numba.cuda.testingr   r   r   r   r   r   numba.cuda.cudadrv.driverr   r	   r
   
numba.cudar   numba.tests.supportr   numbar   r   r   r   r3   r4   r   mainr1   r1   r1   r2   <module>   s"    0 &