
    wg3                        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jD                  dd      d        Z# ejD                  dd      d        Z$d Z%d Z&d Z'd Z(d  Z) G d! d"e      Z*e+d#k(  r ejX                          yy)$    N)unittestCUDATestCaseskip_unless_cc_53skip_on_cudasim)cuda)f2b1)compile_ptx)
from_dtypec                     ||z  | d<   y Nr    aryabs      j/home/mcse/projects/flask/flask-venv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_operator.pysimple_fp16_div_scalarr          UCF    c                     ||z   | d<   y r   r   r   s      r   simple_fp16addr      r   r   c                     | dxx   |z  cc<   y r   r   r   r   s     r   simple_fp16_iaddr          FaKFr   c                     | dxx   |z  cc<   y r   r   r   s     r   simple_fp16_isubr      r   r   c                     | dxx   |z  cc<   y r   r   r   s     r   simple_fp16_imulr       r   r   c                     | dxx   |z  cc<   y r   r   r   s     r   simple_fp16_idivr"       r   r   c                     ||z
  | d<   y r   r   r   s      r   simple_fp16subr$   $   r   r   c                     ||z  | d<   y r   r   r   s      r   simple_fp16mulr&   (   r   r   c                     | | d<   y r   r   r   s     r   simple_fp16negr(   ,   s    RCFr   c                      t        |      | d<   y r   )absr   s     r   simple_fp16absr+   0   s    VCFr   c                     ||kD  | d<   y r   r   r   s      r   simple_fp16_gtr-   4   r   r   c                     ||k\  | d<   y r   r   r   s      r   simple_fp16_ger/   8       !VCFr   c                     ||k  | d<   y r   r   r   s      r   simple_fp16_ltr2   <   r   r   c                     ||k  | d<   y r   r   r   s      r   simple_fp16_ler4   @   r0   r   c                     ||k(  | d<   y r   r   r   s      r   simple_fp16_eqr6   D   r0   r   c                     ||k7  | d<   y r   r   r   s      r   simple_fp16_ner8   H   r0   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_2rB   Q   r@   r   c                 >    t        ||      xr t        ||      | d<   y r   )r?   rB   rr   r   cs       r   test_multiple_hcmp_1rG   V   s    a0
1a 0AaDr   c                 0    t        ||      xr ||k  | d<   y r   r?   rD   s       r   test_multiple_hcmp_2rJ   [   s    a%AAaDr   c                 0    t        ||      xr ||k\  | d<   y r   rI   rD   s       r   test_multiple_hcmp_3rL   `   s    a&QAaDr   c                 "    ||k  xr ||k  | d<   y r   r   rD   s       r   test_multiple_hcmp_4rN   e   s    q5?QUAaDr   c                 "    ||k  xr ||k\  | d<   y r   r   rD   s       r   test_multiple_hcmp_5rP   j   s    q5Q!VAaDr   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                  j	                  d       y r   )supersetUpnprandomseed)self	__class__s    r   rU   zTestOperatorModule.setUpp   s    
		qr   c                    t         j                  fd       }t        j                  d      }t        j                  d      }|j	                         } |d   ||       t        j
                  j                  | ||             y )Nc                 .    d} | |   ||         | |<   y r   r   )r   r   iops      r   fooz1TestOperatorModule.operator_template.<locals>.foox   s    AadAaD>AaDr      r`   r`   )r   jitrV   onescopytestingassert_equal)rY   r^   r_   r   r   ress    `    r   operator_templatez$TestOperatorModule.operator_templatew   sj    		" 
	" GGAJGGAJffhD	#q


R1X.r   c                 B    | j                  t        j                         y r;   )rh   operatoraddrY   s    r   test_addzTestOperatorModule.test_add       x||,r   c                 B    | j                  t        j                         y r;   )rh   rj   subrl   s    r   test_subzTestOperatorModule.test_sub   rn   r   c                 B    | j                  t        j                         y r;   )rh   rj   mulrl   s    r   test_mulzTestOperatorModule.test_mul   rn   r   c                 B    | j                  t        j                         y r;   )rh   rj   truedivrl   s    r   test_truedivzTestOperatorModule.test_truediv   s    x//0r   c                 B    | j                  t        j                         y r;   )rh   rj   floordivrl   s    r   test_floordivz TestOperatorModule.test_floordiv   s    x001r   c                    t         t        t        t        f}t        j
                  t        j                  t        j                  t        j                  f}t        ||      D ]  \  }}| j                  |      5   t        j                  d      |      }t        j                  dt        j                        }t        j                   j!                  d      j#                  t        j                        }t        j                   j!                  d      j#                  t        j                        } |d   ||d   |d           |||      }	t        j$                  j'                  ||	       d d d         y # 1 sw Y   $xY w)Nr^   zvoid(f2[:], f2, f2)r`   dtypera   r   )r   r$   r&   r   rj   rk   rp   rs   rv   zipsubTestr   rb   rV   zerosfloat16rW   astypere   assert_allclose
rY   	functionsopsfnr^   kernelgotarg1arg2expecteds
             r   test_fp16_binaryz#TestOperatorModule.test_fp16_binary   s   #^^+-	||X\\8<<9I9IJ)S) 
	:FB$ 	:8"78<hhq

3yy''*11"**=yy''*11"**=tS$q'473dD>

**39	: 	:
	:	: 	:s   :C5E;;F	z(Compilation unsupported in the simulatorc                    t         t        t        f}d}t        d d  t        t        f}t	        ||      D ]C  \  }}| j                  |      5  t        ||d      \  }}| j                  ||       d d d        E y # 1 sw Y   PxY wN)zadd.f16zsub.f16zmul.f16)instr      cc)r   r$   r&   r   r   r   r
   assertInrY   r   instrsargsr   r   ptx_s           r   test_fp16_binary_ptxz'TestOperatorModule.test_fp16_binary_ptx   s    #^^D	21r2Y/ 	*IBE* *$R&9QeS)* *	** *s   $A;;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 ]  \  \  }}}| j'                  ||      5  t)        j*                  |      }t        j,                  j-                  d      j/                  t        j0                        }t        j,                  j-                  d      dz  j/                  |      }	t        j2                  t        j0                  |      }
t        j4                  d|
      } |d   ||d   |	d           |||	      }t        j6                  j9                  ||       d d d        ! y # 1 sw Y   -xY w)Nr^   tyr`   d   r}   ra   r   )r   r$   r&   r   rj   rk   rp   rs   rv   rV   int8int16int32int64float32float64	itertoolsproductr   r   r   rb   rW   r   r   result_typer   re   r   )rY   r   r   typesr   r^   r   r   r   r   res_tyr   r   s                r   !test_mixed_fp16_binary_arithmeticz4TestOperatorModule.test_mixed_fp16_binary_arithmetic   sf   #^^+-	||X\\8<<9I9IJ"((BHHbhhRZZ)%--c)S.A5I 	:LHRb+ 
:"yy''*11"**=		((+c199"=

B7hhq/tS$q'473dD>

**39
: 
:	:
: 
:s   .C:G44G>	c                    t         t        t        f}d}t        d d  t        f}t	        ||      D ]C  \  }}| j                  |      5  t        ||d      \  }}| j                  ||       d d d        E y # 1 sw Y   PxY w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1r{Y/ 	*IBE* *$R&9QeS)* *	** *s   $A66A?	c                    t         t        t        t        f}t        j
                  t        j                  t        j                  t        j                  f}t        ||      D ]  \  }}| j                  |      5   t        j                  d      |      }t        j                  j                  d      j                  t        j                         }|j#                         }t        j                  j                  d      j                  t        j                         d   } |d   ||        |||       t        j$                  j'                  ||       d d d         y # 1 sw Y   	xY w)Nr|   void(f2[:], f2)r`   r   ra   )r   r   r    r"   rj   iaddisubimulitruedivr   r   r   rb   rV   rW   r   r   rd   re   r   )	rY   r   r   r   r^   r   r   r   args	            r   test_fp16_inplace_binaryz+TestOperatorModule.test_fp16_inplace_binary   s   %'79I%'	}}hmmX]]H<M<MN)S) 		:FB$ :4"34R8ii&&q)00<88:ii&&q)00<Q?tS#&8S!

**39: :		:: :s   9CE  E*	c                 8   t         t        f}t        j                  t        j                  f}t        ||      D ]  \  }}| j                  |      5   t        j                  d      |      }t        j                  dt        j                        }t        j                  j                  d      j                  t        j                        } |d   ||d           ||      }t        j                  j                  ||       d d d         y # 1 sw Y   xY w)Nr|   r   r`   r}   ra   r   )r(   r+   rj   negr*   r   r   r   rb   rV   r   r   rW   r   re   r   )	rY   r   r   r   r^   r   r   r   r   s	            r   test_fp16_unaryz"TestOperatorModule.test_fp16_unary   s    #^4	||X\\*)S) 		:FB$ :4"34R8hhq

3yy''*11"**=tS$q'*d8

**39: :		:: :s   B4DD	c                 p    t         d d  t         f}t        t        |d      \  }}| j                  d|       y )Nr   r   zneg.f16)r   r
   r(   r   rY   r   r   r   s       r   test_fp16_neg_ptxz$TestOperatorModule.test_fp16_neg_ptx   s/    1r{^Tf=Qi%r   c                 p    t         d d  t         f}t        t        |d      \  }}| j                  d|       y )Nr   r   zabs.f16)r   r
   r+   r   r   s       r   test_fp16_abs_ptxz$TestOperatorModule.test_fp16_abs_ptx   s/    1r{^Tf=Qi%r   c                 J   t         t        t        t        t        t
        f}t        j                  t        j                  t        j                  t        j                  t        j                  t        j                  f}t        ||      D ]
  \  }}| j                  |      5   t        j                   d      |      }t#        j$                  dt"        j&                        }t"        j(                  j)                  d      j+                  t"        j,                        }t"        j(                  j)                  d      j+                  t"        j,                        } |d   ||d   |d           |||      }	| j/                  |d   |	       d d d         y # 1 sw Y   xY w)Nr|   zvoid(b1[:], f2, f2)r`   r}   ra   r   )r-   r/   r2   r4   r6   r8   rj   gtgeltleeqner   r   r   rb   rV   r   bool_rW   r   r   assertEqualr   s
             r   test_fp16_comparisonz'TestOperatorModule.test_fp16_comparison   s4   #^#^#^5	 {{HKKhkk{{HKK) )S) 
	3FB$ 	38"78<hhq1yy''*11"**=yy''*11"**=tS$q'473dD>  Q2	3 	3
	3	3 	3s   "C*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 ]  \  \  }}}| j/                  ||      5  t1        j2                  |      }t        j4                  dt        j6                        }t        j8                  j9                  d      j;                  t        j<                        }	t        j8                  j9                  d      dz  j;                  |      }
 |d   ||	d   |
d           ||	|
      }| j?                  |d   |       d d d         y # 1 sw Y   
xY w)Nr   r`   r}   r   ra   r   ) r-   r/   r2   r4   r6   r8   rj   r   r   r   r   r   r   rV   r   r   r   r   r   r   r   r   r   r   r   rb   r   r   rW   r   r   r   )rY   r   r   r   r   r^   r   r   r   r   r   r   s               r   test_mixed_fp16_comparisonz-TestOperatorModule.test_mixed_fp16_comparison  so   #^#^#^5	 {{HKKhkk{{HKK)"((BHHbhhRZZ) &--c)S.A.35 	3LHRb+ 	3"hhq1yy''*11"**=		((+c199"=tS$q'473dD>  Q2	3 	3	3	3 	3s   CG99H	c                    t         t        t        t        t        f}|D ]  }| j                  |      5   t        j                  d      |      }t        j                  dt        j                        }t        j                  d      }t        j                  d      }t        j                  d      } |d   ||||       | j                  |d	          d d d         y # 1 sw Y   xY w)
Nr   void(b1[:], f2, f2, f2)r`   r}          @      @g      @ra   r   )rG   rJ   rL   rN   rP   r   r   rb   rV   r   r   r   
assertTruerY   r   r   compiledr   r   r   arg3s           r   !test_multiple_float16_comparisonsz4TestOperatorModule.test_multiple_float16_comparisons'  s    )))))	+	
  	(B$ (>488$=>rBhhq1zz"~zz"~zz"~sD$5A'( (	(( (   B"C  C)	c                    t         t        t        t        t        f}|D ]  }| j                  |      5   t        j                  d      |      }t        j                  dt        j                        }t        j                  d      }t        j                  d      }t        j                  d      } |d   ||||       | j                  |d	          d d d         y # 1 sw Y   xY w)
Nr   r   r`   r}   r   r   g      ?ra   r   )rG   rJ   rL   rN   rP   r   r   rb   rV   r   r   r   assertFalser   s           r   'test_multiple_float16_comparisons_falsez:TestOperatorModule.test_multiple_float16_comparisons_false8  s    )))))	+	
  	)B$ )>488$=>rBhhq1zz"~zz"~zz"~sD$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 ]D  \  }}}| j!                  |      5  t#        ||d      \  }}	| j%                  ||       d d d        F y # 1 sw Y   QxY w)N)setp.gt.f16setp.ge.f16setp.lt.f16setp.le.f16setp.eq.f16setp.ne.f16r|   r   r   )r-   r/   r2   r4   r6   r8   rj   r   r   r   r   r   r   r	   r   r   r   r
   r   )
rY   r   r   opstringr   r   r^   sr   r   s
             r   test_fp16_comparison_ptxz+TestOperatorModule.test_fp16_comparison_ptxI  s    #^#^#^5	 {{HKKhkk{{HKK)2 1r2YX6 	&IBA$ &$R&9Qa%& &	&& &s   9$C((C1	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 ]m  \  }}| j                  |      5  t        d d  t         t#        t$        j&                        f}t)        ||d	      \  }}| j+                  ||   |       d d d        o y # 1 sw Y   zxY w)
Nr   r   r   r   r   r   r|   r   r   )r-   r/   r2   r4   r6   r8   rj   r   r   r   r   r   r   r   r   r	   r   r   rV   r   r
   r   )	rY   r   r   r   r   r^   r   r   r   s	            r   test_fp16_int8_comparison_ptxz0TestOperatorModule.test_fp16_int8_comparison_ptxZ  s	    $^#^#^5	 {{HKKhkk{{HKK) KKKKKKKKKKKK/ )S) 	1FB$ 11r:bgg#67$R&9QhrlC01 1	11 1s   AEE%	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$                  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 ]  \  \  }}}| j/                  ||      5  t        j0                  t        j2                  |      }	t4        d d  t6        t9        |	      f}
t;        ||
d      \  }}||   ||	   z   }| j=                  ||       d d d         y # 1 sw Y   xY w)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-   r/   r2   r4   r6   r8   rj   r   r   r   r   r   r   rV   r   r   r   r   r   r~   r   r   r   r   r   r   r	   r   r   r
   r   )rY   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	 {{HKKhkk{{HKK) 288RXXRZZ1KK
KK
KK
KK
KK
KK- HHW%uHHW%uHHY'HHY'0
 &--c)S.A.;= 	(LHRb+ (..R81r:g#67$R&9QrlXg%66c3'( (	(( (s   >A,H55H>	)__name__
__module____qualname__rU   rh   rm   rq   rt   rw   rz   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   __classcell__)rZ   s   @r   rR   rR   o   s   /---12 : :" ?@* A* : :& ?@* A* : :  : : ?@& A&
 ?@& A& 3 3& 3 3, ( (  ) )  ?@& A&  ?@1 A1* ?@( A(r   rR   __main__)-numpyrV   numba.cuda.testingr   r   r   r   numbar   numba.core.typesr   r	   
numba.cudar
   rj   r   numba.np.numpy_supportr   r   r   r   r   r    r"   r$   r&   r(   r+   r-   r/   r2   r4   r6   r8   rb   r?   rB   rG   rJ   rL   rN   rP   rR   r   mainr   r   r   <module>r      s    1 1  # "   - 
,t$ % 
,t$ %1
&
'


^( ^(B	 zHMMO r   