o
    8ήc                     @   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rZe  dS 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   s2   t d}|t| k r|| ||  | |< d S d S )N   )r   gridlen)rxyi r   L/tmp/pip-target-vg8gfxp4/lib/python/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 Nr   r   r   r   r   r   add      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+   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_infoB   s   
z!TestCompileToPTX.check_debug_infoc                 C   *   dd }t |g ddd\}}| | d S )Nc                   S      d S r&   r   r   r   r   r   r   U      z;TestCompileToPTX.test_device_function_with_debug.<locals>.fT)r+   debugr   r6   r    r   r"   r#   r   r   r   test_device_function_with_debugN   s   z0TestCompileToPTX.test_device_function_with_debugc                 C   (   dd }t |g dd\}}| | d S )Nc                   S   r8   r&   r   r   r   r   r   r   ]   r9   z2TestCompileToPTX.test_kernel_with_debug.<locals>.fT)r:   r;   r<   r   r   r   test_kernel_with_debug[   s   z'TestCompileToPTX.test_kernel_with_debugc                 C   s   |  |d d S )Nr1   )r4   r5   r   r   r   check_line_infoc   s   z TestCompileToPTX.check_line_infoc                 C   r7   )Nc                   S   r8   r&   r   r   r   r   r   r   j   r9   z?TestCompileToPTX.test_device_function_with_line_info.<locals>.fT)r+   lineinfor   r@   r<   r   r   r   #test_device_function_with_line_infoi   s   z4TestCompileToPTX.test_device_function_with_line_infoc                 C   r>   )Nc                   S   r8   r&   r   r   r   r   r   r   q   r9   z6TestCompileToPTX.test_kernel_with_line_info.<locals>.fT)rA   rB   r<   r   r   r   test_kernel_with_line_infop   s   z+TestCompileToPTX.test_kernel_with_line_infoN)__name__
__module____qualname__r$   r,   r0   r6   r=   r?   r@   rC   rD   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   r%   r&   r   r'   r   r   r   r(   {   r)   zQTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_device.<locals>.addTr*   z.target sm_r   r   )	r   r	   r   get_current_devicecompute_capabilitycudadrvnvvmfind_closest_archr   )r    r(   r!   r"   r#   	device_cccctargetr   r   r   #test_compile_ptx_for_current_devicez   s   
zDTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_deviceN)rE   rF   rG   rQ   r   r   r   r   rH   x   s    rH   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|v r|d7 }qd	}| ||d
| d|  d S )Nc                 S   s   t d t |  d S )N    )r   	nanosleep)r   r   r   r   use_nanosleep   s   
z:TestCompileOnlyTests.test_nanosleep.<locals>.use_nanosleep)   r   )rO   r   
znanosleep.u32r      zGot z" nanosleep instructions, expected )r   r   splitr   )r    rU   r"   r#   nanosleep_countlineexpectedr   r   r   test_nanosleep   s   z#TestCompileOnlyTests.test_nanosleepN)rE   rF   rG   __doc__r]   r   r   r   r   rR      s    rR   __main__N)mathr   numbar   r   r   r   
numba.cudar   r	   numba.cuda.cudadrv.nvvmr
   numba.cuda.testingr   r   r   TestCaser   rH   rR   rE   mainr   r   r   r   <module>   s    n