
    sh                         S SK r S SKrS SKrS SKJrJrJr  S SKJ	r	J
r
JrJr  S SKJr   " S S\5      r\S:X  a  \R"                  " 5         gg)    N)unittestskip_on_cudasimCUDATestCase)cudajitfloat32int32)TypingErrorc                   P   \ rS rSrS rS rS rS r\" S5      S 5       r	S r
\" S5      S	 5       r\" S5      S
 5       r\" S5      S 5       r\" S5      S 5       r\" S5      S 5       rS r\" S5      S 5       r\" S5      S 5       r\" S5      S 5       r\" S5      S 5       rSrg)TestDeviceFunc   c                 >  ^ [         R                  " SSS9S 5       mU4S jn[         R                  " S5      " U5      nSn[        R                  " U[        R                  S9nXD-   nUS	U4   " U5        U R                  [        R                  " XE:H  5      XE45        g )
Nfloat32(float32, float32)Tdevicec                 
    X-   $ N abs     |/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_device_func.pyadd2f,TestDeviceFunc.test_use_add2f.<locals>.add2f   	    5L    c                 P   > [         R                  " S5      nT" X   X   5      X'   g N   r   grid)aryir   s     r   	use_add2f0TestDeviceFunc.test_use_add2f.<locals>.use_add2f   s"    		!A3636*CFr   void(float32[:])
   dtyper   r   r   nparanger   
assertTrueall)selfr$   compilednelemr"   expr   s         @r   test_use_add2fTestDeviceFunc.test_use_add2f   s    	-d	;	 
<		+ 88./	:iiRZZ0iE3sz*SJ7r   c                 |  ^^ [         R                  " SSS9S 5       m[         R                  " SSS9U4S j5       mU4S jn[         R                  " S5      " U5      nSn[        R                  " U[        R                  S	9nXD-   nUS
U4   " U5        U R                  [        R                  " XE:H  5      XE45        g )Nr   Tr   c                 
    X-   $ r   r   r   s     r   r   1TestDeviceFunc.test_indirect_add2f.<locals>.add2f"   r   r   c                    > T" X5      $ r   r   )r   r   r   s     r   indirect4TestDeviceFunc.test_indirect_add2f.<locals>.indirect&   s    ;r   c                 P   > [         R                  " S5      nT" X   X   5      X'   g r   r    )r"   r#   r9   s     r   indirect_add2f:TestDeviceFunc.test_indirect_add2f.<locals>.indirect_add2f*   s"    		!Acfcf-CFr   r&   r'   r(   r   r*   )r/   r<   r0   r1   r"   r2   r   r9   s         @@r   test_indirect_add2f"TestDeviceFunc.test_indirect_add2f    s    	-d	;	 
<	 
-d	;	 
<		. 88./?iiRZZ0iE3sz*SJ7r   c                    ^ [         R                  U4S j5       n[        R                  " S5      nUS-   nUSUR                  4   " U5        [        R
                  R                  XC5        g )Nc                 L   > [         R                  " S5      nT" X   S5      X'   g r   r    )r"   r#   adds     r   
add_kernel8TestDeviceFunc._check_cpu_dispatcher.<locals>.add_kernel8   s    		!A^CFr   r'   r   )r   r   r+   r,   sizetestingassert_equal)r/   rB   rC   r"   expects    `   r   _check_cpu_dispatcher$TestDeviceFunc._check_cpu_dispatcher7   sX    		$ 
	$ iimq1chh;$


,r   c                 >    [         S 5       nU R                  U5        g )Nc                 
    X-   $ r   r   r   s     r   rB   /TestDeviceFunc.test_cpu_dispatcher.<locals>.addD   r   r   )r   rI   )r/   rB   s     r   test_cpu_dispatcher"TestDeviceFunc.test_cpu_dispatcherB   s$    		 
	 	""3'r   znot supported in cudasimc                 @   [        S5      S 5       nU R                  [        5       nU R                  U5        S S S 5        Sn[        R
                  " U5      nU R                  UR                  [        WR                  5      5      S L5        g ! , (       d  f       N\= f)Nz(i4, i4)c                 
    X-   $ r   r   r   s     r   rB   7TestDeviceFunc.test_cpu_dispatcher_invalid.<locals>.addO   r   r   z8Untyped global name 'add':.*using cpu function on device)
r   assertRaisesr
   rI   recompiler-   searchstr	exception)r/   rB   raisesmsgexpecteds        r   test_cpu_dispatcher_invalid*TestDeviceFunc.test_cpu_dispatcher_invalidJ   s    
 
Z	 
	 {+v&&s+ ,H::c?F,<,<(=>dJK	 ,+s   B
Bc                 *  ^ [         S 5       n[        R                  " SS9mUTl        A[        R                   U4S j5       n[
        R                  " S5      nUS-   nUSUR                  4   " U5        [
        R                  R                  XC5        g )Nc                 
    X-   $ r   r   r   s     r   rB   <TestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add[   r   r   mymod)namec                 ^   > [         R                  " S5      nTR                  X   S5      X'   g r   )r   r!   rB   )r"   r#   ra   s     r   rC   CTestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add_kernelc   s$    		!AYYsvq)CFr   r'   r   )
r   types
ModuleTyperB   r   r+   r,   rE   rF   rG   )r/   rB   rC   r"   rH   ra   s        @r    test_cpu_dispatcher_other_module/TestDeviceFunc.test_cpu_dispatcher_other_moduleZ   s    		 
	   g.			* 
	* iimq1chh;$


,r   c                    [         R                  " SS9S 5       n[        [        4nUR                  U5      nUR                  R
                  nU R                  SU5        UR                  U5      nU R                  XE5        g )NTr   c                 
    X-   $ r   r   xys     r   foo-TestDeviceFunc.test_inspect_llvm.<locals>.fooo   r   r   rn   )r   r   r	   compile_devicefndescmangled_nameassertIninspect_llvm)r/   rn   argscresfnamellvms         r   test_inspect_llvm TestDeviceFunc.test_inspect_llvmm   sq    			 
	 u~!!$'((eU#%e"r   c                    [         R                  " SS9S 5       n[        [        4nUR                  U5      nUR                  R
                  nU R                  SU5        UR                  U5      nU R                  XE5        g )NTr   c                 
    X-   $ r   r   rk   s     r   rn   ,TestDeviceFunc.test_inspect_asm.<locals>.foo   r   r   rn   )r   r   r	   rp   rq   rr   rs   inspect_asm)r/   rn   ru   rv   rw   ptxs         r   test_inspect_asmTestDeviceFunc.test_inspect_asm~   so    			 
	 u~!!$'((eU#ood#e!r   c                    [         R                  " SS9S 5       nU R                  [        5       nUR	                  [
        [
        45        S S S 5        U R                  S[        WR                  5      5        g ! , (       d  f       N4= f)NTr   c                 
    X-   $ r   r   rk   s     r   rn   8TestDeviceFunc.test_inspect_sass_disallowed.<locals>.foo   r   r   z(Cannot inspect SASS of a device function)	r   r   rS   RuntimeErrorinspect_sassr	   rs   rW   rX   )r/   rn   rY   s      r   test_inspect_sass_disallowed+TestDeviceFunc.test_inspect_sass_disallowed   sn    			 
	 |,eU^, - 	@&**+	- -,s   A::
Bz'cudasim will allow calling any functionc                     [         R                  " SS9S 5       nU R                  [        5       nUS   " 5         S S S 5        U R	                  S[        WR                  5      5        g ! , (       d  f       N4= f)NTr   c                      g r   r   r   r   r   f?TestDeviceFunc.test_device_func_as_kernel_disallowed.<locals>.f   s    r   r   r   z,Cannot compile a device function as a kernel)r   r   rS   r   rs   rW   rX   )r/   r   rY   s      r   %test_device_func_as_kernel_disallowed4TestDeviceFunc.test_device_func_as_kernel_disallowed   se    			 
	 |,dGI - 	D&**+	- -,s   A))
A7z2cudasim ignores casting by jit decorator signaturec                 f  ^ [         R                  " SSS9S 5       m[         R                  U4S j5       n[         R                  " S[        R                  S9n[         R
                  " [        R                  " / SQ[        R                  S95      nUS	   " X#5        U R                  S
US   5        g )Nz!int32(int32, int32, int32, int32)Tr   c                 H    U S-  S-  US-  S-  -  US-  S-  -  US-  S-  -  $ )N         r      r   )rgr   r   s       r   rgba0TestDeviceFunc.test_device_casting.<locals>.rgba   sF    $h2%$h1_&$h1_& $h2%' (r   c                 :   > T" US   US   US   US   5      U S'   g )Nr   r         r   )rl   channelsr   s     r   rgba_caller7TestDeviceFunc.test_device_casting.<locals>.rgba_caller   s'    Xa[(1+x{KAaDr   r   r(   )g      ?g       @g      @g      @r   ir   )	r   r   device_arrayr+   r	   	to_deviceasarrayr   assertEqual)r/   r   rl   r   r   s       @r   test_device_casting"TestDeviceFunc.test_device_casting   s     
5d	C	( 
D	( 
	L 
	L arxx0>>"**-A35::#? @ 	D!&QqT*r   c                     U R                  UR                  S5        U R                  UR                  R                  [        S S  45        U R                  UR                  R
                  [        5        g Nf1)r   rb   sigru   r   return_typer	   )r/   decls     r   _test_declare_device#TestDeviceFunc._test_declare_device   sN    D)6--u5r   z!cudasim does not check signaturesc                 t    [         R                  " S[        [        S S  5      5      nU R	                  U5        g r   )r   declare_devicer	   r   r   r/   r   s     r   test_declare_device_signature,TestDeviceFunc.test_declare_device_signature   s+      uWQZ'89!!"%r   c                 T    [         R                  " SS5      nU R                  U5        g )Nr   zint32(float32[:]))r   r   r   r   s     r   test_declare_device_string)TestDeviceFunc.test_declare_device_string   s#      ':;!!"%r   c                     U R                  [        S5         [        R                  " S[        S S  45        S S S 5        g ! , (       d  f       g = f)NReturn typer   )assertRaisesRegex	TypeErrorr   r   r   r/   s    r   test_bad_declare_device_tuple,TestDeviceFunc.test_bad_declare_device_tuple   s4    ##I}=wqzm4 >==s    A  
Ac                     U R                  [        S5         [        R                  " SS5        S S S 5        g ! , (       d  f       g = f)Nr   r   z(float32[:],))r   r   r   r   r   s    r   test_bad_declare_device_string-TestDeviceFunc.test_bad_declare_device_string   s.    ##I}=o6 >==s	   8
Ar   N)__name__
__module____qualname____firstlineno__r3   r>   rI   rN   r   r\   rg   ry   r   r   r   r   r   r   r   r   r   __static_attributes__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__)rT   re   numpyr+   numba.cuda.testingr   r   r   numbar   r   r   r	   numba.core.errorsr
   r   r   mainr   r   r   <module>r      sD    	   F F + + )O7\ O7d zMMO r   