U
    hâËdW  ã                   @   s¼   d dl mZ d dlmZmZmZmZmZmZ d dl	m
Z
mZ d dlmZ d dlmZmZmZ edƒG dd„ dejƒƒZedƒG d	d
„ d
eƒƒZedƒG dd„ dejƒƒZedkr¸e ¡  dS )é    ©Úsqrt)ÚcudaÚfloat32Úint16Úint32Úuint32Úvoid)Úcompile_ptxÚcompile_ptx_for_current_device)ÚNVVM)Úskip_on_cudasimÚunittestÚCUDATestCasez(Compilation unsupported in the simulatorc                   @   s\   e Z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S )ÚTestCompileToPTXc                 C   sl   dd„ }t d d … t d d … t d d … f}t||ƒ\}}|  d|¡ |  d|¡ |  d|¡ |  |t¡ d S )Nc                 S   s.   t  d¡}|t| ƒk r*|| ||  | |< d S )Né   )r   ÚgridÚlen)ÚrÚxÚyÚi© r   ú^/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_compiler.pyÚf   s    
z.TestCompileToPTX.test_global_kernel.<locals>.fÚfunc_retvalú.visible .funcú.visible .entry)r   r
   ÚassertNotInÚassertInÚassertEqualr	   ©Úselfr   ÚargsÚptxÚrestyr   r   r   Útest_global_kernel   s    "z#TestCompileToPTX.test_global_kernelc                 C   sÈ   dd„ }t t f}t||dd�\}}|  d|¡ |  d|¡ |  d|¡ |  |t ¡ tttƒ}t||dd�\}}|  |t¡ tttƒ}t||dd�\}}|  |t¡ d}t||dd�\}}|  |t¡ d S )	Nc                 S   s   | | S ©Nr   ©r   r   r   r   r   Úadd   s    z2TestCompileToPTX.test_device_function.<locals>.addT©Údevicer   r   r   zuint32(uint32, uint32))r   r
   r   r   r    r   r   r   )r"   r)   r#   r$   r%   Z	sig_int32Z	sig_int16Z
sig_stringr   r   r   Útest_device_function   s     

z%TestCompileToPTX.test_device_functionc                 C   s†   dd„ }t t t t f}t||dd�\}}|  d|¡ |  d|¡ |  d|¡ t||ddd�\}}|  d	|¡ |  d
|¡ |  d|¡ d S )Nc                 S   s   t | | | | ƒS r'   r   )r   r   ÚzÚdr   r   r   r   <   s    z)TestCompileToPTX.test_fastmath.<locals>.fTr*   z
fma.rn.f32z
div.rn.f32zsqrt.rn.f32)r+   Zfastmathzfma.rn.ftz.f32zdiv.approx.ftz.f32zsqrt.approx.ftz.f32)r   r
   r   r!   r   r   r   Útest_fastmath;   s    zTestCompileToPTX.test_fastmathc                 C   s.   t ƒ js|  d¡ |  |d¡ |  |d¡ d S )Nz$debuginfo not generated for NVVM 3.4z\.section\s+\.debug_infoú\.file.*test_compiler.py")r   Ú	is_nvvm70ÚskipTestÚassertRegex©r"   r$   r   r   r   Úcheck_debug_infoO   s    
z!TestCompileToPTX.check_debug_infoc                 C   s*   dd„ }t |dddd�\}}|  |¡ d S )Nc                   S   s   d S r'   r   r   r   r   r   r   b   s    z;TestCompileToPTX.test_device_function_with_debug.<locals>.fr   T)r+   Údebug©r
   r5   ©r"   r   r$   r%   r   r   r   Útest_device_function_with_debug[   s    z0TestCompileToPTX.test_device_function_with_debugc                 C   s(   dd„ }t |ddd�\}}|  |¡ d S )Nc                   S   s   d S r'   r   r   r   r   r   r   j   s    z2TestCompileToPTX.test_kernel_with_debug.<locals>.fr   T)r6   r7   r8   r   r   r   Útest_kernel_with_debugh   s    z'TestCompileToPTX.test_kernel_with_debugc                 C   s   |   |d¡ d S )Nr0   )r3   r4   r   r   r   Úcheck_line_infop   s    z TestCompileToPTX.check_line_infoc                 C   s<   t ƒ js|  d¡ dd„ }t|dddd�\}}|  |¡ d S )Nú#lineinfo not generated for NVVM 3.4c                   S   s   d S r'   r   r   r   r   r   r   z   s    z?TestCompileToPTX.test_device_function_with_line_info.<locals>.fr   T)r+   Úlineinfo©r   r1   r2   r
   r;   r8   r   r   r   Ú#test_device_function_with_line_infov   s
    
z4TestCompileToPTX.test_device_function_with_line_infoc                 C   s:   t ƒ js|  d¡ dd„ }t|ddd�\}}|  |¡ d S )Nr<   c                   S   s   d S r'   r   r   r   r   r   r   „   s    z6TestCompileToPTX.test_kernel_with_line_info.<locals>.fr   T)r=   r>   r8   r   r   r   Útest_kernel_with_line_info€   s
    
z+TestCompileToPTX.test_kernel_with_line_infoc              	   C   sF   dd„ }|   td¡�( t|td d d… td d d… fƒ W 5 Q R X d S )Nc                 S   s   | d |d  S )Nr   r   r(   r   r   r   r   ‹   s    z5TestCompileToPTX.test_non_void_return_type.<locals>.fzmust have void return typer   )ÚassertRaisesRegexÚ	TypeErrorr
   r   )r"   r   r   r   r   Útest_non_void_return_typeŠ   s    z*TestCompileToPTX.test_non_void_return_typeN)Ú__name__Ú
__module__Ú__qualname__r&   r,   r/   r5   r9   r:   r;   r?   r@   rC   r   r   r   r   r   	   s   

r   c                   @   s   e Zd Zdd„ ZdS )Ú TestCompileToPTXForCurrentDevicec                 C   s`   dd„ }t t f}t||dd�\}}t ¡ j}tjj |¡}d|d › |d › �}|  ||¡ d S )Nc                 S   s   | | S r'   r   r(   r   r   r   r)   •   s    zQTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_device.<locals>.addTr*   z.target sm_r   r   )	r   r   r   Zget_current_deviceZcompute_capabilityZcudadrvZnvvmZfind_closest_archr   )r"   r)   r#   r$   r%   Z	device_ccÚccÚtargetr   r   r   Ú#test_compile_ptx_for_current_device”   s    
zDTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_deviceN)rD   rE   rF   rJ   r   r   r   r   rG   ’   s   rG   c                   @   s   e Zd ZdZdd„ ZdS )ÚTestCompileOnlyTestsz‚For tests where we can only check correctness by examining the compiler
    output rather than observing the effects of execution.c                 C   sb   dd„ }t |tfdd�\}}d}| d¡D ]}d|kr*|d7 }q*d	}|  ||d
|› d|› �¡ d S )Nc                 S   s   t  d¡ t  | ¡ d S )Né    )r   Z	nanosleep)r   r   r   r   Úuse_nanosleep©   s    
z:TestCompileOnlyTests.test_nanosleep.<locals>.use_nanosleep)é   r   )rH   r   Ú
znanosleep.u32r   é   zGot z" nanosleep instructions, expected )r
   r   Úsplitr    )r"   rM   r$   r%   Znanosleep_countÚlineÚexpectedr   r   r   Útest_nanosleep¨   s    
ÿz#TestCompileOnlyTests.test_nanosleepN)rD   rE   rF   Ú__doc__rT   r   r   r   r   rK   £   s   rK   Ú__main__N)Úmathr   Znumbar   r   r   r   r   r	   Z
numba.cudar
   r   Znumba.cuda.cudadrv.nvvmr   Znumba.cuda.testingr   r   r   ZTestCaser   rG   rK   rD   Úmainr   r   r   r   Ú<module>   s     	