
    wg                         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e      Zedk(  r ej"                          yy)    N)unittestskip_on_cudasimCUDATestCase)cudajitfloat32int32)TypingErrorc                   H   e Zd Zd Zd Zd Zd Z ed      d        Zd Z	 ed      d        Z
 ed      d	        Z ed      d
        Z ed      d        Z ed      d        Zd Z ed      d        Z ed      d        Z ed      d        Z ed      d        Zy)TestDeviceFuncc                 F   t        j                  dd      d        fd} t        j                  d      |      }d}t        j                  |t        j                        }||z   } |d	|f   |       | j                  t        j                  ||k(        ||f       y )
Nfloat32(float32, float32)Tdevicec                     | |z   S N abs     m/home/mcse/projects/flask/flask-venv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_device_func.pyadd2fz,TestDeviceFunc.test_use_add2f.<locals>.add2f       q5L    c                 T    t        j                  d      } | |   | |         | |<   y N   r   grid)aryir   s     r   	use_add2fz0TestDeviceFunc.test_use_add2f.<locals>.use_add2f   s(    		!A3q63q6*CFr   void(float32[:])
   dtyper   r   r   nparanger   
assertTrueall)selfr"   compilednelemr    expr   s         @r   test_use_add2fzTestDeviceFunc.test_use_add2f   s    	-d	;	 
<		+ 0488./	:iiRZZ0CiE3scz*S#J7r   c                    t        j                  dd      d        t        j                  dd      fd       fd} t        j                  d      |      }d}t        j                  |t        j                  	      }||z   } |d
|f   |       | j                  t        j                  ||k(        ||f       y )Nr   Tr   c                     | |z   S r   r   r   s     r   r   z1TestDeviceFunc.test_indirect_add2f.<locals>.add2f"   r   r   c                      | |      S r   r   )r   r   r   s     r   indirectz4TestDeviceFunc.test_indirect_add2f.<locals>.indirect&   s    A;r   c                 T    t        j                  d      } | |   | |         | |<   y r   r   )r    r!   r4   s     r   indirect_add2fz:TestDeviceFunc.test_indirect_add2f.<locals>.indirect_add2f*   s(    		!Ac!fc!f-CFr   r#   r$   r%   r   r'   )r,   r6   r-   r.   r    r/   r   r4   s         @@r   test_indirect_add2fz"TestDeviceFunc.test_indirect_add2f    s    	-d	;	 
<	 
-d	;	 
<		. 0488./?iiRZZ0CiE3scz*S#J7r   c                     t         j                  fd       }t        j                  d      }|dz   } |d|j                  f   |       t        j
                  j                  ||       y )Nc                 N    t        j                  d      } | |   d      | |<   y r   r   )r    r!   adds     r   
add_kernelz8TestDeviceFunc._check_cpu_dispatcher.<locals>.add_kernel8   s#    		!AQ^CFr   r$   r   )r   r   r(   r)   sizetestingassert_equal)r,   r:   r;   r    expects    `   r   _check_cpu_dispatcherz$TestDeviceFunc._check_cpu_dispatcher7   s\    		$ 
	$ iimq
1chh;$


,r   c                 >    t         d        }| j                  |       y )Nc                     | |z   S r   r   r   s     r   r:   z/TestDeviceFunc.test_cpu_dispatcher.<locals>.addD   r   r   )r   r@   )r,   r:   s     r   test_cpu_dispatcherz"TestDeviceFunc.test_cpu_dispatcherB   s$    		 
	 	""3'r   znot supported in cudasimc                 4   t        d      d        }| j                  t              5 }| j                  |       d d d        d}t	        j
                  |      }| j                  |j                  t        j                              d u       y # 1 sw Y   VxY w)Nz(i4, i4)c                     | |z   S r   r   r   s     r   r:   z7TestDeviceFunc.test_cpu_dispatcher_invalid.<locals>.addO   r   r   z8Untyped global name 'add':.*using cpu function on device)
r   assertRaisesr
   r@   recompiler*   searchstr	exception)r,   r:   raisesmsgexpecteds        r   test_cpu_dispatcher_invalidz*TestDeviceFunc.test_cpu_dispatcher_invalidJ   s    
 
Z	 
	 {+ 	,v&&s+	,H::c?F,<,<(=>dJK		, 	,s   BBc                 ,   t         d        }t        j                  d      |_        ~t        j                   fd       }t        j                  d      }|dz   } |d|j                  f   |       t
        j                  j                  ||       y )Nc                     | |z   S r   r   r   s     r   r:   z<TestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add[   r   r   mymod)namec                 `    t        j                  d      }j                  | |   d      | |<   y r   )r   r   r:   )r    r!   rR   s     r   r;   zCTestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add_kernelc   s(    		!AYYs1vq)CFr   r$   r   )
r   types
ModuleTyper:   r   r(   r)   r<   r=   r>   )r,   r:   r;   r    r?   rR   s        @r    test_cpu_dispatcher_other_modulez/TestDeviceFunc.test_cpu_dispatcher_other_moduleZ   s    		 
	   g.			* 
	* iimq
1chh;$


,r   c                    t        j                  d      d        }t        t        f}|j                  |      }|j                  j
                  }| j                  d|       |j                  |      }| j                  ||       y )NTr   c                     | |z   S r   r   xys     r   fooz-TestDeviceFunc.test_inspect_llvm.<locals>.fooo   r   r   r]   )r   r   r	   compile_devicefndescmangled_nameassertIninspect_llvm)r,   r]   argscresfnamellvms         r   test_inspect_llvmz TestDeviceFunc.test_inspect_llvmm   ss    			 
	 u~!!$'((eU#%eT"r   c                    t        j                  d      d        }t        t        f}|j                  |      }|j                  j
                  }| j                  d|       |j                  |      }| j                  ||       y )NTr   c                     | |z   S r   r   rZ   s     r   r]   z,TestDeviceFunc.test_inspect_asm.<locals>.foo   r   r   r]   )r   r   r	   r^   r_   r`   ra   inspect_asm)r,   r]   rc   rd   re   ptxs         r   test_inspect_asmzTestDeviceFunc.test_inspect_asm~   sq    			 
	 u~!!$'((eU#ood#eS!r   c                    t        j                  d      d        }| j                  t              5 }|j	                  t
        t
        f       d d d        | j                  dt        j                               y # 1 sw Y   /xY w)NTr   c                     | |z   S r   r   rZ   s     r   r]   z8TestDeviceFunc.test_inspect_sass_disallowed.<locals>.foo   r   r   z(Cannot inspect SASS of a device function)	r   r   rF   RuntimeErrorinspect_sassr	   ra   rJ   rK   )r,   r]   rL   s      r   test_inspect_sass_disallowedz+TestDeviceFunc.test_inspect_sass_disallowed   su    			 
	 |, 	-eU^,	- 	@&**+	-	- 	-s   A<<Bz'cudasim will allow calling any functionc                     t        j                  d      d        }| j                  t              5 } |d           d d d        | j	                  dt        j                               y # 1 sw Y   /xY w)NTr   c                       y r   r   r   r   r   fz?TestDeviceFunc.test_device_func_as_kernel_disallowed.<locals>.f   s    r   r   r   z,Cannot compile a device function as a kernel)r   r   rF   ro   ra   rJ   rK   )r,   rt   rL   s      r   %test_device_func_as_kernel_disallowedz4TestDeviceFunc.test_device_func_as_kernel_disallowed   sn    			 
	 |, 	AdGI	 	D&**+	-	 	s   A++A4z2cudasim ignores casting by jit decorator signaturec                 p   t        j                  dd      d        t         j                  fd       }t        j                  dt        j                        }t        j
                  t        j                  g dt        j                              } |d	   ||       | j                  d
|d          y )Nz!int32(int32, int32, int32, int32)Tr   c                 H    | dz  dz  |dz  dz  z  |dz  dz  z  |dz  dz  z  S )N         r      r   )rgr   r   s       r   rgbaz0TestDeviceFunc.test_device_casting.<locals>.rgba   sF    $h2%$h1_&$h1_& $h2%' (r   c                 :     |d   |d   |d   |d         | d<   y )Nr   r         r   )r[   channelsr   s     r   rgba_callerz7TestDeviceFunc.test_device_casting.<locals>.rgba_caller   s'    Xa[(1+x{KAaDr   r   r%   )g      ?g       @g      @g      @ru   ir   )	r   r   device_arrayr(   r	   	to_deviceasarrayr   assertEqual)r,   r   r[   r   r   s       @r   test_device_castingz"TestDeviceFunc.test_device_casting   s     
5d	C	( 
D	( 
	L 
	L arxx0>>"**-A35::#? @ 	D!X&QqT*r   c                     | j                  |j                  d       | j                  |j                  j                  t        d d  f       | j                  |j                  j
                  t               y Nf1)r   rS   sigrc   r   return_typer	   )r,   decls     r   _test_declare_devicez#TestDeviceFunc._test_declare_device   sN    D)6--u5r   z!cudasim does not check signaturesc                 r    t        j                  dt        t        d d              }| j	                  |       y r   )r   declare_devicer	   r   r   r,   r   s     r   test_declare_device_signaturez,TestDeviceFunc.test_declare_device_signature   s+      uWQZ'89!!"%r   c                 R    t        j                  dd      }| j                  |       y )Nr   zint32(float32[:]))r   r   r   r   s     r   test_declare_device_stringz)TestDeviceFunc.test_declare_device_string   s#      ':;!!"%r   c                     | j                  t        d      5  t        j                  dt        d d  f       d d d        y # 1 sw Y   y xY w)NReturn typer   )assertRaisesRegex	TypeErrorr   r   r   r,   s    r   test_bad_declare_device_tuplez,TestDeviceFunc.test_bad_declare_device_tuple   s=    ##I}= 	5wqzm4	5 	5 	5s	   ?Ac                     | j                  t        d      5  t        j                  dd       d d d        y # 1 sw Y   y xY w)Nr   r   z(float32[:],))r   r   r   r   r   s    r   test_bad_declare_device_stringz-TestDeviceFunc.test_bad_declare_device_string   s7    ##I}= 	7o6	7 	7 	7s	   7A N)__name__
__module____qualname__r0   r7   r@   rC   r   rO   rW   rg   rl   rq   rv   r   r   r   r   r   r   r   r   r   r   r      s(   8&8.	-( /0L 1L-& /0# 1#  /0" 1"  /0	- 1	- >?	- @	- IJ+ K+66
 89& :& 89& :& 895 :5 897 :7r   r   __main__)rG   rU   numpyr(   numba.cuda.testingr   r   r   numbar   r   r   r	   numba.core.errorsr
   r   r   mainr   r   r   <module>r      sF    	   F F + + )O7\ O7d zHMMO r   