
    `i3                        d dl Zd dlmZmZmZmZ d dlmZ d dl	m
Z
mZ d dlmZ d dlZd dlZd dlmZ d Zd Zd	 Zd
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z d Z! ej"        dd          d             Z# ej"        dd          d             Z$d Z%d Z&d Z'd Z(d  Z) G d! d"e          Z*e+d#k    r ej,                     dS dS )$    N)unittestCUDATestCaseskip_unless_cc_53skip_on_cudasim)cuda)f2b1)compile_ptx)
from_dtypec                     ||z  | d<   d S Nr    aryabs      y/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_operator.pysimple_fp16_div_scalarr          UCFFF    c                     ||z   | d<   d S r   r   r   s      r   simple_fp16addr      r   r   c                 &    | dxx         |z  cc<   d S r   r   r   r   s     r   simple_fp16_iaddr          FFFaKFFFFFr   c                 &    | dxx         |z  cc<   d S r   r   r   s     r   simple_fp16_isubr      r   r   c                 &    | dxx         |z  cc<   d S r   r   r   s     r   simple_fp16_imulr       r   r   c                 &    | dxx         |z  cc<   d S r   r   r   s     r   simple_fp16_idivr"       r   r   c                     ||z
  | d<   d S r   r   r   s      r   simple_fp16subr$   $   r   r   c                     ||z  | d<   d S r   r   r   s      r   simple_fp16mulr&   (   r   r   c                     | | d<   d S r   r   r   s     r   simple_fp16negr(   ,   s    RCFFFr   c                 *    t          |          | d<   d S r   )absr   s     r   simple_fp16absr+   0   s    VVCFFFr   c                     ||k    | d<   d S r   r   r   s      r   simple_fp16_gtr-   4       UCFFFr   c                     ||k    | d<   d S r   r   r   s      r   simple_fp16_ger0   8       !VCFFFr   c                     ||k     | d<   d S r   r   r   s      r   simple_fp16_ltr3   <   r.   r   c                     ||k    | d<   d S r   r   r   s      r   simple_fp16_ler5   @   r1   r   c                     ||k    | d<   d S r   r   r   s      r   simple_fp16_eqr7   D   r1   r   c                     ||k    | d<   d S r   r   r   s      r   simple_fp16_ner9   H   r1   r   z
b1(f2, f2)T)devicec                     | |k     S Nr   xys     r   
hlt_func_1r@   L       q5Lr   c                     | |k     S r<   r   r=   s     r   
hlt_func_2rC   Q   rA   r   c                 L    t          ||          ot          ||          | d<   d S r   )r@   rC   rr   r   cs       r   test_multiple_hcmp_1rH   V   s(    a0
1a 0 0AaDDDr   c                 8    t          ||          o||k     | d<   d S r   r@   rE   s       r   test_multiple_hcmp_2rK   [   s"    a%AAaDDDr   c                 8    t          ||          o||k    | d<   d S r   rJ   rE   s       r   test_multiple_hcmp_3rM   `   s"    a&QAaDDDr   c                 $    ||k     o||k     | d<   d S r   r   rE   s       r   test_multiple_hcmp_4rO   e   s    q5?QUAaDDDr   c                 $    ||k     o||k    | d<   d S r   r   rE   s       r   test_multiple_hcmp_5rQ   j   s    q5Q!VAaDDDr   c                       e Zd Z fdZ	 d Zd Zd Zd Zd Zd Z	e
d             Z ed	          d
             Ze
d             Z ed	          d             Ze
d             Ze
d             Z ed	          d             Z ed	          d             Ze
d             Ze
d             Ze
d             Ze
d             Z ed	          d             Z ed	          d             Z ed	          d             Z xZS )TestOperatorModulec                     t                                                       t          j                            d           d S r   )supersetUpnprandomseed)self	__class__s    r   rV   zTestOperatorModule.setUpp   s.    
	qr   c                 &   t           j        fd            }t          j        d          }t          j        d          }|                                } |d         ||           t          j                            | ||                     d S )Nc                 B    d} | |         ||                   | |<   d S r   r   )r   r   iops      r   fooz1TestOperatorModule.operator_template.<locals>.foox   s'    A2adAaD>>AaDDDr      ra   ra   )r   jitrW   onescopytestingassert_equal)rZ   r_   r`   r   r   ress    `    r   operator_templatez$TestOperatorModule.operator_templatew   s    		" 	" 	" 	" 
	" GAJJGAJJffhhD	#q

RR1XX.....r   c                 D    |                      t          j                   d S r<   )ri   operatoraddrZ   s    r   test_addzTestOperatorModule.test_add       x|,,,,,r   c                 D    |                      t          j                   d S r<   )ri   rk   subrm   s    r   test_subzTestOperatorModule.test_sub   ro   r   c                 D    |                      t          j                   d S r<   )ri   rk   mulrm   s    r   test_mulzTestOperatorModule.test_mul   ro   r   c                 D    |                      t          j                   d S r<   )ri   rk   truedivrm   s    r   test_truedivzTestOperatorModule.test_truediv   s    x/00000r   c                 D    |                      t          j                   d S r<   )ri   rk   floordivrm   s    r   test_floordivz TestOperatorModule.test_floordiv   s    x011111r   c                 "   t           t          t          t          f}t          j        t          j        t          j        t          j        f}t          ||          D ]4\  }}| 
                    |          5   t          j        d          |          }t          j        dt          j                  }t          j                            d                              t          j                  }t          j                            d                              t          j                  } |d         ||d         |d                     |||          }	t          j                            ||	           d d d            n# 1 swxY w Y   6d S )Nr_   zvoid(f2[:], f2, f2)ra   dtyperb   r   )r   r$   r&   r   rk   rl   rq   rt   rw   zipsubTestr   rc   rW   zerosfloat16rX   astyperf   assert_allclose
rZ   	functionsopsfnr_   kernelgotarg1arg2expecteds
             r   test_fp16_binaryz#TestOperatorModule.test_fp16_binary   s|   #^^+-	|X\8<9IJ)S)) 
	: 
	:FB$$ 	: 	:8"788<<hq
333y''**11"*==y''**11"*==tS$q'473332dD>>
**3999	: 	: 	: 	: 	: 	: 	: 	: 	: 	: 	: 	: 	: 	: 	:
	: 
	:s   5DFF		F	z(Compilation unsupported in the simulatorc                 H   t           t          t          f}d}t          d d          t          t          f}t	          ||          D ]^\  }}|                     |          5  t          ||d          \  }}|                     ||           d d d            n# 1 swxY w Y   _d S N)zadd.f16zsub.f16zmul.f16)instr      cc)r   r$   r&   r   r   r   r
   assertInrZ   r   instrsargsr   r   ptx_s           r   test_fp16_binary_ptxz'TestOperatorModule.test_fp16_binary_ptx   s    #^^D	2111r2Y// 	* 	*IBE** * *$R&999QeS)))* * * * * * * * * * * * * * *	* 	*s   ,BB	B	c                    t           t          t          t          f}t          j        t          j        t          j        t          j        f}t          j
        t          j        t          j        t          j        t          j        t          j        f}t!          j        t%          ||          |          D ]=\  \  }}}|                     ||          5  t)          j        |          }t          j                            d                              t          j                  }t          j                            d          dz                      |          }	t          j        t          j        |          }
t          j        d|
          } |d         ||d         |	d                     |||	          }t          j                            ||           d d d            n# 1 swxY w Y   ?d S )Nr_   tyra   d   r~   rb   r   )r   r$   r&   r   rk   rl   rq   rt   rw   rW   int8int16int32int64float32float64	itertoolsproductr   r   r   rc   rX   r   r   result_typer   rf   r   )rZ   r   r   typesr   r_   r   r   r   r   res_tyr   r   s                r   !test_mixed_fp16_binary_arithmeticz4TestOperatorModule.test_mixed_fp16_binary_arithmetic   s   #^^+-	|X\8<9IJ"(BHbhRZ)%-c)S.A.A5II 	: 	:LHRb++ 
: 
:"y''**11"*==	((++c199"==
B77hq///tS$q'473332dD>>
**3999
: 
: 
: 
: 
: 
: 
: 
: 
: 
: 
: 
: 
: 
: 
:	: 	:s   DG""G&	)G&	c                 <   t           t          t          f}d}t          d d          t          f}t	          ||          D ]^\  }}|                     |          5  t          ||d          \  }}|                     ||           d d d            n# 1 swxY w Y   _d S r   )r   r   r    r   r   r   r
   r   r   s           r   test_fp16_inplace_binary_ptxz/TestOperatorModule.test_fp16_inplace_binary_ptx   s    %'79IJ	2111r{Y// 	* 	*IBE** * *$R&999QeS)))* * * * * * * * * * * * * * *	* 	*s   ,BB	B	c                    t           t          t          t          f}t          j        t          j        t          j        t          j        f}t          ||          D ]!\  }}| 
                    |          5   t          j        d          |          }t          j                            d                              t          j                  }|                                }t          j                            d                              t          j                  d         } |d         ||            |||           t          j                            ||           d d d            n# 1 swxY w Y   #d S )Nr}   void(f2[:], f2)ra   r   rb   )r   r   r    r"   rk   iaddisubimulitruedivr   r   r   rc   rW   rX   r   r   re   rf   r   )	rZ   r   r   r   r_   r   r   r   args	            r   test_fp16_inplace_binaryz+TestOperatorModule.test_fp16_inplace_binary   ss   %'79I%'	}hmX]H<MN)S)) 		: 		:FB$$ : :4"344R88i&&q))00<<88::i&&q))00<<Q?tS#&&&8S!!!
**3999: : : : : : : : : : : : : : :		: 		:s   5C.E//E3	6E3	c                 R   t           t          f}t          j        t          j        f}t          ||          D ]\  }}|                     |          5   t          j        d          |          }t          j
        dt          j                  }t          j                            d                              t          j                  } |d         ||d                     ||          }t          j                            ||           d d d            n# 1 swxY w Y   d S )Nr}   r   ra   r~   rb   r   )r(   r+   rk   negr*   r   r   r   rc   rW   r   r   rX   r   rf   r   )	rZ   r   r   r   r_   r   r   r   r   s	            r   test_fp16_unaryz"TestOperatorModule.test_fp16_unary   s?   #^4	|X\*)S)) 		: 		:FB$$ : :4"344R88hq
333y''**11"*==tS$q'***2d88
**3999: : : : : : : : : : : : : : :		: 		:s   B=DD	"D	c                     t           d d          t           f}t          t          |d          \  }}|                     d|           d S )Nr   r   zneg.f16)r   r
   r(   r   rZ   r   r   r   s       r   test_fp16_neg_ptxz$TestOperatorModule.test_fp16_neg_ptx   sE    111r{^Tf===Qi%%%%%r   c                     t           d d          t           f}t          t          |d          \  }}|                     d|           d S )Nr   r   zabs.f16)r   r
   r+   r   r   s       r   test_fp16_abs_ptxz$TestOperatorModule.test_fp16_abs_ptx   sE    111r{^Tf===Qi%%%%%r   c                 ^   t           t          t          t          t          t
          f}t          j        t          j        t          j	        t          j
        t          j        t          j        f}t          ||          D ]0\  }}|                     |          5   t          j        d          |          }t#          j        dt"          j                  }t"          j                            d                              t"          j                  }t"          j                            d                              t"          j                  } |d         ||d         |d                     |||          }	|                     |d         |	           d d d            n# 1 swxY w Y   2d S )Nr}   zvoid(b1[:], f2, f2)ra   r~   rb   r   )r-   r0   r3   r5   r7   r9   rk   gtgeltleeqner   r   r   rc   rW   r   bool_rX   r   r   assertEqualr   s
             r   test_fp16_comparisonz'TestOperatorModule.test_fp16_comparison   s   #^#^#^5	 {HKhk{HK) )S)) 
	3 
	3FB$$ 	3 	38"788<<hq111y''**11"*==y''**11"*==tS$q'473332dD>>  Q222	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3
	3 
	3s   C=F  F$	'F$	c                    t           t          t          t          t          t
          f}t          j        t          j        t          j	        t          j
        t          j        t          j        f}t          j        t          j        t          j        t          j        t          j        t          j        f}t)          j        t-          ||          |          D ]$\  \  }}}|                     ||          5  t1          j        |          }t          j        dt          j                  }t          j                            d                              t          j                  }	t          j                            d          dz                      |          }
 |d         ||	d         |
d                     ||	|
          }|                     |d         |           d d d            n# 1 swxY w Y   &d S )Nr   ra   r~   r   rb   r   ) r-   r0   r3   r5   r7   r9   rk   r   r   r   r   r   r   rW   r   r   r   r   r   r   r   r   r   r   r   rc   r   r   rX   r   r   r   )rZ   r   r   r   r   r_   r   r   r   r   r   r   s               r   test_mixed_fp16_comparisonz-TestOperatorModule.test_mixed_fp16_comparison  s   #^#^#^5	 {HKhk{HK)"(BHbhRZ) &-c)S.A.A.35 5 	3 	3LHRb++ 	3 	3"hq111y''**11"*==	((++c199"==tS$q'473332dD>>  Q222	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3 	3	3 	3s   2C-G++G/	2G/	c                    t           t          t          t          t          f}|D ]}|                     |          5   t          j        d          |          }t          j	        dt          j
                  }t          j        d          }t          j        d          }t          j        d          } |d         ||||           |                     |d	                    d d d            n# 1 swxY w Y   d S )
Nr   void(b1[:], f2, f2, f2)ra   r~          @      @g      @rb   r   )rH   rK   rM   rO   rQ   r   r   rc   rW   r   r   r   
assertTruerZ   r   r   compiledr   r   r   arg3s           r   !test_multiple_float16_comparisonsz4TestOperatorModule.test_multiple_float16_comparisons'  s$   )))))	+	
  	( 	(B$$ ( (>48$=>>rBBhq111z"~~z"~~z"~~sD$555A'''( ( ( ( ( ( ( ( ( ( ( ( ( ( (	( 	(   B)C00C4	7C4	c                    t           t          t          t          t          f}|D ]}|                     |          5   t          j        d          |          }t          j	        dt          j
                  }t          j        d          }t          j        d          }t          j        d          } |d         ||||           |                     |d	                    d d d            n# 1 swxY w Y   d S )
Nr   r   ra   r~   r   r   g      ?rb   r   )rH   rK   rM   rO   rQ   r   r   rc   rW   r   r   r   assertFalser   s           r   'test_multiple_float16_comparisons_falsez:TestOperatorModule.test_multiple_float16_comparisons_false8  s&   )))))	+	
  	) 	)B$$ ) )>48$=>>rBBhq111z"~~z"~~z"~~sD$555  Q((() ) ) ) ) ) ) ) ) ) ) ) ) ) )	) 	)r   c                    t           t          t          t          t          t
          f}t          j        t          j        t          j	        t          j
        t          j        t          j        f}d}t          d d          t          t          f}t          |||          D ]_\  }}}|                     |          5  t#          ||d          \  }}	|                     ||           d d d            n# 1 swxY w Y   `d S )N)setp.gt.f16setp.ge.f16setp.lt.f16setp.le.f16setp.eq.f16setp.ne.f16r}   r   r   )r-   r0   r3   r5   r7   r9   rk   r   r   r   r   r   r   r	   r   r   r   r
   r   )
rZ   r   r   opstringr   r   r_   sr   r   s
             r   test_fp16_comparison_ptxz+TestOperatorModule.test_fp16_comparison_ptxI  s   #^#^#^5	 {HKhk{HK)2 111r2YX66 	& 	&IBA$$ & &$R&999Qa%%%& & & & & & & & & & & & & & &	& 	&s   6,C..C2	5C2	c                    t           t          t          t          t          t
          f}t          j        t          j        t          j	        t          j
        t          j        t          j        f}t          j        dt          j        dt          j	        dt          j
        dt          j        dt          j        di}t          ||          D ]\  }}|                     |          5  t          d d          t           t#          t$          j                  f}t)          ||d	          \  }}|                     ||         |           d d d            n# 1 swxY w Y   d S )
Nr   r   r   r   r   r   r}   r   r   )r-   r0   r3   r5   r7   r9   rk   r   r   r   r   r   r   r   r   r	   r   r   rW   r   r
   r   )	rZ   r   r   r   r   r_   r   r   r   s	            r   test_fp16_int8_comparison_ptxz0TestOperatorModule.test_fp16_int8_comparison_ptxZ  sQ    $^#^#^5	 {HKhk{HK) KKKKKK/ )S)) 	1 	1FB$$ 1 1111r:bg#6#67$R&999QhrlC0001 1 1 1 1 1 1 1 1 1 1 1 1 1 1	1 	1s    A EE	E	c                 8   t           t          t          t          t          t
          f}t          j        t          j        t          j	        t          j
        t          j        t          j        f}t          j        t          j        t          j        t          j        t          j        f}t          j        dt          j        dt          j	        dt          j
        dt          j        dt          j        di}t          j        d          dt          j        d	          dt          j        d
          dt          j        d          di}t)          j        t-          ||          |          D ]\  \  }}}|                     ||          5  t          j        t          j        |          }	t4          d d          t6          t9          |	          f}
t;          ||
d          \  }}||         ||	         z   }|                     ||           d d d            n# 1 swxY w Y   d S )Nzsetp.gt.zsetp.ge.zsetp.lt.zsetp.le.zsetp.eq.z	setp.neu.r   f64r   r   f32r   r   r   r   )r-   r0   r3   r5   r7   r9   rk   r   r   r   r   r   r   rW   r   r   r   r   r   r   r   r   r   r   r   r   r	   r   r   r
   r   )rZ   r   r   types_promoter   opsuffixr   r_   r   arg2_tyr   r   r   s                r   (test_mixed_fp16_comparison_promotion_ptxz;TestOperatorModule.test_mixed_fp16_comparison_promotion_ptxp  s   #^#^#^5	 {HKhk{HK) 28RXRZ1K
K
K
K
K
K- HW%%uHW%%uHY''HY''0
 &-c)S.A.A.;= = 	( 	(LHRb++ ( (.R88111r:g#6#67$R&999QrlXg%66c3'''( ( ( ( ( ( ( ( ( ( ( ( ( ( (	( 	(s   B HH	H	)__name__
__module____qualname__rV   ri   rn   rr   ru   rx   r{   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   __classcell__)r[   s   @r   rS   rS   o   s           / / /- - -- - -- - -1 1 12 2 2 : : :" _?@@* * A@* : : :& _?@@* * A@* : : :  : : : _?@@& & A@&
 _?@@& & A@& 3 3 3& 3 3 3, ( ( (  ) ) )  _?@@& & A@&  _?@@1 1 A@1* _?@@( ( A@( ( ( ( (r   rS   __main__)-numpyrW   numba.cuda.testingr   r   r   r   numbar   numba.core.typesr   r	   
numba.cudar
   rk   r   numba.np.numpy_supportr   r   r   r   r   r    r"   r$   r&   r(   r+   r-   r0   r3   r5   r7   r9   rc   r@   rC   rH   rK   rM   rO   rQ   rS   r   mainr   r   r   <module>r      s      1 1 1 1 1 1 1 1 1 1 1 1       # # # # # # # # " " " " " "      - - - - - -                                 
,t$$$  %$ 
,t$$$  %$1 1 1
& & &
' ' '
  
  
^( ^( ^( ^( ^( ^( ^( ^(B	 zHMOOOOO r   