
    shs                        S SK r S SKrS SKrS SKrS SKJrJr  S SKJ	r	  S SK
Jr  S SKJr  S SKJrJrJrJr  S rS rS	 rS
 rS rS rS rS rS rS rS rS rS rS r S r!S r"S r#S r$S r%S r&S r'S r(S r)S r*S r+S  r,S! r-S" r.S# r/S$ r0\Rb                  " S%S&9S' 5       r2\Rb                  " S%S&9S( 5       r3S) r4S* r5S+ r6S, r7S- r8S. r9S/ r:S0 r;S1 r<S2 r=S3 r>S4 r?S5 r@S6 rAS7 rBS8 rCS9 rDS: rES; rFS< rGS= rHS> rIS? rJS@ rKSA rLSB rMSC rNSD rOSE rPSF rQSG rRSH rSSI rTSJ rU " SK SL\5      rV\WSM:X  a  \R                  " 5         gg)N    N)cudaint64)compile_ptx)TypingError)f2)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_53c                 @    [         R                  R                  nXS'   g Nr   r   	threadIdxxaryis     {/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pysimple_threadidxr          AF    c                 @    [         R                  R                  nXU'   g Nr   r   s     r   fill_threadidxr      r   r   c                     [         R                  R                  n[         R                  R                  n[         R                  R                  nUS-   US-   -  US-   -  XX#4'   g N   )r   r   r   yz)r   r   jks       r   fill3d_threadidxr"      sP    AAAEa!e$A.C1Lr   c                 8    [         R                  " S5      nXU'   g r   r   gridr   s     r   simple_grid1dr&      s    		!AFr   c                 D    [         R                  " S5      u  pX-   XU4'   g N   r$   )r   r   r    s      r   simple_grid2dr*   $   s    99Q<DAC1Ir   c                 r    [         R                  " S5      n[         R                  " S5      nUS:X  a  X S'   g g Nr   r   r   r%   gridsize)r   r   r   s      r   simple_gridsize1dr/   )   s0    		!AaAAvA r   c                     [         R                  " S5      u  p[         R                  " S5      u  p4US:X  a  US:X  a	  X0S'   X@S'   g g g )Nr)   r   r   r-   )r   r   r    r   r   s        r   simple_gridsize2dr1   0   sC    99Q<DA==DAAv!q&AA vr   c                    [         R                  " S5      u  p[         R                  R                  [         R                  R                  -  n[         R                  R
                  [         R                  R
                  -  nU R                  u  pV[        XU5       H  n[        X%U5       H  nXx-   XU4'   M     M      g r(   )r   r%   gridDimr   blockDimr   shaperange)	cstartXstartYgridXgridYheightwidthr   r   s	            r   intrinsic_forloop_stepr>   8   s    YYq\NFLLNNT]]__,ELLNNT]]__,EGGMF6%(vu-AeAdG . )r   c                 6    [         R                  " U5      U S'   g r   )r   popcr   r7   s     r   simple_popcrB   C       YYq\CFr   c                 8    [         R                  " XU5      U S'   g r   )r   fmar   abr7   s       r   
simple_fmarI   G   s    XXaACFr   c                 V    [         R                  R                  US   US   5      U S'   g r   r   fp16haddr   rG   rH   s      r   simple_haddrO   K   "    YY^^AaD!A$'CFr   c                 H    [         R                  R                  X5      U S'   g r   rK   rN   s      r   simple_hadd_scalarrR   O       YY^^A!CFr   c                 ^    [         R                  R                  US   US   US   5      U S'   g r   r   rL   hfmarF   s       r   simple_hfmarW   S   s(    YY^^AaD!A$!-CFr   c                 J    [         R                  R                  XU5      U S'   g r   rU   rF   s       r   simple_hfma_scalarrY   W   s    YY^^A!$CFr   c                 V    [         R                  R                  US   US   5      U S'   g r   r   rL   hsubrN   s      r   simple_hsubr]   [   rP   r   c                 H    [         R                  R                  X5      U S'   g r   r[   rN   s      r   simple_hsub_scalarr_   _   rS   r   c                 V    [         R                  R                  US   US   5      U S'   g r   r   rL   hmulrN   s      r   simple_hmulrc   c   rP   r   c                 H    [         R                  R                  X5      U S'   g r   ra   rN   s      r   simple_hmul_scalarre   g   rS   r   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   hdivrN   s      r   simple_hdiv_scalarrh   k   rS   r   c                     [         R                  " S5      nX0R                  :  a*  X   nX#   n[         R                  R	                  XE5      X'   g g r   )r   r%   sizerL   rg   )r   array_aarray_br   rG   rH   s         r   simple_hdiv_kernelrm   o   s?    		!A88|JJ% r   c                 N    [         R                  R                  US   5      U S'   g r   r   rL   hnegr   rG   s     r   simple_hnegrr   w       YY^^AaD!CFr   c                 H    [         R                  R                  U5      U S'   g r   ro   rq   s     r   simple_hneg_scalarru   {       YY^^ACFr   c                 N    [         R                  R                  US   5      U S'   g r   r   rL   habsrq   s     r   simple_habsrz      rs   r   c                 H    [         R                  R                  U5      U S'   g r   rx   rq   s     r   simple_habs_scalarr|      rv   r   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   heqrN   s      r   simple_heq_scalarr          YY]]1 CFr   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   hnerN   s      r   simple_hne_scalarr      r   r   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   hgerN   s      r   simple_hge_scalarr      r   r   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   hgtrN   s      r   simple_hgt_scalarr      r   r   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   hlerN   s      r   simple_hle_scalarr      r   r   c                 H    [         R                  R                  X5      U S'   g r   r   rL   hltrN   s      r   simple_hlt_scalarr      r   r   T)devicec                 @    [         R                  R                  X5      $ r   r   r   r   s     r   
hlt_func_1r          99==r   c                 @    [         R                  R                  X5      $ r   r   r   s     r   
hlt_func_2r      r   r   c                 D    [        X5      =(       a    [        X#5      U S'   g r   )r   r   rrG   rH   r7   s       r   test_multiple_hcmp_1r      s    a0
1 0AaDr   c                 l    [        X5      =(       a    [        R                  R                  X#5      U S'   g r   )r   r   rL   r   r   s       r   test_multiple_hcmp_2r      "    a3		a 3AaDr   c                 l    [        X5      =(       a    [        R                  R                  X25      U S'   g r   )r   r   rL   r   r   s       r   test_multiple_hcmp_3r      r   r   c                     [         R                  R                  X5      =(       a    [         R                  R                  X#5      U S'   g r   r   r   s       r   test_multiple_hcmp_4r      *    99==6499==#6AaDr   c                     [         R                  R                  X5      =(       a    [         R                  R                  X25      U S'   g r   )r   rL   r   r   r   s       r   test_multiple_hcmp_5r      r   r   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   hmaxrN   s      r   simple_hmax_scalarr      rS   r   c                 H    [         R                  R                  X5      U S'   g r   )r   rL   hminrN   s      r   simple_hmin_scalarr      rS   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   lenrL   hsinr   r   r   s      r   simple_hsinr      5    		!A3q6zyy~~ad# r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hcosr   s      r   simple_hcosr      r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hlogr   s      r   simple_hlogr      r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hlog2r   s      r   simple_hlog2r      5    		!A3q6zyyqt$ r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hlog10r   s      r   simple_hlog10r      7    		!A3q6zyy% r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hexpr   s      r   simple_hexpr      r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hexp2r   s      r   simple_hexp2r      r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hsqrtr   s      r   simple_hsqrtr      r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hrsqrtr   s      r   simple_hrsqrtr     s7    		!A3q6zyy% r   c                     U S-  $ )Ng      ࿩ )r   dtypes     r   numpy_hrsqrtr   
  s    9r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hceilr   s      r   simple_hceilr     r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hfloorr   s      r   simple_hfloorr     r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hrcpr   s      r   simple_hrcpr     r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   htruncr   s      r   simple_htruncr   #  r   r   c                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hrintr   s      r   simple_hrintr   *  r   r   c                 6    [         R                  " U5      U S'   g r   )r   cbrtrq   s     r   simple_cbrtr   1  rC   r   c                 6    [         R                  " U5      U S'   g r   )r   brevrA   s     r   simple_brevr   5  rC   r   c                 6    [         R                  " U5      U S'   g r   )r   clzrA   s     r   
simple_clzr   9      XXa[CFr   c                 6    [         R                  " U5      U S'   g r   )r   ffsrA   s     r   
simple_ffsr   =  r   r   c                      [        U5      U S'   g r   roundrA   s     r   simple_roundr   A  s    1XCFr   c                      [        X5      U S'   g r   r   )r   r7   ndigitss      r   simple_round_tor   E  s    1CFr   c                 r    [         R                  " S5      nX   S:  a  US-  S:X  a  X#   X'   g SX'   g SX'   g )Nr      r)   r         r$   )rG   rH   r7   r   s       r   branching_with_ifsr   I  s:    		!Ataxq5A:4ADADr   c                     [         R                  " S5      n[         R                  " US-  S:H  X#   S5      n[         R                  " X   S:  US5      X'   g )Nr   r)   r   r   r   r   )r   r%   selp)rG   rH   r7   r   inners        r   branching_with_selpsr   U  sE    		!AIIa!eqj!$+E99QTAXua(ADr   c                 T    [         R                  " S5      n[         R                  X'   g r   )r   r%   laneidr   s     r   simple_laneidr   \  s    		!A[[CFr   c                 *    [         R                  U S'   g r   )r   warpsize)r   s    r   simple_warpsizer  a  s    ]]CFr   c                 0    [         R                  " U 5        g r   r$   r   s    r   nonliteral_gridr  e  s    IIaLr   c                 0    [         R                  " U 5        g r   )r   r.   r  s    r   nonliteral_gridsizer  i  s    MM!r   c                     ^  \ rS rSrU 4S jrS rS rS r\" S5      S 5       r	\" S5      S 5       r
S	 rS
 rS r\" S5      S 5       r\" S5      S 5       rS rS rS rS rS rS rS rS r\S 5       r\S 5       r\" S5      S 5       r\S 5       r\S 5       r\" S5      S 5       r\S 5       r\S  5       r \" S5      S! 5       r!\S" 5       r"\S# 5       r#\" S5      S$ 5       r$\S% 5       r%\S& 5       r&\S' 5       r'\S( 5       r(\" S5      S) 5       r)\S* 5       r*\S+ 5       r+\" S5      S, 5       r,\S- 5       r-\S. 5       r.\S/ 5       r/\S0 5       r0\S1 5       r1\S2 5       r2S3 r3S4 r4S5 r5\" S65      S7 5       r6S8 r7S9 r8S: r9S; r:\" S65      S< 5       r;S= r<S> r=S? r>S@ r?\" S65      SA 5       r@SB rASC rBSD rCSE rDSF rE\" SG5      SH 5       rFSI rGSJ rH\" SG5      SK 5       rISL rJSMrKU =rL$ )NTestCudaIntrinsicim  c                 `   > [         TU ]  5         [        R                  R	                  S5        g r   )supersetUpnprandomseed)self	__class__s    r   r  TestCudaIntrinsic.setUpn  s    
		qr   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " U5        U R                  US   S:H  5        g )Nvoid(int32[:])r   r   r   r   r   )r   jitr   r  onesint32
assertTruer  compiledr   s      r   test_simple_threadidx'TestCudaIntrinsic.test_simple_threadidxr  sK    88,-.>?ggarxx(sA!$r   c                 <   [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9nUSU4   " U5        U R                  [        R                  " X4:H  5      5        g )Nr  
   r  r   )	r   r  r   r  r  r  aranger  all)r  r  Nr   exps        r   test_fill_threadidx%TestCudaIntrinsic.test_fill_threadidxx  sj    88,-n=ggarxx(ii*Assz*+r   c                    ^^^ Su  mmmUUU4S jnUUU4S jnU" 5       nU" 5       nU R                  [        R                  " X4:H  5      5        g )N)r         c                     > [         R                  " S5      " [        5      n [        R                  " TTT4[        R
                  S9nU STTT44   " U5        U$ )Nzvoid(int32[:,:,::1])r  r   )r   r  r"   r  zerosr  r  r   XYZs     r   c_contigous<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.c_contigous  sN    xx 678HIH((Aq!9BHH5CQAq	\"3'Jr   c                     > [         R                  " S5      " [        5      n [        R                  " [        R
                  " TTT4[        R                  S95      nU STTT44   " U5        U$ )Nzvoid(int32[::1,:,:])r  r   )r   r  r"   r  asfortranarrayr*  r  r+  s     r   f_contigous<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.f_contigous  sY    xx 678HIH##BHHaAYbhh$GHCQAq	\"3'Jr   )r  r  r!  )r  r/  r3  c_resf_resr,  r-  r.  s        @@@r   test_fill3d_threadidx'TestCudaIntrinsic.test_fill3d_threadidx  s?    1a		 u~./r   zCudasim does not check typesc                     U R                  [        S5         [        R                  " S5      " [        5        S S S 5        g ! , (       d  f       g = fNRequireLiteralValuezvoid(int32))assertRaisesRegexr   r   r  r  r  s    r   test_nonliteral_grid_error,TestCudaIntrinsic.test_nonliteral_grid_error  s1    ##K1FGHH]#O4 HGG   !A
Ac                     U R                  [        S5         [        R                  " S5      " [        5        S S S 5        g ! , (       d  f       g = fr:  )r<  r   r   r  r  r=  s    r   test_nonliteral_gridsize_error0TestCudaIntrinsic.test_nonliteral_gridsize_error  s2    ##K1FGHH]#$78 HGGr@  c                 *   [         R                  " S5      " [        5      nSu  p#X#-  n[        R                  " U[        R
                  S9nXU4   " U5        U R                  [        R                  " U[        R                  " U5      :H  5      5        g )Nvoid(int32[::1])r      r  )	r   r  r&   r  emptyr  r  r!  r   )r  r  ntidnctaidnelemr   s         r   test_simple_grid1d$TestCudaIntrinsic.test_simple_grid1d  sj    88./>hhuBHH-s#sbii&6678r   c                    [         R                  " S5      " [        5      nSnSnUS   US   -  US   US   -  4n[        R                  " U[        R
                  S9nUR                  5       nXU4   " U5        [        UR                  S   5       H*  n[        UR                  S   5       H  nXx-   XgU4'   M     M,     U R                  [        R                  " XV:H  5      5        g Nzvoid(int32[:,::1])r   r   r'  r(  r   r   r  )r   r  r*   r  rH  r  copyr6   r5   r  r!  )	r  r  rI  rJ  r5   r   r#  r   r    s	            r   test_simple_grid2d$TestCudaIntrinsic.test_simple_grid2d  s    8801-@a6!9$d1gq	&9:hhuBHH-hhjs#syy|$A399Q<(EqD	 ) % 	sz*+r   c                     [         R                  " S5      " [        5      nSu  p#[        R                  " S[        R
                  S9nXU4   " U5        U R                  US   X2-  5        g )NrE  rF  r   r  r   )r   r  r/   r  r*  r  assertEqualr  r  rI  rJ  r   s        r   test_simple_gridsize1d(TestCudaIntrinsic.test_simple_gridsize1d  sW    88./0ABhhq)s#Q/r   zRequires too many threadsc                 &   [         R                  S 5       n[        R                  " S[        R                  S9n[        R                  " S[        R                  S9nUS   " X#5        U R                  US   S5        U R                  US   S5        g )Nc                    [         R                  " S5      n[         R                  R                  [         R                  R                  -  [         R
                  R                  -   n[         R                  " S5      n[         R                  R                  [         R                  R                  -  nX#:w  a  SU S'   XE:w  a  SUS'   g g r,   )r   r%   blockIdxr   r4   r   r.   r3   )
grid_errorgridsize_errori1i2gs1gs2s         r   f,TestCudaIntrinsic.test_issue_9229.<locals>.f  s    1B4==??2T^^5E5EEB--"C--//DLLNN2Cx !
1z$%q! r   r   r  )i Q   r   )r   r  r  r*  uint64rV  )r  rc  r]  r^  s       r   test_issue_9229!TestCudaIntrinsic.test_issue_9229  s|     
	& 
	& XXaryy1
!2995 	
-4A**A.r   zTests PTX emissionc           	      h   [         S S  [         [         S S  4n[        R                  " U5      " [        5      n[        R                  " U5      " [        5      nSnSn[
        R                  " SS[
        R                   S9nUR                  5       nSUS S& [
        R                  " U[
        R                   S9nX$S4   " XU5        UR                  U5      n	U R                  S	[        [        R                  " S
U	5      5      5        [
        R                  R                  XSS9  [
        R                  " U[
        R                   S9nX4S4   " XU5        UR                  U5      n	U R                  S[        [        R                  " S
U	5      5      5        [
        R                  R                  XSS9  g )N    r(     )r5   
fill_valuer   r   r'  r  r   r)   z	\s+bra\s+	branching)err_msgr   r   )r   r   r  r   r   r  fullrR  r   inspect_asmrV  r   refindalltestingassert_array_equal)
r  sigcu_branching_with_ifscu_branching_with_selpsnrH   r7   expectedrG   ptxs
             r   	test_selpTestCudaIntrinsic.test_selp  sL   Qxa) $.@ A"&((3-0D"EGG"288<668!IIarxx(d#A!,#//4C

< =>?


%%a;%GIIarxx(1%aA.%11#6C

< =>?


%%a6%Br   c                 "   [         R                  " S5      " [        5      nSnSn[        R                  " S[        R
                  S9nXU4   " U5        U R                  US   US   US   -  5        U R                  US   US   US   -  5        g )NrE  rP  rQ  r)   r  r   r   )r   r  r1   r  r*  r  rV  rW  s        r   test_simple_gridsize2d(TestCudaIntrinsic.test_simple_gridsize2d  s    88./0ABhhq)s#QT!W!45QT!W!45r   c           	         [         R                  " S5      " [        5      nSnSnUS   US   -  US   US   -  4n[        R                  " U[        R
                  S9nXU4   " U5        Uu  pgUR                  u  p[        [        US   5      [        US   5      5       HU  u  pXj-   X{-   p[        XU5       H9  n[        XU5       H&  nU R                  X_U4   X-   :H  X_U4   X-   45        M(     M;     MW     g rO  )
r   r  r>   r  rH  r  r5   zipr6   r  )r  r  rI  rJ  r5   r   r:   r;   r<   r=   r   r    r8   r9   r   r   s                   r   test_intrinsic_forloop_step-TestCudaIntrinsic.test_intrinsic_forloop_step  s    88012HIa6!9$d1gq	&9:hhuBHH-s#		d1gd1g7DA"Y	F6%0vu5AOOC1I$6TAE8JK 6 1 8r   c                     [         R                  S 5       n[        R                  " S[        R                  S9R                  SSS5      nUS   " U5        [        R                  R                  US5        g )Nc                 ~    [         R                  " S5      u  pn[         R                  " S5      u  pEnXE-  U-  XX#4'   g Nr   r-   )outr   r   r   rG   rH   r7   s          r   foo*TestCudaIntrinsic.test_3dgrid.<locals>.foo	  s6    iilGA!mmA&GA!519C1Lr   i  r  	   )r   r   r   r  )r   r  r  r*  r  reshapers  assert_equal)r  r  arrs      r   test_3dgridTestCudaIntrinsic.test_3dgrid  s_    		% 
	%
 hhvRXX.66q!Q? !#&


V,r   c                 
   [         R                  S 5       nSu  p#n[        R                  " X#-  U-  [        R                  S9R                  X#U5      nUS   " U5        U R                  [        R                  " U5      5        g )Nc                     [         R                  " S5      u  pn[         R                  " S5      u  pEnU[         R                  R                  [         R
                  R                  [         R                  R                  -  -   :H  =(       a    U[         R                  R                  [         R
                  R                  [         R                  R                  -  -   :H  =(       aS    U[         R                  R                  [         R
                  R                  [         R                  R                  -  -   :H  nU[         R                  R                  [         R                  R                  -  :H  =(       aw    U[         R                  R                  [         R                  R                  -  :H  =(       a8    U[         R                  R                  [         R                  R                  -  :H  nU=(       a    UXX#4'   g r  )
r   r%   r.   r   r   r\  r4   r   r   r3   )	r  r   r   r   rG   rH   r7   grid_is_rightgridsize_is_rights	            r   r  ,TestCudaIntrinsic.test_3dgrid_2.<locals>.foo  s7   iilGA!mmA&GA!T^^%%$--//(III JT^^%%$--//(IIIJT^^%%$--//(III 
 "#dmmoo&F!F "G!"dmmoo&F!F"G!"dmmoo&F!F  )>->C1Lr   )   r(     r  ))r   r   r)   )r   r)   r   )r   r  r  r*  bool_r  r  r!  )r  r  r   r   r   r  s         r   test_3dgrid_2TestCudaIntrinsic.test_3dgrid_2  sl    		? 
	? (ahh	"((3;;A!D !#&s$r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g )Nvoid(int32[:], uint32)r   r  r     r   r   r   r  rB   r  r*  r  rV  r  s      r   test_popc_u4TestCudaIntrinsic.test_popc_u4)  sL    8845kBhhq)sD!Q#r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g )Nzvoid(int32[:], uint64)r   r  r  l        @ r   r   r  r  s      r   test_popc_u8TestCudaIntrinsic.test_popc_u8/  sL    8845kBhhq)sN+Q#r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " USSS5        [        R                  R                  US   S	5        g )
Nzvoid(f4[:], f4, f4, f4)r   r  r         @      @      @r   r  )r   r  rI   r  r*  float32rs  assert_allcloser  s      r   test_fma_f4TestCudaIntrinsic.test_fma_f45  T    8856zBhhq

+sBB'


""3q695r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " USSS5        [        R                  R                  US   S	5        g )
Nzvoid(f8[:], f8, f8, f8)r   r  r  r  r  r  r   r  )r   r  rI   r  r*  float64rs  r  r  s      r   test_fma_f8TestCudaIntrinsic.test_fma_f8;  r  r   c                 |   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  " S/[        R
                  S9n[        R                  " S/[        R
                  S9nUS   " X#U5        [        R                  R                  US   X4-   5        g Nvoid(f2[:], f2[:], f2[:])r   r  r  r  r  r   )	r   r  rO   r  r*  float16arrayrs  r  r  r  r   arg1arg2s        r   	test_haddTestCudaIntrinsic.test_haddA  ~    8878Ehhq

+xxBJJ/xxBJJ/s$'


""3q64;7r   c                 H   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      n[        R
                  " S5      nUS   " X#U5        X4-   n[        R                  R                  US   U5        g )Nvoid(f2[:], f2, f2)r   r  JM!	@r  r  r   )r   r  rR   r  r*  r  rs  r  r  r  r   r  r  refs         r   test_hadd_scalar"TestCudaIntrinsic.test_hadd_scalarJ  ss    88123EFhhq

+zz)$zz"~s$'k


""3q63/r   z(Compilation unsupported in the simulatorc                 r    [         S S  [         [         4n[        [        USS9u  p#U R                  SU5        g )Nr'  r   cczadd.f16)r   r   rR   assertInr  argsrz  _s       r   test_hadd_ptxTestCudaIntrinsic.test_hadd_ptxT  0    1r2/&Ai%r   c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  " S/[        R
                  S9n[        R                  " S/[        R
                  S9n[        R                  " S/[        R
                  S9nUS   " X#XE5        [        R                  R                  US   X4-  U-   5        g )	Nz void(f2[:], f2[:], f2[:], f2[:])r   r  r  r  r  r  r   )	r   r  rW   r  r*  r  r  rs  r  )r  r  r   r  r  arg3s         r   	test_hfmaTestCudaIntrinsic.test_hfmaZ  s    88>?Lhhq

+xxBJJ/xxBJJ/xxBJJ/s$-


""3q64;+=>r   c                 z   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      n[        R
                  " S5      n[        R
                  " S5      nUS   " X#XE5        X4-  U-   n[        R                  R                  US   U5        g )	Nzvoid(f2[:], f2, f2, f2)r   r  r  r  r  r  r   )r   r  rY   r  r*  r  rs  r  )r  r  r   r  r  r  r  s          r   test_hfma_scalar"TestCudaIntrinsic.test_hfma_scalard  s    88567IJhhq

+zz"~zz"~zz"~s$-kD 


""3q63/r   c                 |    [         S S  [         [         [         4n[        [        USS9u  p#U R                  SU5        g )Nr  r  z
fma.rn.f16)r   r   rY   r  r  s       r   test_hfma_ptxTestCudaIntrinsic.test_hfma_ptxo  s3    1r2r"/&AlC(r   c                 |   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  " S/[        R
                  S9n[        R                  " S/[        R
                  S9nUS   " X#U5        [        R                  R                  US   X4-
  5        g r  )	r   r  r]   r  r*  r  r  rs  r  r  s        r   	test_hsubTestCudaIntrinsic.test_hsubu  r  r   c                 H   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      n[        R
                  " S5      nUS   " X#U5        X4-
  n[        R                  R                  US   U5        g Nr  r   r  r  gQ?r  r   )r   r  r_   r  r*  r  rs  r  r  s         r   test_hsub_scalar"TestCudaIntrinsic.test_hsub_scalar~  t    88123EFhhq

+zz)$zz$s$'k


""3q63/r   c                 r    [         S S  [         [         4n[        [        USS9u  p#U R                  SU5        g )Nr  r  zsub.f16)r   r   r_   r  r  s       r   test_hsub_ptxTestCudaIntrinsic.test_hsub_ptx  r  r   c                 z   [         R                  " 5       " [        5      n[        R                  " S[        R
                  S9n[        R                  " S/[        R
                  S9n[        R                  " S/[        R
                  S9nUS   " X#U5        [        R                  R                  US   X4-  5        g )Nr   r  r  r  r  r   )	r   r  rc   r  r*  r  r  rs  r  r  s        r   	test_hmulTestCudaIntrinsic.test_hmul  sz    88:k*hhq

+xxBJJ/xxBJJ/s$'


""3q64;7r   c                 H   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      n[        R
                  " S5      nUS   " X#U5        X4-  n[        R                  R                  US   U5        g r  )r   r  re   r  r*  r  rs  r  r  s         r   test_hmul_scalar"TestCudaIntrinsic.test_hmul_scalar  r  r   c                 r    [         S S  [         [         4n[        [        USS9u  p#U R                  SU5        g )Nr  r  zmul.f16)r   r   re   r  r  s       r   test_hmul_ptxTestCudaIntrinsic.test_hmul_ptx  r  r   c                 H   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      n[        R
                  " S5      nUS   " X#U5        X4-  n[        R                  R                  US   U5        g r  )r   r  rh   r  r*  r  rs  r  r  s         r   test_hdiv_scalar"TestCudaIntrinsic.test_hdiv_scalar  st    88123EFhhq

+zz)$zz$s$'k


""3q63/r   c                    [         R                  " S5      " [        5      n[        R                  R                  SSSS9R                  [        R                  5      n[        R                  R                  SSSS9R                  [        R                  5      n[        R                  " U[        R                  S9nUR                  UR                  5      " XBU5        X#-  n[        R                  R                  XE5        g )Nr  i    i  rj   r  )r   r  rm   r  r  randintastyper  
zeros_likeforallrj   rs  r  )r  r  arry1arry2r   r  s         r   	test_hdivTestCudaIntrinsic.test_hdiv  s    88789KL		!!&%c!:AA"**M		!!&%c!:AA"**MmmE4!#e4m


""3,r   c                 0   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  " S/[        R
                  S9nUS   " X#5        [        R                  R                  US   U* 5        g )Nvoid(f2[:], f2[:])r   r  r  r  r   )	r   r  rr   r  r*  r  r  rs  r  r  r  r   r  s       r   	test_hnegTestCudaIntrinsic.test_hneg  sf    8801+>hhq

+xxBJJ/s!


""3q6D51r   c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      nUS   " X#5        U* n[        R                  R                  US   U5        g )Nvoid(f2[:], f2)r   r  r  r  r   )r   r  ru   r  r*  r  rs  r  r  r  r   r  r  s        r   test_hneg_scalar"TestCudaIntrinsic.test_hneg_scalar  sd    88-./ABhhq

+zz)$s!e


""3q63/r   c                 h    [         S S  [         4n[        [        USS9u  p#U R                  SU5        g )Nr  r  zneg.f16)r   r   ru   r  r  s       r   test_hneg_ptxTestCudaIntrinsic.test_hneg_ptx  .    1r{/&Ai%r   c                 >   [         R                  " 5       " [        5      n[        R                  " S[        R
                  S9n[        R                  " S/[        R
                  S9nUS   " X#5        [        R                  R                  US   [        U5      5        g )Nr   r        r  r   )
r   r  rz   r  r*  r  r  rs  r  absr  s       r   	test_habsTestCudaIntrinsic.test_habs  sd    88:k*hhq

+xxRZZ0s!


""3q63t95r   c                 (   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      nUS   " X#5        [        U5      n[        R                  R                  US   U5        g )Nr   r   r  gJM!	r  r   )	r   r  r|   r  r*  r  r
  rs  r  r  s        r   test_habs_scalar"TestCudaIntrinsic.test_habs_scalar  sf    88-./ABhhq

+zz*%s!$i


""3q63/r   c                 h    [         S S  [         4n[        [        USS9u  p#U R                  SU5        g )Nr  r  zabs.f16)r   r   r|   r  r  s       r   test_habs_ptxTestCudaIntrinsic.test_habs_ptx  r  r   c                    [         [        [        [        [        [
        [        [        [        [        [        [        4n[        [        4n[        R                  [        R                   [        R"                  [        R$                  [        R&                  [        R(                  [        R*                  [        R,                  [        R.                  [        R0                  [        R2                  [4        4n[        R6                  [        R8                  4nSn[        R:                  R=                  S5        [        R:                  R?                  SSUS9RA                  [        RB                  5      n[        RD                  " U5      n[G        X5       Hz  u  pU RI                  U	S9   [J        RL                  " S5      " U5      nUSU4   " Xv5        U	" U[        RB                  S9n
[        RN                  RQ                  Xz5        S S S 5        M|     [        R:                  R?                  SSUS9RA                  [        RB                  5      n[G        X$5       Hz  u  pU RI                  U	S9   [J        RL                  " S5      " U5      nUSU4   " X{5        U	" U[        RB                  S9n
[        RN                  RQ                  Xz5        S S S 5        M|     g ! , (       d  f       GMT  = f! , (       d  f       M  = f)	Nrj  r   r  r  fnr  r  r  ))r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  sincosloglog2log10sqrtceilfloor
reciprocaltruncrintr   r#  exp2r  r  r  r  r  r  r  subTestr   r  rs  r  )r  kernelsexp_kernelsexpected_functionsexpected_exp_functionsr"  r   r   kernelr  ry  x2s               r   test_fp16_intrinsics_common-TestCudaIntrinsic.test_fp16_intrinsics_common  s    m}| 	"
 #L1 ffbff ffbggrxx ggrww mmRXXrww*	,
 #%&&"''!2 
		qIIaQ/66rzzBMM!g:JF$"67?qsA!arzz2

**17	 %$ ; YYq"1-44RZZ@kBJF$"67?qsA"b

3

**17	 %$ C %$ %$s   )AK.AK,
K)	,
K;	c                    [         R                  " 5       S 5       nSn[        R                  R	                  S5        [        R                  R                  U5      R                  [        R                  5      n[        R                  " U5      nUSU4   " XC5        [        R                  R                  USU-  5        g )Nc                     [         R                  " S5      nU[        U 5      :  a$  [         R                  R	                  X   5      X'   g g r   )r   r%   r   rL   hexp10r   s      r   hexp10_vectors5TestCudaIntrinsic.test_hexp10.<locals>.hexp10_vectors  s7    		!A3q6zyy''- r   rj  r   r  )r   r  r  r  r  randr  r  r  rs  r  )r  r.  r"  r   r   s        r   test_hexp10TestCudaIntrinsic.test_hexp10  s    		. 
	. 
		qIINN1$$RZZ0MM! 	q!tQ"


""1bAg.r   c                    [         [        [        [        [        [
        4n[        R                  [        R                  [        R                  [        R                  [        R                  [        R                  4n[        X5       GH:  u  p4U R                  US9   [        R                   " S5      " U5      n["        R$                  " S["        R&                  S9n["        R$                  " S["        R&                  S9n["        R(                  " S5      n["        R(                  " S5      n	["        R(                  " S5      n
US   " XyU	5        U" X5      nU R+                  XgS	   5        US   " XyU
5        U" X5      nU R+                  XgS	   5        US   " XyU5        U" X5      nU R+                  XgS	   5        S S S 5        GM=     g ! , (       d  f       GMP  = f)
N)opzvoid(b1[:], f2, f2)r   r  r)   r   r   r  r   )r   r   r   r   r   r   operatoreqnegegtleltr  r"  r   r  r  r*  r  r  rV  )r  fnsopsr  r4  r'  ry  gotr  r  arg4s              r   test_fp16_comparison&TestCudaIntrinsic.test_fp16_comparison!  s\    "35F "35FH{{HKK{{HKK6 #mFB$"78<88ARXX6hhq1zz!}zz!}zz!} tS-d>  q62 tS-d>  q62 tS-d>  q62- %$ $$$s   DG
G	c                    [         [        [        [        [        4nU H  nU R                  US9   [        R                  " S5      " U5      n[        R                  " S[        R                  S9n[        R                  " S5      n[        R                  " S5      n[        R                  " S5      nUS   " XEXg5        U R                  US	   5        S S S 5        M     g ! , (       d  f       M  = f)
Nr  zvoid(b1[:], f2, f2, f2)r   r  r  r  r  r  r   )r   r   r   r   r   r"  r   r  r  r*  r  r  r  )r  	functionsr  r  r   r  r  r  s           r   !test_multiple_float16_comparisons3TestCudaIntrinsic.test_multiple_float16_comparisonsA  s    )))))	+	
 B$88$=>rBhhq1zz"~zz"~zz"~s$5A' %$ $$s   B"C
C.	c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      n[        R
                  " S5      nUS   " X#U5        [        R                  R                  US   U5        [        R
                  " S5      nUS   " X#U5        [        R                  R                  US   U5        g 	Nr  r   r  r  r  r  r   g      @)r   r  r   r  r*  r  rs  r  r  s        r   	test_hmaxTestCudaIntrinsic.test_hmaxR      88123EFhhq

+zz"~zz"~s$'


""3q640zz"~s$'


""3q640r   c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R
                  " S5      n[        R
                  " S5      nUS   " X#U5        [        R                  R                  US   U5        [        R
                  " S5      nUS   " X#U5        [        R                  R                  US   U5        g rG  )r   r  r   r  r*  r  rs  r  r  s        r   	test_hminTestCudaIntrinsic.test_hmin^  rJ  r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nSnUS   " X#5        [        R                  R                  US   US-  5        g )Nzvoid(float32[:], float32)r   r  r  r  r   UUUUUU?)r   r  r   r  r*  r  rs  r  r  r  r   cbrt_args       r   test_cbrt_f32TestCudaIntrinsic.test_cbrt_f32j  X    8878Ehhq

+s%


""3q68+>?r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nSnUS   " X#5        [        R                  R                  US   US-  5        g )Nzvoid(float64[:], float64)r   r  g      @r  r   rO  )r   r  r   r  r*  r  rs  r  rP  s       r   test_cbrt_f64TestCudaIntrinsic.test_cbrt_f64q  rT  r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g )Nzvoid(uint32[:], uint32)r   r  r  i0  r   i  )r   r  r   r  r*  uint32rV  r  s      r   test_brev_u4TestCudaIntrinsic.test_brev_u4x  sL    8856{Chhq		*sJ'Q,r   z.only get given a Python "int", assumes 32 bitsc                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g )Nzvoid(uint64[:], uint64)r   r  r  l   0  C r   l       `x)r   r  r   r  r*  rf  rV  r  s      r   test_brev_u8TestCudaIntrinsic.test_brev_u8~  sN    8856{Chhq		*s./Q!34r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g )Nvoid(int32[:], int32)r   r  r     r      r   r  r   r  r*  r  rV  r  s      r   test_clz_i4TestCudaIntrinsic.test_clz_i4  sL    8834Z@hhq)sJ'Q$r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g)	a  
Although the CUDA Math API
(http://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html)
only says int32 & int64 arguments are supported in C code, the LLVM
IR input supports i8, i16, i32 & i64 (LLVM doesn't have a concept of
unsigned integers, just unsigned operations on integers).
http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics
r  r   r  r  ra  r   rb  Nrc  r  s      r   test_clz_u4TestCudaIntrinsic.test_clz_u4  sN     8845jAhhq)sJ'Q$r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g Nr`  r   r  r  l    r   rc  r  s      r   test_clz_i4_1s TestCudaIntrinsic.test_clz_i4_1s  L    8834Z@hhq)sJ'Q#r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   SS5        g )Nr`  r   r  r  r   rj  CUDA semanticsrc  r  s      r   test_clz_i4_0s TestCudaIntrinsic.test_clz_i4_0s  sO    8834Z@hhq)sC Q%56r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g )Nvoid(int32[:], int64)r   r  r     r   /   rc  r  s      r   test_clz_i8TestCudaIntrinsic.test_clz_i8  sM    8834Z@hhq)s-.Q$r   c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        US   " US5        U R                  US   S	5        g )
Nr`  r   r  r  ra  r              rj  r   r  r   r  r*  r  rV  r  s      r   test_ffs_i4TestCudaIntrinsic.test_ffs_i4  sn    8834Z@hhq)sJ'Q$sJ'Q$r   c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        US   " US5        U R                  US   S	5        g )
Nr  r   r  r  ra  r   ry  rz  rj  r{  r  s      r   test_ffs_u4TestCudaIntrinsic.test_ffs_u4  sn    8845jAhhq)sJ'Q$sJ'Q$r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g rj  r{  r  s      r   test_ffs_i4_1s TestCudaIntrinsic.test_ffs_i4_1s  rm  r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        g )Nr`  r   r  r  r   r{  r  s      r   test_ffs_i4_0s TestCudaIntrinsic.test_ffs_i4_0s  sL    8834Z@hhq)sC Q#r   c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " US5        U R                  US   S5        US   " US5        U R                  US   S	5        g )
Nrs  r   r  r  rt  r   rk  l        !   r{  r  s      r   test_ffs_i8TestCudaIntrinsic.test_ffs_i8  so    8834Z@hhq)s-.Q$sK(Q$r   c                 r   [         R                  " S5      " [        5      nSn[        R                  " US-  [        R
                  S9n[        R                  " [        R                  " S[        R
                  S9U5      nUSUS-  4   " U5        U R                  [        R                  " X4:H  5      5        g )Nr  r)   rj  r  r   )
r   r  r   r  r*  r  tiler   r  r!  )r  r  countr   r#  s        r   test_simple_laneid$TestCudaIntrinsic.test_simple_laneid  s~    88,-m<hhurz2ggbii"((3U;EBJ$sz*+r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " U5        U R                  US   SS5        g )Nr  r   r  r  r   rj  ro  )r   r  r  r  r*  r  rV  r  s      r   test_simple_warpsize&TestCudaIntrinsic.test_simple_warpsize  sM    88,-o>hhq)sQ%56r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nS H,  nUS   " X#5        U R                  US   [        U5      5        M.     g )Nzvoid(int64[:], float32)r   r  r	  g      g      g      g      ?g      @g      @g      @r  r   r   r  r   r  r*  r   rV  r   r  r  r   r   s       r   test_round_f4TestCudaIntrinsic.test_round_f4  Y    8856|Dhhq)@ATN3"SVU1X. Ar   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nS H,  nUS   " X#5        U R                  US   [        U5      5        M.     g )Nzvoid(int64[:], float64)r   r  r  r  r   r  r  s       r   test_round_f8TestCudaIntrinsic.test_round_f8  r  r   c           	         [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  R                  S5        [        R                  R                  S5      R                  [        R
                  5      n[        R                  " U[        R                  " [        R                  [        R                  * [        R                  /5      45        Sn[        R                  " X45       HF  u  pVU R                  XVS9   US   " X%U5        U R!                  US	   [#        XV5      S
S9  S S S 5        MH     g ! , (       d  f       MZ  = f)N void(float32[:], float32, int32)r   r  {   rj  )r   r   r)   r   r   r'  r   valr   r  r   singleprec)r   r  r   r  r*  r  r  r  r  concatenater  infnan	itertoolsproductr"  assertPreciseEqualr   r  r  r   valsdigitsr  r   s          r   test_round_to_f4"TestCudaIntrinsic.test_round_to_f4  s    88>?Phhq

+
		syy#**2::6
bhh'@ABC
 &--d;LC#7s1''Ac0C-5 ( 7 87 <77s   %*E
E+	z$Overflow behavior differs on CPythonc                 ,   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  " [        R
                  5      R                  nSnUS   " X#U5        U R                  US   U5        g )Nr  r   r  i,  r  r   )	r   r  r   r  r*  r  finfomaxrV  r  r  r   r  r   s        r   test_round_to_f4_overflow+TestCudaIntrinsic.test_round_to_f4_overflow	  sm     88>?Phhq

+hhrzz"&& s)Q%r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nSnSnUS   " X#U5        U R                  US   [        X45      SS	9  g )
Nr  r   r  gQ?r   r  r   r  r  )r   r  r   r  r*  r  r  r   r  s        r   test_round_to_f4_halfway*TestCudaIntrinsic.test_round_to_f4_halfway  sb    88>?Phhq

+ s)Ac(;(Kr   c           	      N   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  R                  S5        [        R                  R                  S5      n[        R                  " U[        R                  " [        R                  [        R                  * [        R                  /5      45        Sn[        R                  " X45       HF  u  pVU R                  XVS9   US   " X%U5        U R                  US	   [!        XV5      S
S9  S S S 5        MH     SnSnU R                  XVS9   US   " X%U5        U R                  US	   [!        XV5      SS9  S S S 5        g ! , (       d  f       M  = f! , (       d  f       g = f)N void(float64[:], float64, int32)r   r  r  rj  )r  r  r  r  r  r   r   r)   r   r   r'  r  r  r   exactr  g`8p=<   double)r   r  r   r  r*  r  r  r  r  r  r  r  r  r  r"  r  r   r  s          r   test_round_to_f8"TestCudaIntrinsic.test_round_to_f8!  sE   88>?Phhq

+
		syy#
bhh'@ABC7%--d;LC#7s1''Ac0C-4 ( 6 87 < +\\c\3TN3W-##CFE#,?)1 $ 3 43 87 43s   *F*F
F	
F$c                 ,   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  " [        R
                  5      R                  nSnUS   " X#U5        U R                  US   U5        g )Nr  r   r  r  r  r   )	r   r  r   r  r*  r  r  r  rV  r  s        r   test_round_to_f8_overflow+TestCudaIntrinsic.test_round_to_f8_overflow8  sm     88>?Phhq

+hhrzz"&& s)Q%r   c                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nSnSnUS   " X#U5        U R                  US   [        X45      SS	9  g )
Nr  r   r  g\(\?r   r  r   r  r  )r   r  r   r  r*  r  r  r   r  s        r   test_round_to_f8_halfway*TestCudaIntrinsic.test_round_to_f8_halfwayE  sb    88>?Phhq

+ s)Ac(;(Kr   r   )M__name__
__module____qualname____firstlineno__r  r  r$  r7  r
   r>  rB  rL  rS  rX  rg  r{  r~  r  r  r  r  r  r  r  r   r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r)  r1  r@  rD  rH  rL  rR  rV  rZ  r]  rd  rg  rk  rp  rv  r|  r  r  r  r  r  r  r  r  r  r  r  r  r  r  __static_attributes____classcell__)r  s   @r   r  r  m  s   %,0& 345 55 349 599,0 01/ 2/0 )*C +C06L"
-%*$$66 8 8 0 0 ?@& A&
 ? ? 0 0 ?@) A)
 8 8 0 0 ?@& A&
 8 8 0 0 ?@& A&
 0 0 - - 2 2 0 0 ?@& A&
 6 6 0 0 ?@& A&
  8  8D / /$ 3 3> ( (  	1 	1 	1 	1@@- EF5 G5%%$7 EF% G%%%$$ EF% G%,7//74 ;<& =&L3. ;<
& =
&	L 	Lr   r  __main__)Yr  numpyr  r5  rq  numbar   r   
numba.cudar   numba.core.errorsr   numba.core.typesr   numba.cuda.testingr   r	   r
   r   r   r   r"   r&   r*   r/   r1   r>   rB   rI   rO   rR   rW   rY   r]   r_   rc   re   rh   rm   rr   ru   rz   r|   r   r   r   r   r   r   r  r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r  r  r  r  mainr   r   r   <module>r     s      	  " ) 3 3

/

(".%("(""&""!!!!!!    1
4
4
7
7
""$$$%&$%%&%&$&%	)
aL aLH zMMO r   