o
    iE*                     @   s   d dl mZ d dlmZmZmZmZmZmZ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 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rhe  dS dS )    sqrt)cudafloat32int16int32int64uint32void)compilecompile_for_current_devicecompile_ptxcompile_ptx_for_current_device)runtime)skip_on_cudasimunittestCUDATestCasec                 C      | | S N xyr   r   b/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_compiler.pyf_module      r   z(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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 )&TestCompilec                 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)rr   r   ir   r   r   f   s   
z)TestCompile.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TestCompile.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   r   r   r   r   r   r   r   add$   r   z-TestCompile.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-   	sig_int32	sig_int16
sig_stringr   r   r   test_device_function#   s    

z TestCompile.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"   B   s   z$TestCompile.test_fastmath.<locals>.fTr0   z
fma.rn.f32z
div.rn.f32zsqrt.rn.f32)r1   fastmathzfma.rn.ftz.f32zdiv.approx.ftz.f32zsqrt.approx.ftz.f32)r   r   r'   r)   r   r   r   test_fastmathA   s   zTestCompile.test_fastmathc                 C   s   |  |d |  |d d S )Nz\.section\s+\.debug_info\.file.*test_compiler.py"assertRegexr*   r,   r   r   r   check_debug_infoU   s   zTestCompile.check_debug_infoc                 C   *   dd }t |dddd\}}| | d S )Nc                   S      d S r   r   r   r   r   r   r"   e      z6TestCompile.test_device_function_with_debug.<locals>.fr   T)r1   debugr   r>   r*   r"   r,   r-   r   r   r   test_device_function_with_debug^   s   z+TestCompile.test_device_function_with_debugc                 C   (   dd }t |ddd\}}| | d S )Nc                   S   r@   r   r   r   r   r   r   r"   m   rA   z-TestCompile.test_kernel_with_debug.<locals>.fr   T)rB   rC   rD   r   r   r   test_kernel_with_debugk   s   z"TestCompile.test_kernel_with_debugc                 C   s   |  |d d S )Nr:   r;   r=   r   r   r   check_line_infos   s   zTestCompile.check_line_infoc                 C   r?   )Nc                   S   r@   r   r   r   r   r   r   r"   z   rA   z:TestCompile.test_device_function_with_line_info.<locals>.fr   T)r1   lineinfor   rH   rD   r   r   r   #test_device_function_with_line_infoy   s   z/TestCompile.test_device_function_with_line_infoc                 C   rF   )Nc                   S   r@   r   r   r   r   r   r   r"      rA   z1TestCompile.test_kernel_with_line_info.<locals>.fr   T)rI   rJ   rD   r   r   r   test_kernel_with_line_info   s   z&TestCompile.test_kernel_with_line_infoc                 C   s\   dd }|  td t|td d d td d d f W d    d S 1 s'w   Y  d S )Nc                 S   s   | d |d  S )Nr   r   r   r   r   r   r"      s   z0TestCompile.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%TestCompile.test_non_void_return_typec                 C   L   dd }|  td t|ttfdd W d    d S 1 sw   Y  d S )Nc                 S   r   r   r   r   r   r   r   r"      r   z7TestCompile.test_c_abi_disallowed_for_kernel.<locals>.fz&The C ABI is not supported for kernelscabirM   NotImplementedErrorr   r   rO   r   r   r    test_c_abi_disallowed_for_kernel      "z,TestCompile.test_c_abi_disallowed_for_kernelc                 C   rQ   )Nc                 S   r   r   r   r   r   r   r   r"      r   z+TestCompile.test_unsupported_abi.<locals>.fzUnsupported ABI: fastcallfastcallrS   rU   rO   r   r   r   test_unsupported_abi   rX   z TestCompile.test_unsupported_abic                 C   sd   dd }t |tttddd\}}| |d | |d t |tttddd\}}| |d d S )	Nc                 S   r   r   r   r   r   r   r   r"      r   z1TestCompile.test_c_abi_device_function.<locals>.fTrR   r1   rT   param_2z=\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+f\(z&\.visible\s+\.func\s+\(\.param\s+\.b64)r   r   r&   r<   r   rD   r   r   r   test_c_abi_device_function   s   z&TestCompile.test_c_abi_device_functionc                 C   s*   t ttttddd\}}| |d d S )NTrR   r[   D\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+f_module\(r   r   r   r<   r*   r,   r-   r   r   r   'test_c_abi_device_function_module_scope   s   
z3TestCompile.test_c_abi_device_function_module_scopec                 C   s4   ddi}t ttttdd|d\}}| |d d S )Nabi_name	_Z4funciiTrR   )r1   rT   abi_infozE\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+_Z4funcii\(r_   )r*   rd   r,   r-   r   r   r   test_c_abi_with_abi_name   s
   
z$TestCompile.test_c_abi_with_abi_namec                 C   s(   t ttttdd\}}| |d d S )NTr0   r^   )r   r   r   r<   r`   r   r   r   test_compile_defaults_to_c_abi   s   z*TestCompile.test_compile_defaults_to_c_abic                 C   sf   t  dk r| d tttttddd\}}d}tj|d d dd	}| || | |t d S )
N)      z,-gen-lto unavailable in this toolkit versionTltoirr1   outputiCN   little)	byteorder)	r   get_versionskipTestr   r   r   int
from_bytesr(   )r*   ri   r-   LTOIR_MAGICheaderr   r   r   test_compile_to_ltoir   s   

z!TestCompile.test_compile_to_ltoirc                 C   sV   d}d| }|  t| tttttd|d W d    d S 1 s$w   Y  d S )NillegalzUnsupported output type: Trj   )rM   rV   r   r   r   )r*   illegal_outputmsgr   r   r   test_compile_to_invalid_error   s   
"z)TestCompile.test_compile_to_invalid_errorN)__name__
__module____qualname__r.   r5   r9   r>   rE   rG   rH   rK   rL   rP   rW   rZ   r]   ra   re   rf   ru   ry   r   r   r   r   r      s&    		
r   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )TestCompileForCurrentDevicec           	      C   s`   dd }t t f}|||dd\}}t j}tjj|}d|d  |d  }| || d S )Nc                 S   r   r   r   r   r   r   r   r/      r   zFTestCompileForCurrentDevice._check_ptx_for_current_device.<locals>.addTr0   z.target sm_r   r   )r   r   get_current_devicecompute_capabilitycudadrvnvvmfind_closest_archr'   )	r*   compile_functionr/   r+   r,   r-   	device_cccctargetr   r   r   _check_ptx_for_current_device   s   
z9TestCompileForCurrentDevice._check_ptx_for_current_devicec                 C      |  t d S r   )r   r   r*   r   r   r   #test_compile_ptx_for_current_device      z?TestCompileForCurrentDevice.test_compile_ptx_for_current_devicec                 C   r   r   )r   r   r   r   r   r   test_compile_for_current_device   r   z;TestCompileForCurrentDevice.test_compile_for_current_deviceN)rz   r{   r|   r   r   r   r   r   r   r   r}      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|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   )r   r   
znanosleep.u32r      zGot z" nanosleep instructions, expected )r   r	   splitr(   )r*   r   r,   r-   nanosleep_countlineexpectedr   r   r   test_nanosleep   s   z#TestCompileOnlyTests.test_nanosleepN)rz   r{   r|   __doc__r   r   r   r   r   r      s    r   __main__N)mathr   numbar   r   r   r   r   r	   r
   
numba.cudar   r   r   r   numba.cuda.cudadrvr   numba.cuda.testingr   r   r   r   TestCaser   r}   r   rz   mainr   r   r   r   <module>   s     $ T