U
    ,d                     @   s   d dl mZ d dl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re  dS )    sqrt)cudafloat32uint32void)compile_ptxcompile_ptx_for_current_device)NVVM)skip_on_cudasimunittestCUDATestCasez(Compilation unsupported in the simulatorc                   @   sT   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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   Zgridlen)rxyi r   I/tmp/pip-unpacked-wheel-eu7e0c37/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   sV   dd }t t f}t||dd\}}| d| | 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   )r   r   r   r   r   )r   r&   r    r!   r"   r   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
   Z	is_nvvm70ZskipTestassertRegexr   r!   r   r   r   check_debug_infoB   s    
z!TestCompileToPTX.check_debug_infoc                 C   s*   dd }t |g ddd\}}| | d S )Nc                   S   s   d S r$   r   r   r   r   r   r   U   s    z;TestCompileToPTX.test_device_function_with_debug.<locals>.fT)r(   debugr   r0   r   r   r!   r"   r   r   r   test_device_function_with_debugN   s    z0TestCompileToPTX.test_device_function_with_debugc                 C   s(   dd }t |g dd\}}| | d S )Nc                   S   s   d S r$   r   r   r   r   r   r   ]   s    z2TestCompileToPTX.test_kernel_with_debug.<locals>.fT)r1   r2   r3   r   r   r   test_kernel_with_debug[   s    z'TestCompileToPTX.test_kernel_with_debugc                 C   s   |  |d d S )Nr-   )r.   r/   r   r   r   check_line_infoc   s    z TestCompileToPTX.check_line_infoc                 C   s*   dd }t |g ddd\}}| | d S )Nc                   S   s   d S r$   r   r   r   r   r   r   j   s    z?TestCompileToPTX.test_device_function_with_line_info.<locals>.fT)r(   lineinfor   r6   r3   r   r   r   #test_device_function_with_line_infoi   s    z4TestCompileToPTX.test_device_function_with_line_infoc                 C   s(   dd }t |g dd\}}| | d S )Nc                   S   s   d S r$   r   r   r   r   r   r   q   s    z6TestCompileToPTX.test_kernel_with_line_info.<locals>.fT)r7   r8   r3   r   r   r   test_kernel_with_line_infop   s    z+TestCompileToPTX.test_kernel_with_line_infoN)__name__
__module____qualname__r#   r)   r,   r0   r4   r5   r6   r9   r:   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_devicez   s    
zDTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_deviceN)r;   r<   r=   rA   r   r   r   r   r>   x   s   r>   c                   @   s   e Zd ZdZdd ZdS )TestCompileOnlyTestszFor 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   )r?   r   
znanosleep.u32r      zGot z" nanosleep instructions, expected )r   r   splitr   )r   rD   r!   r"   Znanosleep_countlineexpectedr   r   r   test_nanosleep   s    
z#TestCompileOnlyTests.test_nanosleepN)r;   r<   r=   __doc__rK   r   r   r   r   rB      s   rB   __main__N)mathr   Znumbar   r   r   r   Z
numba.cudar   r	   Znumba.cuda.cudadrv.nvvmr
   Znumba.cuda.testingr   r   r   ZTestCaser   r>   rB   r;   mainr   r   r   r   <module>   s   n