U
    ,d                     @   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re  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 ]}||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 )Ng      ?
      )ranger   grid)xabcdefZa1Za2a3Za4Za5Zb1Zb2Zb3Zb4Zb5Zc1c2c3Zc4Zc5d1Zd2Zd3Zd4Zd5i r!   H/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudadrv/test_linker.pyfunc_with_lots_of_registers   sZ    
&&r#   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)selfZlinkerr!   r!   r"   test_linker_basicB   s    zTestLinker.test_linker_basicc                 C   s   t ddatjtjtdd}|r0d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   r(   )r   yr    r!   r!   r"   fooT   s    
z%TestLinker._test_linking.<locals>.foo{   ZdtypeiA  )r   r   r   i  )r   declare_devicer(   ospathjoindirname__file__jitnparrayint32
assertTrue)r&   eagerr+   argsr.   ABr!   r!   r"   _test_linkingI   s    
zTestLinker._test_linkingc                 C   s   | j dd d S )NFr<   r@   r&   r!   r!   r"   test_linking_lazy_compile`   s    z$TestLinker.test_linking_lazy_compilec                 C   s   | j dd d S )NTrA   rB   rC   r!   r!   r"   test_linking_eager_compilec   s    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 )Nr(   r)   r*   
jitlink.cur+   c                    s*   t d}|t| k r& || | |< d S r,   )r   r   len)rr   r    r(   r!   r"   kernell   s    
z*TestLinker.test_linking_cu.<locals>.kernelr   r0   )r          )r   r1   r2   r3   r4   r5   r6   r7   r8   Zaranger:   Z
zeros_likeZtestingZassert_array_equal)r&   r+   rK   r   rI   expectedr!   rJ   r"   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 5 Q R X | t|dd | dt|d j | dt|d j d S )Nr(   r)   r*   zwarn.cuT)recordvoid(int32)rG   c                    s    |  d S Nr!   r   rJ   r!   r"   rK      s    z6TestLinker.test_linking_cu_log_warning.<locals>.kernelr   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   r1   r2   r3   r4   r5   r6   warningscatch_warningsr   r7   assertEqualrH   assertInstrmessage)r&   r+   wrK   r!   rJ   r"   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 5 Q R X |j
jd	 }| d
| | d| | d| d S )Nr(   r)   r*   zerror.curQ   rG   c                    s    |  d S rR   r!   rS   rJ   r!   r"   rK      s    z0TestLinker.test_linking_cu_error.<locals>.kernelr   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")r   r1   r2   r3   r4   r5   r6   assertRaisesr
   r7   	exceptionr=   rW   )r&   r+   r   rK   msgr!   rJ   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 )NzBLinking CUDA source files is not supported with the ctypes bindingvoid()rF   rG   c                   S   s   d S rR   r!   r!   r!   r!   r"   r      s    z8TestLinker.test_linking_cu_ctypes_unsupported.<locals>.f)assertRaisesRegexNotImplementedErrorr   r7   )r&   r^   r   r!   r!   r"   "test_linking_cu_ctypes_unsupported   s    z-TestLinker.test_linking_cu_ctypes_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.cuhrG   c                   S   s   d S rR   r!   r!   r!   r!   r"   rK      s    z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernelra   RuntimeErrorr   r7   r&   Zexpected_errrK   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`   r*   rG   c                   S   s   d S rR   r!   r!   r!   r!   r"   rK      s    zDTestLinker.test_linking_file_with_no_extension_error.<locals>.kernelrd   rf   r!   r!   r"   )test_linking_file_with_no_extension_error   s    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`   rG   c                   S   s   d S rR   r!   r!   r!   r!   r"   rK      s    z7TestLinker.test_linking_cu_cuda_include.<locals>.kernel)r2   r3   r4   r5   r6   r   r7   )r&   r+   rK   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 )Nzvoid(int32[::1])znonexistent.arG   c                 S   s   d| d< d S )Nr   r!   rS   r!   r!   r"   r      s    z2TestLinker.test_try_to_link_nonexistent.<locals>.fznonexistent.a not found)r\   r	   r   r7   rW   r]   r=   )r&   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.rL      9   N)	r   r7   r#   
specializer8   emptyr   ZassertGreaterget_regs_per_threadr&   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 )Nrl   Zmax_registersrL   rk   	r   r7   r#   rm   r8   rn   r   assertLessEqualro   rp   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&   rs   rL   rk   rt   rp   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   rw   rs   )r   r   r   r   r7   r#   ru   ro   )r&   sigrq   r!   r!   r"   test_set_registers_eager   s    z#TestLinker.test_set_registers_eagerN)__name__
__module____qualname__r   r'   r@   rD   rE   r   rO   r[   r_   r   rc   rg   rh   r   ri   rj   rr   rv   rx   rz   r!   r!   r!   r"   r$   ?   s.   





	r$   __main__)os.pathr2   Znumpyr8   rT   Znumba.cuda.testingr   r   r   r   r   r   Znumba.cuda.cudadrv.driverr   r	   r
   Z
numba.cudar   Znumba.tests.supportr   Znumbar   r   r   r   r#   r$   r{   mainr!   r!   r!   r"   <module>   s   0 &