
    `i                         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                     dS dS )    N)unittestskip_on_cudasimCUDATestCase)cudajitfloat32int32)TypingErrorc                      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dS )TestDeviceFuncc                 Z   t          j        dd          d             fd} t          j        d          |          }d}t          j        |t          j                  }||z   } |d	|f         |           |                     t          j        ||k              ||f           d S )
Nfloat32(float32, float32)Tdevicec                     | |z   S N abs     |/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_device_func.pyadd2fz,TestDeviceFunc.test_use_add2f.<locals>.add2f       q5L    c                 f    t          j        d          } | |         | |                   | |<   d S N   r   grid)aryir   s     r   	use_add2fz0TestDeviceFunc.test_use_add2f.<locals>.use_add2f   s1    	!AU3q63q6**CFFFr   void(float32[:])
   dtyper   r   r   nparanger   
assertTrueall)selfr"   compilednelemr    expr   s         @r   test_use_add2fzTestDeviceFunc.test_use_add2f   s    	-d	;	;	;	 	 
<	;		+ 	+ 	+ 	+ 	+ 048.//	::iRZ000CiE3scz**S#J7777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         |           |                     t          j        ||k              ||f           d S )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    5A;;r   c                 f    t          j        d          } | |         | |                   | |<   d S r   r   )r    r!   r4   s     r   indirect_add2fz:TestDeviceFunc.test_indirect_add2f.<locals>.indirect_add2f*   s1    	!AXc!fc!f--CFFFr   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	;	;	;	 	 	 	 
<	;		. 	. 	. 	. 	. 048.//??iRZ000CiE3scz**S#J77777r   c                     t           j        fd            }t          j        d          }|dz   } |d|j        f         |           t          j                            ||           d S )Nc                 Z    t          j        d          } | |         d          | |<   d S r   r   )r    r!   adds     r   
add_kernelz8TestDeviceFunc._check_cpu_dispatcher.<locals>.add_kernel8   s+    	!ASQ^^CFFFr   r$   r   )r   r   r(   r)   sizetestingassert_equal)r,   r:   r;   r    expects    `   r   _check_cpu_dispatcherz$TestDeviceFunc._check_cpu_dispatcher7   sy    		$ 	$ 	$ 	$ 
	$ immq
1ch;$$$

,,,,,r   c                 P    t           d             }|                     |           d S )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   s7    		 	 
	 	""3'''''r   znot supported in cudasimc                 h   t          d          d             }|                     t                    5 }|                     |           d d d            n# 1 swxY w Y   d}t	          j        |          }|                     |                    t          |j	                            d u           d S )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KKKKs   AAAc                 2   t           d             }t          j        d          |_        ~t          j         fd            }t          j        d          }|dz   } |d|j        f         |           t
          j        	                    ||           d S )Nc                     | |z   S r   r   r   s     r   r:   z<TestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add[   r   r   mymod)namec                 n    t          j        d          }                    | |         d          | |<   d S r   )r   r   r:   )r    r!   rR   s     r   r;   zCTestDeviceFunc.test_cpu_dispatcher_other_module.<locals>.add_kernelc   s/    	!AYYs1vq))CFFFr   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...			* 	* 	* 	* 
	* immq
1ch;$$$

,,,,,r   c                 "   t          j        d          d             }t          t          f}|                    |          }|j        j        }|                     d|           |                    |          }|                     ||           d S )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   s    					 	 
		 u~!!$''(eU###%%eT"""""r   c                 "   t          j        d          d             }t          t          f}|                    |          }|j        j        }|                     d|           |                    |          }|                     ||           d S )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~   s    					 	 
		 u~!!$''(eU###ood##eS!!!!!r   c                 8   t          j        d          d             }|                     t                    5 }|                    t
          t
          f           d d d            n# 1 swxY w Y   |                     dt          |j                             d S )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   s    					 	 
		 |,, 	-eU^,,,	- 	- 	- 	- 	- 	- 	- 	- 	- 	- 	- 	- 	- 	- 	- 	@&*++	- 	- 	- 	- 	-s   "A''A+.A+z'cudasim will allow calling any functionc                    t          j        d          d             }|                     t                    5 } |d                      d d d            n# 1 swxY w Y   |                     dt          |j                             d S )NTr   c                      d S r   r   r   r   r   fz?TestDeviceFunc.test_device_func_as_kernel_disallowed.<locals>.f   s    Dr   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   s    					 	 
		 |,, 	AdGIII	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	D&*++	- 	- 	- 	- 	-s   AAAz2cudasim ignores casting by jit decorator signaturec                 x   t          j        dd          d             t           j        fd            }t          j        dt          j                  }t          j        t          j        g dt          j                            } |d	         ||           |                     d
|d                    d S )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                 Z     |d         |d         |d         |d                   | d<   d S )Nr   r         r   )r[   channelsr   s     r   rgba_callerz7TestDeviceFunc.test_device_casting.<locals>.rgba_caller   s0    4Xa[(1+x{KKAaDDDr   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	C	C	( 	( 
D	C	( 
	L 	L 	L 	L 
	L arx000>"*-A-A-A35:#? #? #? @ @ 	D!X&&&QqT*****r   c                     |                      |j        d           |                      |j        j        t          d d          f           |                      |j        j        t                     d S Nf1)r   rS   sigrc   r   return_typer	   )r,   decls     r   _test_declare_devicez#TestDeviceFunc._test_declare_device   s_    D)))666-u55555r   z!cudasim does not check signaturesc                     t          j        dt          t          d d                              }|                     |           d S r   )r   declare_devicer	   r   r   r,   r   s     r   test_declare_device_signaturez,TestDeviceFunc.test_declare_device_signature   s>     uWQQQZ'8'899!!"%%%%%r   c                 Z    t          j        dd          }|                     |           d S )Nr   zint32(float32[:]))r   r   r   r   s     r   test_declare_device_stringz)TestDeviceFunc.test_declare_device_string   s/     ':;;!!"%%%%%r   c                     |                      t          d          5  t          j        dt          d d          f           d d d            d S # 1 swxY w Y   d S )NReturn typer   )assertRaisesRegex	TypeErrorr   r   r   r,   s    r   test_bad_declare_device_tuplez,TestDeviceFunc.test_bad_declare_device_tuple   s    ##I}== 	5 	5wqqqzm444	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5 	5s   $AAAc                     |                      t          d          5  t          j        dd           d d d            d S # 1 swxY w Y   d S )Nr   r   z(float32[:],))r   r   r   r   r   s    r   test_bad_declare_device_stringz-TestDeviceFunc.test_bad_declare_device_string   s    ##I}== 	7 	7o666	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7 	7s   ?A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 8&8 8 8.	- 	- 	-( ( ( _/00L L 10L- - -& _/00# # 10#  _/00" " 10"  _/00	- 	- 10	- _>??	- 	- @?	- _IJJ+ + KJ+66 6 6
 _899& & :9& _899& & :9& _8995 5 :95 _8997 7 :97 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      s    				      F F F F F F F F F F + + + + + + + + + + + + ) ) ) ) ) )O7 O7 O7 O7 O7\ O7 O7 O7d zHMOOOOO r   