U
    ?h                     @   st   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rpe  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Zdevicec                 S   s   | | S N abr   r   Z/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_device_func.pyadd2f   s    z,TestDeviceFunc.test_use_add2f.<locals>.add2fc                    s$   t d} | | | | | |< d S N   r   gridaryir   r   r   	use_add2f   s    
z0TestDeviceFunc.test_use_add2f.<locals>.use_add2fvoid(float32[:])
   Z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   s   | | S r   r   r   r   r   r   r   "   s    z1TestDeviceFunc.test_indirect_add2f.<locals>.add2fc                    s
    | |S r   r   r   r   r   r   indirect&   s    z4TestDeviceFunc.test_indirect_add2f.<locals>.indirectc                    s$   t d} | | | | | |< d S r   r   r   )r*   r   r   indirect_add2f*   s    
z:TestDeviceFunc.test_indirect_add2f.<locals>.indirect_add2fr   r   r   r   r    )r%   r+   r&   r'   r   r(   r   )r   r*   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%   r.   r/   r   expectr   r-   r   _check_cpu_dispatcher7   s    
z$TestDeviceFunc._check_cpu_dispatcherc                 C   s   t dd }| | d S )Nc                 S   s   | | S r   r   r   r   r   r   r.   D   s    z/TestDeviceFunc.test_cpu_dispatcher.<locals>.add)r   r5   )r%   r.   r   r   r   test_cpu_dispatcherB   s    
z"TestDeviceFunc.test_cpu_dispatcherznot supported in cudasimc              	   C   s\   t ddd }| t}| | W 5 Q R X d}t|}| |t|j	d k	 d S )Nz(i4, i4)c                 S   s   | | S r   r   r   r   r   r   r.   O   s    z7TestDeviceFunc.test_cpu_dispatcher_invalid.<locals>.addz8Untyped global name 'add':.*using cpu function on device)
r   assertRaisesr	   r5   recompiler#   searchstr	exception)r%   r.   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   s   | | S r   r   r   r   r   r   r.   [   s    z<TestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.addmymod)namec                    s"   t d} | | d| |< d S r   )r   r   r.   r   rA   r   r   r/   c   s    
zCTestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add_kernelr   r   )
r   types
ModuleTyper.   r   r!   r"   r0   r1   r2   r3   r   rC   r    test_cpu_dispatcher_other_moduleZ   s    

z/TestDeviceFunc.test_cpu_dispatcher_other_modulec                 C   sT   t jdddd }ttf}||}|jj}| d| ||}| || d S )NTr   c                 S   s   | | S r   r   xyr   r   r   fooo   s    z-TestDeviceFunc.test_inspect_llvm.<locals>.foorJ   )r   r   r   compile_devicefndescmangled_nameassertInZinspect_llvm)r%   rJ   argscresfnamellvmr   r   r   test_inspect_llvmm   s    



z TestDeviceFunc.test_inspect_llvmc                 C   sT   t jdddd }ttf}||}|jj}| d| ||}| || d S )NTr   c                 S   s   | | S r   r   rG   r   r   r   rJ      s    z,TestDeviceFunc.test_inspect_asm.<locals>.foorJ   )r   r   r   rK   rL   rM   rN   Zinspect_asm)r%   rJ   rO   rP   rQ   ptxr   r   r   test_inspect_asm~   s    



zTestDeviceFunc.test_inspect_asmc              	   C   sN   t jdddd }| t}|ttf W 5 Q R X | dt|j d S )NTr   c                 S   s   | | S r   r   rG   r   r   r   rJ      s    z8TestDeviceFunc.test_inspect_sass_disallowed.<locals>.fooz(Cannot inspect SASS of a device function)	r   r   r7   RuntimeErrorZinspect_sassr   rN   r;   r<   )r%   rJ   r=   r   r   r   test_inspect_sass_disallowed   s    

z+TestDeviceFunc.test_inspect_sass_disallowedz'cudasim will allow calling any functionc              	   C   sJ   t jdddd }| t}|d   W 5 Q R X | 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   r7   rV   rN   r;   r<   )r%   rX   r=   r   r   r   %test_device_func_as_kernel_disallowed   s    

z4TestDeviceFunc.test_device_func_as_kernel_disallowedz2cudasim ignores casting by jit decorator signaturec                    sx   t jddddd  t j fdd}t jdtjd	}t tjd
dddg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   )rH   channelsra   r   r   rgba_caller   s    z7TestDeviceFunc.test_device_casting.<locals>.rgba_callerr   r   g      ?g       @g      @g      @rY   ir   )	r   r   Zdevice_arrayr!   r   Z	to_deviceZasarrayr   assertEqual)r%   rf   rH   rd   r   re   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)rg   rB   sigrO   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 ri   )r   declare_devicer   r   rn   r%   rj   r   r   r   test_declare_device_signature   s    z,TestDeviceFunc.test_declare_device_signaturec                 C   s   t dd}| | d S )Nrj   zint32(float32[:]))r   ro   rn   rp   r   r   r   test_declare_device_string   s    z)TestDeviceFunc.test_declare_device_stringc              	   C   s2   |  td tdtd d  f W 5 Q R X d S )NReturn typerj   )assertRaisesRegex	TypeErrorr   ro   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 5 Q R X d S )Nrs   rj   z(float32[:],))rt   ru   r   ro   rv   r   r   r   test_bad_declare_device_string   s    z-TestDeviceFunc.test_bad_declare_device_stringN)__name__
__module____qualname__r)   r,   r5   r6   r   r@   rF   rS   rU   rW   rZ   rh   rn   rq   rr   rw   rx   r   r   r   r   r
      s4   








r
   __main__)r8   rD   numpyr!   Znumba.cuda.testingr   r   r   Znumbar   r   r   r   Znumba.core.errorsr	   r
   ry   mainr   r   r   r   <module>   s    S