o
    ۾i                     @   sx   d dl Z d dlZd dlZd dlmZmZmZ d dlm	Z	m
Z
mZmZ d dlmZ G dd deZedkr:e  dS dS )    N)unittestskip_on_cudasimCUDATestCase)cudajitfloat32int32)TypingErrorc                   @   s   e Zd Zdd Zdd Zdd Zdd Zed	d
d Zdd Z	ed	dd Z
ed	dd Zed	dd Zeddd Zeddd Zdd Zeddd Zeddd  Zedd!d" Zedd#d$ Zd%S )&TestDeviceFuncc                    sz   t jddddd   fdd}t d|}d	}tj|tjd
}|| }|d|f | | t||k||f d S )Nfloat32(float32, float32)Tdevicec                 S      | | S N abr   r   \/home/ubuntu/.local/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_device_func.pyadd2f      z,TestDeviceFunc.test_use_add2f.<locals>.add2fc                    $   t d} | | | | | |< d S N   r   gridaryir   r   r   	use_add2f      
z0TestDeviceFunc.test_use_add2f.<locals>.use_add2fvoid(float32[:])
   dtyper   r   r   nparanger   
assertTrueall)selfr    compilednelemr   expr   r   r   test_use_add2f   s   
zTestDeviceFunc.test_use_add2fc                    s   t jddddd  t jddd fddfdd	}t d
|}d}tj|tjd}|| }|d|f | | t||k||f d S )Nr   Tr   c                 S   r   r   r   r   r   r   r   r   "   r   z1TestDeviceFunc.test_indirect_add2f.<locals>.add2fc                    s
    | |S r   r   r   r   r   r   indirect&   s   
z4TestDeviceFunc.test_indirect_add2f.<locals>.indirectc                    r   r   r   r   )r0   r   r   indirect_add2f*   r!   z:TestDeviceFunc.test_indirect_add2f.<locals>.indirect_add2fr"   r#   r$   r   r&   )r+   r1   r,   r-   r   r.   r   )r   r0   r   test_indirect_add2f    s   
z"TestDeviceFunc.test_indirect_add2fc                    sH   t j fdd}td}|d }|d|jf | tj|| d S )Nc                    s    t d} | | d| |< d S r   r   r   addr   r   
add_kernel8   s   
z8TestDeviceFunc._check_cpu_dispatcher.<locals>.add_kernelr#   r   )r   r   r'   r(   sizetestingassert_equalr+   r4   r5   r   expectr   r3   r   _check_cpu_dispatcher7   s   
z$TestDeviceFunc._check_cpu_dispatcherc                 C   s   t dd }| | d S )Nc                 S   r   r   r   r   r   r   r   r4   D   r   z/TestDeviceFunc.test_cpu_dispatcher.<locals>.add)r   r;   )r+   r4   r   r   r   test_cpu_dispatcherB   s   
z"TestDeviceFunc.test_cpu_dispatcherznot supported in cudasimc                 C   sp   t ddd }| t}| | W d    n1 sw   Y  d}t|}| |t|j	d u d S )Nz(i4, i4)c                 S   r   r   r   r   r   r   r   r4   O   r   z7TestDeviceFunc.test_cpu_dispatcher_invalid.<locals>.addz8Untyped global name 'add':.*using cpu function on device)
r   assertRaisesr	   r;   recompiler)   searchstr	exception)r+   r4   raisesmsgexpectedr   r   r   test_cpu_dispatcher_invalidJ   s   

z*TestDeviceFunc.test_cpu_dispatcher_invalidc                    sh   t dd }tjdd | _~tj  fdd}td}|d }|d|jf | tj	|| d S )	Nc                 S   r   r   r   r   r   r   r   r4   [   r   z<TestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.addmymod)namec                    s"   t d} | | d| |< d S r   )r   r   r4   r   rG   r   r   r5   c   s   
zCTestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add_kernelr#   r   )
r   types
ModuleTyper4   r   r'   r(   r6   r7   r8   r9   r   rI   r    test_cpu_dispatcher_other_moduleZ   s   

z/TestDeviceFunc.test_cpu_dispatcher_other_modulec                 C   T   t jdddd }ttf}||}|jj}| d| ||}| || d S )NTr   c                 S   r   r   r   xyr   r   r   fooo   r   z-TestDeviceFunc.test_inspect_llvm.<locals>.foorQ   )r   r   r   compile_devicefndescmangled_nameassertIninspect_llvm)r+   rQ   argscresfnamellvmr   r   r   test_inspect_llvmm      



z TestDeviceFunc.test_inspect_llvmc                 C   rM   )NTr   c                 S   r   r   r   rN   r   r   r   rQ      r   z,TestDeviceFunc.test_inspect_asm.<locals>.foorQ   )r   r   r   rR   rS   rT   rU   inspect_asm)r+   rQ   rW   rX   rY   ptxr   r   r   test_inspect_asm~   r\   zTestDeviceFunc.test_inspect_asmc                 C   sb   t jdddd }| t}|ttf W d    n1 s!w   Y  | dt|j d S )NTr   c                 S   r   r   r   rN   r   r   r   rQ      r   z8TestDeviceFunc.test_inspect_sass_disallowed.<locals>.fooz(Cannot inspect SASS of a device function)	r   r   r=   RuntimeErrorinspect_sassr   rU   rA   rB   )r+   rQ   rC   r   r   r   test_inspect_sass_disallowed   s   

z+TestDeviceFunc.test_inspect_sass_disallowedz'cudasim will allow calling any functionc                 C   s^   t jdddd }| t}|d   W d    n1 sw   Y  | dt|j d S )NTr   c                   S   s   d S r   r   r   r   r   r   f   s   z?TestDeviceFunc.test_device_func_as_kernel_disallowed.<locals>.fr   r   z,Cannot compile a device function as a kernel)r   r   r=   r`   rU   rA   rB   )r+   rc   rC   r   r   r   %test_device_func_as_kernel_disallowed   s   

z4TestDeviceFunc.test_device_func_as_kernel_disallowedz2cudasim ignores casting by jit decorator signaturec                    st   t jddddd  t j fdd}t jdtjd	}t tjg d
tjd	}|d || | d|d  d S )Nz!int32(int32, int32, int32, int32)Tr   c                 S   s0   | d@ d> |d@ d> B |d@ d> B |d@ d> B S )N         r      r   )rgr   r   r   r   r   rgba   s   



z0TestDeviceFunc.test_device_casting.<locals>.rgbac                    s&    |d |d |d |d | d< d S )Nr   r         r   )rO   channelsrl   r   r   rgba_caller   s   &z7TestDeviceFunc.test_device_casting.<locals>.rgba_callerr   r$   )g      ?g       @g      @g      @rd   ir   )	r   r   device_arrayr'   r   	to_deviceasarrayr   assertEqual)r+   rq   rO   ro   r   rp   r   test_device_casting   s   	
z"TestDeviceFunc.test_device_castingc                 C   s<   |  |jd |  |jjtd d  f |  |jjt d S Nf1)ru   rH   sigrW   r   return_typer   )r+   declr   r   r   _test_declare_device   s   z#TestDeviceFunc._test_declare_devicez!cudasim does not check signaturesc                 C   s&   t dttd d  }| | d S rw   )r   declare_devicer   r   r|   r+   rx   r   r   r   test_declare_device_signature   s   z,TestDeviceFunc.test_declare_device_signaturec                 C   s   t dd}| | d S )Nrx   zint32(float32[:]))r   r}   r|   r~   r   r   r   test_declare_device_string   s   z)TestDeviceFunc.test_declare_device_stringc                 C   sH   |  td tdtd d  f W d    d S 1 sw   Y  d S )NReturn typerx   )assertRaisesRegex	TypeErrorr   r}   r   r+   r   r   r   test_bad_declare_device_tuple   s   "z,TestDeviceFunc.test_bad_declare_device_tuplec                 C   s>   |  td tdd W d    d S 1 sw   Y  d S )Nr   rx   z(float32[:],))r   r   r   r}   r   r   r   r   test_bad_declare_device_string   s   "z-TestDeviceFunc.test_bad_declare_device_stringN)__name__
__module____qualname__r/   r2   r;   r<   r   rF   rL   r[   r_   rb   re   rv   r|   r   r   r   r   r   r   r   r   r
      s6    








r
   __main__)r>   rJ   numpyr'   numba.cuda.testingr   r   r   numbar   r   r   r   numba.core.errorsr	   r
   r   mainr   r   r   r   <module>   s     S