+
    :is                        ^ RI t ^ RIt^ RIt^ RIt^ RIHtHt ^ RIH	t	 ^ RI
Ht ^ RIHt ^ RIHtHtHtHt R tR tR	 tR
 tR tR tR tR tR tR tR tR tR tR t R t!R t"R t#R t$R t%R t&R t'R t(R t)R t*R t+R  t,R! t-R" t.R# t/R$ t0]Pb                  ! R%R&7      R' 4       t2]Pb                  ! R%R&7      R( 4       t3R) t4R* t5R+ t6R, t7R- t8R. t9R/ t:R0 t;R1 t<R2 t=R3 t>R4 t?R5 t@R6 tAR7 tBR8 tCR9 tDR: tER; tFR< tGR= tHR> tIR? tJR@ tKRA tLRB tMRC tNRD tORE tPRF tQRG tRRH tSRI tTRJ tU ! RK RL]4      tV]WRM8X  d   ]P                  ! 4        R# R# )N    N)cudaint64)compile_ptx)TypingError)f2)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_53c                 B    \         P                  P                  pW^ &   R# r   Nr   	threadIdxxaryis   & e/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pysimple_threadidxr          AF    c                 B    \         P                  P                  pWV&   R # Nr   r   s   & r   fill_threadidxr      r   r   c                     \         P                  P                  p\         P                  P                  p\         P                  P                  pV^,           V^,           ,          V^,           ,          WW#3&   R#    N)r   r   r   yz)r   r   jks   &   r   fill3d_threadidxr"      sP    AAAEa!e$A.C1Lr   c                 :    \         P                  ! ^4      pWV&   R# r   r   gridr   s   & r   simple_grid1dr&      s    		!AFr   c                 N    \         P                  ! ^4      w  rW,           WV3&   R#    Nr$   )r   r   r    s   &  r   simple_grid2dr*   $   s    99Q<DAC1Ir   c                 x    \         P                  ! ^4      p\         P                  ! ^4      pV^ 8X  d   W ^ &   R# R# r   r   r%   gridsize)r   r   r   s   &  r   simple_gridsize1dr.   )   s0    		!AaAAvA r   c                     \         P                  ! ^4      w  r\         P                  ! ^4      w  r4V^ 8X  d   V^ 8X  d   W0^ &   W@^&   R# R# R# r(   r,   )r   r   r    r   r   s   &    r   simple_gridsize2dr0   0   sC    99Q<DA==DAAv!q&AA vr   c                    \         P                  ! ^4      w  r\         P                  P                  \         P                  P                  ,          p\         P                  P
                  \         P                  P
                  ,          pV P                  w  rV\        WV4       F"  p\        W%V4       F  pWx,           WV3&   K  	  K$  	  R# 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                 8    \         P                  ! V4      V ^ &   R# r   )r   popcr   r6   s   &&r   simple_popcrA   C       YYq\CFr   c                 :    \         P                  ! WV4      V ^ &   R# r   )r   fmar   abr6   s   &&&&r   
simple_fmarH   G   s    XXaACFr   c                 h    \         P                  P                  V^ ,          V^ ,          4      V ^ &   R# r   r   fp16haddr   rF   rG   s   &&&r   simple_haddrN   K   "    YY^^AaD!A$'CFr   c                 J    \         P                  P                  W4      V ^ &   R# r   rJ   rM   s   &&&r   simple_hadd_scalarrQ   O       YY^^A!CFr   c                 x    \         P                  P                  V^ ,          V^ ,          V^ ,          4      V ^ &   R# r   r   rK   hfmarE   s   &&&&r   simple_hfmarV   S   s(    YY^^AaD!A$!-CFr   c                 L    \         P                  P                  WV4      V ^ &   R# r   rT   rE   s   &&&&r   simple_hfma_scalarrX   W   s    YY^^A!$CFr   c                 h    \         P                  P                  V^ ,          V^ ,          4      V ^ &   R# r   r   rK   hsubrM   s   &&&r   simple_hsubr\   [   rO   r   c                 J    \         P                  P                  W4      V ^ &   R# r   rZ   rM   s   &&&r   simple_hsub_scalarr^   _   rR   r   c                 h    \         P                  P                  V^ ,          V^ ,          4      V ^ &   R# r   r   rK   hmulrM   s   &&&r   simple_hmulrb   c   rO   r   c                 J    \         P                  P                  W4      V ^ &   R# r   r`   rM   s   &&&r   simple_hmul_scalarrd   g   rR   r   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   hdivrM   s   &&&r   simple_hdiv_scalarrg   k   rR   r   c                     \         P                  ! ^4      pW0P                  8  d4   W,          pW#,          p\         P                  P	                  WE4      W&   R# R# r   )r   r%   sizerK   rf   )r   array_aarray_br   rF   rG   s   &&&   r   simple_hdiv_kernelrl   o   s?    		!A88|JJ% r   c                 X    \         P                  P                  V^ ,          4      V ^ &   R# r   r   rK   hnegr   rF   s   &&r   simple_hnegrq   w       YY^^AaD!CFr   c                 J    \         P                  P                  V4      V ^ &   R# r   rn   rp   s   &&r   simple_hneg_scalarrt   {       YY^^ACFr   c                 X    \         P                  P                  V^ ,          4      V ^ &   R# r   r   rK   habsrp   s   &&r   simple_habsry      rr   r   c                 J    \         P                  P                  V4      V ^ &   R# r   rw   rp   s   &&r   simple_habs_scalarr{      ru   r   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   heqrM   s   &&&r   simple_heq_scalarr~          YY]]1 CFr   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   hnerM   s   &&&r   simple_hne_scalarr      r   r   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   hgerM   s   &&&r   simple_hge_scalarr      r   r   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   hgtrM   s   &&&r   simple_hgt_scalarr      r   r   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   hlerM   s   &&&r   simple_hle_scalarr      r   r   c                 J    \         P                  P                  W4      V ^ &   R# r   r   rK   hltrM   s   &&&r   simple_hlt_scalarr      r   r   T)devicec                 @    \         P                  P                  W4      # r   r   r   r   s   &&r   
hlt_func_1r          99==r   c                 @    \         P                  P                  W4      # r   r   r   s   &&r   
hlt_func_2r      r   r   c                 H    \        W4      ;'       d    \        W#4      V ^ &   R# r   )r   r   rrF   rG   r6   s   &&&&r   test_multiple_hcmp_1r      s    a00
1 0AaDr   c                 p    \        W4      ;'       d     \        P                  P                  W#4      V ^ &   R# r   )r   r   rK   r   r   s   &&&&r   test_multiple_hcmp_2r      %    a33		a 3AaDr   c                 p    \        W4      ;'       d     \        P                  P                  W24      V ^ &   R# r   )r   r   rK   r   r   s   &&&&r   test_multiple_hcmp_3r      r   r   c                     \         P                  P                  W4      ;'       d     \         P                  P                  W#4      V ^ &   R# r   r   r   s   &&&&r   test_multiple_hcmp_4r      -    99==66499==#6AaDr   c                     \         P                  P                  W4      ;'       d     \         P                  P                  W24      V ^ &   R# r   )r   rK   r   r   r   s   &&&&r   test_multiple_hcmp_5r      r   r   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   hmaxrM   s   &&&r   simple_hmax_scalarr      rR   r   c                 J    \         P                  P                  W4      V ^ &   R# r   )r   rK   hminrM   s   &&&r   simple_hmin_scalarr      rR   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   lenrK   hsinr   r   r   s   && r   simple_hsinr      5    		!A3q6zyy~~ad# r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hcosr   s   && r   simple_hcosr      r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hlogr   s   && r   simple_hlogr      r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hlog2r   s   && r   simple_hlog2r      5    		!A3q6zyyqt$ r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hlog10r   s   && r   simple_hlog10r      7    		!A3q6zyy% r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hexpr   s   && r   simple_hexpr      r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hexp2r   s   && r   simple_hexp2r      r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hsqrtr   s   && r   simple_hsqrtr      r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hrsqrtr   s   && r   simple_hrsqrtr     s7    		!A3q6zyy% r   c                     V R,          # )g      ?g      ࿩ )r   dtypes   &&r   numpy_hrsqrtr   
  s    9r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hceilr   s   && r   simple_hceilr     r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hfloorr   s   && r   simple_hfloorr     r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hrcpr   s   && r   simple_hrcpr     r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   htruncr   s   && r   simple_htruncr   #  r   r   c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hrintr   s   && r   simple_hrintr   *  r   r   c                 8    \         P                  ! V4      V ^ &   R# r   )r   cbrtrp   s   &&r   simple_cbrtr   1  rB   r   c                 8    \         P                  ! V4      V ^ &   R# r   )r   brevr@   s   &&r   simple_brevr   5  rB   r   c                 8    \         P                  ! V4      V ^ &   R# r   )r   clzr@   s   &&r   
simple_clzr   9      XXa[CFr   c                 8    \         P                  ! V4      V ^ &   R# r   )r   ffsr@   s   &&r   
simple_ffsr   =  r   r   c                 "    \        V4      V ^ &   R# r   roundr@   s   &&r   simple_roundr   A  s    1XCFr   c                 "    \        W4      V ^ &   R# r   r   )r   r6   ndigitss   &&&r   simple_round_tor   E  s    1CFr   c                     \         P                  ! ^4      pW,          ^8  d!   V^,          ^ 8X  d   W#,          W&   R# ^W&   R# ^W&   R# r   r$   )rF   rG   r6   r   s   &&& r   branching_with_ifsr   I  s:    		!Ataxq5A:4ADADr   c                     \         P                  ! ^4      p\         P                  ! V^,          ^ 8H  W#,          ^4      p\         P                  ! W,          ^8  V^4      W&   R# r   )r   r%   selp)rF   rG   r6   r   inners   &&&  r   branching_with_selpsr   U  sE    		!AIIa!eqj!$+E99QTAXua(ADr   c                 V    \         P                  ! ^4      p\         P                  W&   R# r   )r   r%   laneidr   s   & r   simple_laneidr   \  s    		!A[[CFr   c                 ,    \         P                  V ^ &   R# r   )r   warpsize)r   s   &r   simple_warpsizer   a  s    ]]CFr   c                 2    \         P                  ! V 4       R # r   r$   r   s   &r   nonliteral_gridr   e  s    IIaLr   c                 2    \         P                  ! V 4       R # r   )r   r-   r   s   &r   nonliteral_gridsizer  i  s    MM!r   c                     a a ] tR tRt oV 3R ltR tR tR t]! R4      R 4       t	]! R4      R 4       t
R	 tR
 tR t]! R4      R 4       t]! R4      R 4       tR tR tR tR tR tR tR tR t]R 4       t]R 4       t]! R4      R 4       t]R 4       t]R 4       t]! R4      R 4       t]R 4       t]R  4       t ]! R4      R! 4       t!]R" 4       t"]R# 4       t#]! R4      R$ 4       t$]R% 4       t%]R& 4       t&]R' 4       t']R( 4       t(]! R4      R) 4       t)]R* 4       t*]R+ 4       t+]! R4      R, 4       t,]R- 4       t-]R. 4       t.]R/ 4       t/]R0 4       t0]R1 4       t1]R2 4       t2R3 t3R4 t4R5 t5]! R64      R7 4       t6R8 t7R9 t8R: t9R; t:]! R64      R< 4       t;R= t<R> t=R? t>R@ t?]! R64      RA 4       t@RB tARC tBRD tCRE tDRF tE]! RG4      RH 4       tFRI tGRJ tH]! RG4      RK 4       tIRL tJRMtKVtLV ;tM# )NTestCudaIntrinsicim  c                b   < \         SV `  4        \        P                  P	                  ^ 4       R# r   )supersetUpnprandomseed)self	__class__s   &r   r  TestCudaIntrinsic.setUpn  s    
		qr   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! V4       V P                  V^ ,          ^ 8H  4       R# )void(int32[:])r   N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   \         P                  ! R 4      ! \        4      p^
p\        P                  ! V\        P
                  R7      p\        P                  ! V\        P
                  R7      pV^V3,          ! V4       V P                  \        P                  ! W48H  4      4       R# r  r  N)	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                   aaa ^^^uoooVVV3R lpVVV3R lpV! 4       pV! 4       pV P                  \        P                  ! W48H  4      4       R# )   c                     < \         P                  ! R 4      ! \        4      p \        P                  ! SSS3\        P
                  R7      pV ^SSS33,          ! V4       V# )zvoid(int32[:,:,::1])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                     < \         P                  ! R 4      ! \        4      p \        P                  ! \        P
                  ! SSS3\        P                  R7      4      pV ^SSS33,          ! V4       V# )zvoid(int32[::1,:,:])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   N)r  r  r  )r  r*  r.  c_resf_resr'  r(  r)  s   &    @@@r   test_fill3d_threadidx'TestCudaIntrinsic.test_fill3d_threadidx  sC    Q1a		 u~./r   zCudasim does not check typesc                    V P                  \        R 4      ;_uu_ 4        \        P                  ! R4      ! \        4       RRR4       R#   + '       g   i     R# ; iRequireLiteralValuezvoid(int32)N)assertRaisesRegexr   r   r  r   r  s   &r   test_nonliteral_grid_error,TestCudaIntrinsic.test_nonliteral_grid_error  s7    ##K1FGGHH]#O4 HGGG   !A

A	c                    V P                  \        R 4      ;_uu_ 4        \        P                  ! R4      ! \        4       RRR4       R#   + '       g   i     R# ; ir5  )r7  r   r   r  r  r8  s   &r   test_nonliteral_gridsize_error0TestCudaIntrinsic.test_nonliteral_gridsize_error  s8    ##K1FGGHH]#$78 HGGGr;  c                @   \         P                  ! R 4      ! \        4      p^^r2W#,          p\        P                  ! V\        P
                  R7      pWV3,          ! V4       V P                  \        P                  ! V\        P                  ! V4      8H  4      4       R# void(int32[::1])r  N)	r   r  r&   r  emptyr  r  r  r  )r  r  ntidnctaidnelemr   s   &     r   test_simple_grid1d$TestCudaIntrinsic.test_simple_grid1d  sj    88./>!fhhuBHH-s#sbii&6678r   c                (   \         P                  ! R 4      ! \        4      pRpRpV^ ,          V^ ,          ,          V^,          V^,          ,          3p\        P                  ! V\        P
                  R7      pVP                  4       pWV3,          ! V4       \        VP                  ^ ,          4       F2  p\        VP                  ^,          4       F  pWx,           WgV3&   K  	  K4  	  V P                  \        P                  ! WV8H  4      4       R# zvoid(int32[:,::1])r  Nr#           )r   r  r*   r  rB  r  copyr5   r4   r  r  )	r  r  rC  rD  r4   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                    \         P                  ! R 4      ! \        4      p^^r2\        P                  ! ^\        P
                  R7      pWV3,          ! V4       V P                  V^ ,          W2,          4       R# r@  )r   r  r.   r  r%  r  assertEqualr  r  rC  rD  r   s   &    r   test_simple_gridsize1d(TestCudaIntrinsic.test_simple_gridsize1d  sW    88./0AB!fhhq)s#Q/r   zRequires too many threadsc                L   \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P                  ! ^\        P                  R7      pVR,          ! W#4       V P                  V^ ,          ^ 4       V P                  V^ ,          ^ 4       R# )c                    \         P                  ! ^4      p\         P                  P                  \         P                  P                  ,          \         P
                  P                  ,           p\         P                  ! ^4      p\         P                  P                  \         P                  P                  ,          pW#8w  d   ^V ^ &   WE8w  d   ^V^ &   R# R# r   )r   r%   blockIdxr   r3   r   r-   r2   )
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  N)i Q   )r   r  r  r%  uint64rS  )r  r`  rZ  r[  s   &   r   test_issue_9229!TestCudaIntrinsic.test_issue_9229  s|     
	& 
	& XXaryy1
!2995 	
-4A**A.r   zTests PTX emissionc           	        \         R ,          \         \         R ,          3p\        P                  ! V4      ! \        4      p\        P                  ! V4      ! \        4      p^ p^p\
        P                  ! ^ ^\
        P                   R7      pVP                  4       p^VR&   \
        P                  ! V\
        P                   R7      pW$^3,          ! WV4       VP                  V4      p	V P                  ^\        \        P                  ! RV	4      4      4       \
        P                  P                  WRR7       \
        P                  ! V\
        P                   R7      pW4^3,          ! WV4       VP                  V4      p	V P                  ^ \        \        P                  ! RV	4      4      4       \
        P                  P                  WRR7       R# )	NNN)r4   
fill_valuer   :NrM  Nr  z	\s+bra\s+	branching)err_msgr   N)r   r   r  r   r   r  fullrO  r  inspect_asmrS  r   refindalltestingassert_array_equal)
r  sigcu_branching_with_ifscu_branching_with_selpsnrG   r6   expectedrF   ptxs
   &         r   	test_selpTestCudaIntrinsic.test_selp  sJ   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   \         P                  ! R 4      ! \        4      pRpRp\        P                  ! ^\        P
                  R7      pWV3,          ! V4       V P                  V^ ,          V^ ,          V^ ,          ,          4       V P                  V^,          V^,          V^,          ,          4       R# )rA  r  NrJ  rL  )r   r  r0   r  r%  r  rS  rT  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           	     t   \         P                  ! R 4      ! \        4      pRpRpV^ ,          V^ ,          ,          V^,          V^,          ,          3p\        P                  ! V\        P
                  R7      pWV3,          ! V4       Vw  rgVP                  w  r\        \        V^ ,          4      \        V^,          4      4       Fm  w  rWj,           W{,           r\        WV4       FI  p\        WV4       F6  pV P                  W_V3,          W,           8H  W_V3,          W,           34       K8  	  KK  	  Ko  	  R# rI  )
r   r  r=   r  rB  r  r4   zipr5   r  )r  r  rC  rD  r4   r   r9   r:   r;   r<   r   r    r7   r8   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                    \         P                  R  4       p\        P                  ! R\        P                  R7      P                  ^	^	^	4      pVR,          ! V4       \        P                  P                  VR4       R# )c                     \         P                  ! ^4      w  rp\         P                  ! ^4      w  rEpWE,          V,          WW#3&   R# rK  Nr,   )outr   r   r   rF   rG   r6   s   &      r   foo*TestCudaIntrinsic.test_3dgrid.<locals>.foo	  s6    iilGA!mmA&GA!519C1Lr   r  Ni  )rK  rK  rK  r  )r   r  r  r%  r  reshapero  assert_equal)r  r  arrs   &  r   test_3dgridTestCudaIntrinsic.test_3dgrid  s_    		% 
	%
 hhvRXX.66q!Q? !#&


V,r   c                *   \         P                  R  4       p^^^rCp\        P                  ! W#,          V,          \        P                  R7      P                  W#V4      pVR,          ! V4       V P                  \        P                  ! V4      4       R# )c                 T   \         P                  ! ^4      w  rp\         P                  ! ^4      w  rEpV\         P                  P                  \         P
                  P                  \         P                  P                  ,          ,           8H  ;'       d    V\         P                  P                  \         P
                  P                  \         P                  P                  ,          ,           8H  ;'       d\    V\         P                  P                  \         P
                  P                  \         P                  P                  ,          ,           8H  pV\         P                  P                  \         P                  P                  ,          8H  ;'       d    V\         P                  P                  \         P                  P                  ,          8H  ;'       d=    V\         P                  P                  \         P                  P                  ,          8H  pT;'       d    TWW#3&   R# r  )
r   r%   r-   r   r   rY  r3   r   r   r2   )	r  r   r   r   rF   rG   r6   grid_is_rightgridsize_is_rights	   &        r   r  ,TestCudaIntrinsic.test_3dgrid_2.<locals>.foo  sR   iilGA!mmA&GA!T^^%%$--//(III J JT^^%%$--//(IIIJ JT^^%%$--//(III 
 "#dmmoo&F!F "G "G!"dmmoo&F!F"G "G!"dmmoo&F!F  )>>->C1Lr   r  N))r#  rK  r)   )rK  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  sn    		? 
	? %ahh	"((3;;A!D !#&s$r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! V^4       V P                  V^ ,          ^4       R# )void(int32[:], uint32)r  Nr  r   r  rA   r  r%  r  rS  r  s   &  r   test_popc_u4TestCudaIntrinsic.test_popc_u4)  sL    8845kBhhq)sD!Q#r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^4       R# )zvoid(int32[:], uint64)r  l        @ Nr  r  r  s   &  r   test_popc_u8TestCudaIntrinsic.test_popc_u8/  sL    8845kBhhq)sN+Q#r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VRRR4       \        P                  P                  V^ ,          ^
4       R# )zvoid(f4[:], f4, f4, f4)r         @      @      @Nr  )r   r  rH   r  r%  float32ro  assert_allcloser  s   &  r   test_fma_f4TestCudaIntrinsic.test_fma_f45  T    8856zBhhq

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


""3q695r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VRRR4       \        P                  P                  V^ ,          ^
4       R# )zvoid(f8[:], f8, f8, f8)r  r  r  r  Nr  )r   r  rH   r  r%  float64ro  r  r  s   &  r   test_fma_f8TestCudaIntrinsic.test_fma_f8;  r  r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! R.\        P
                  R7      p\        P                  ! R.\        P
                  R7      pVR,          ! W#V4       \        P                  P                  V^ ,          W4,           4       R# void(f2[:], f2[:], f2[:])r  r  r  Nr  )	r   r  rN   r  r%  float16arrayro  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   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      p\        P
                  ! R4      pVR,          ! W#V4       W4,           p\        P                  P                  V^ ,          V4       R# )void(f2[:], f2, f2)r  JM!	@r  Nr  )r   r  rQ   r  r%  r  ro  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 ,          \         \         3p\        \        VRR7      w  r#V P                  RV4       R# )rg  cczadd.f16NrM  rK  )r   r   rQ   assertInr  argsrv  _s   &   r   test_hadd_ptxTestCudaIntrinsic.test_hadd_ptxT  0    1r2/&Ai%r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! R.\        P
                  R7      p\        P                  ! R.\        P
                  R7      p\        P                  ! R.\        P
                  R7      pVR,          ! W#WE4       \        P                  P                  V^ ,          W4,          V,           4       R# )z void(f2[:], f2[:], f2[:], f2[:])r  r  r  r  Nr  )	r   r  rV   r  r%  r  r  ro  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                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      p\        P
                  ! R4      p\        P
                  ! R4      pVR,          ! W#WE4       W4,          V,           p\        P                  P                  V^ ,          V4       R# )zvoid(f2[:], f2, f2, f2)r  r  r  r  Nr  )r   r  rX   r  r%  r  ro  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                    \         R ,          \         \         \         3p\        \        VRR7      w  r#V P                  RV4       R# )rg  r  z
fma.rn.f16Nr  )r   r   rX   r  r  s   &   r   test_hfma_ptxTestCudaIntrinsic.test_hfma_ptxo  s3    1r2r"/&AlC(r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! R.\        P
                  R7      p\        P                  ! R.\        P
                  R7      pVR,          ! W#V4       \        P                  P                  V^ ,          W4,
          4       R# r  )	r   r  r\   r  r%  r  r  ro  r  r  s   &    r   	test_hsubTestCudaIntrinsic.test_hsubu  r  r   c                h   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      p\        P
                  ! R4      pVR,          ! W#V4       W4,
          p\        P                  P                  V^ ,          V4       R# r  r  r  gQ?Nr  )r   r  r^   r  r%  r  ro  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 ,          \         \         3p\        \        VRR7      w  r#V P                  RV4       R# )rg  r  zsub.f16Nr  )r   r   r^   r  r  s   &   r   test_hsub_ptxTestCudaIntrinsic.test_hsub_ptx  r  r   c                   \         P                  ! 4       ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! R.\        P
                  R7      p\        P                  ! R.\        P
                  R7      pVR,          ! W#V4       \        P                  P                  V^ ,          W4,          4       R# )r   r  r  r  Nr  )	r   r  rb   r  r%  r  r  ro  r  r  s   &    r   	test_hmulTestCudaIntrinsic.test_hmul  sz    88:k*hhq

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


""3q64;7r   c                h   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      p\        P
                  ! R4      pVR,          ! W#V4       W4,          p\        P                  P                  V^ ,          V4       R# r  )r   r  rd   r  r%  r  ro  r  r  s   &     r   test_hmul_scalar"TestCudaIntrinsic.test_hmul_scalar  r  r   c                    \         R ,          \         \         3p\        \        VRR7      w  r#V P                  RV4       R# )rg  r  zmul.f16Nr  )r   r   rd   r  r  s   &   r   test_hmul_ptxTestCudaIntrinsic.test_hmul_ptx  r  r   c                h   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      p\        P
                  ! R4      pVR,          ! W#V4       W4,          p\        P                  P                  V^ ,          V4       R# r  )r   r  rg   r  r%  r  ro  r  r  s   &     r   test_hdiv_scalar"TestCudaIntrinsic.test_hdiv_scalar  st    88123EFhhq

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


""3q63/r   c                    \         P                  ! R 4      ! \        4      p\        P                  P                  RRRR7      P                  \        P                  4      p\        P                  P                  RRRR7      P                  \        P                  4      p\        P                  ! V\        P                  R7      pVP                  VP                  4      ! WBV4       W#,          p\        P                  P                  WE4       R# )r    i  ri   r  Ni  )r   r  rl   r  r	  randintastyper  
zeros_likeforallri   ro  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                N   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! R.\        P
                  R7      pVR,          ! W#4       \        P                  P                  V^ ,          V) 4       R# )void(f2[:], f2[:])r  r  Nr  )	r   r  rq   r  r%  r  r  ro  r  r  r  r   r  s   &   r   	test_hnegTestCudaIntrinsic.test_hneg  sf    8801+>hhq

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


""3q6D51r   c                0   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      pVR,          ! W#4       V) p\        P                  P                  V^ ,          V4       R# )void(f2[:], f2)r  r  Nr  )r   r  rt   r  r%  r  ro  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                x    \         R ,          \         3p\        \        VRR7      w  r#V P                  RV4       R# )rg  r  zneg.f16Nr  )r   r   rt   r  r  s   &   r   test_hneg_ptxTestCudaIntrinsic.test_hneg_ptx  .    1r{/&Ai%r   c                \   \         P                  ! 4       ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! R.\        P
                  R7      pVR,          ! W#4       \        P                  P                  V^ ,          \        V4      4       R# )r   r  N      r  )
r   r  ry   r  r%  r  r  ro  r  absr  s   &   r   	test_habsTestCudaIntrinsic.test_habs  sd    88:k*hhq

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


""3q63t95r   c                @   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      pVR,          ! W#4       \        V4      p\        P                  P                  V^ ,          V4       R# )r  r  NgJM!	r  )	r   r  r{   r  r%  r  r  ro  r  r  s   &    r   test_habs_scalar"TestCudaIntrinsic.test_habs_scalar  sf    88-./ABhhq

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


""3q63/r   c                x    \         R ,          \         3p\        \        VRR7      w  r#V P                  RV4       R# )rg  r  zabs.f16Nr  )r   r   r{   r  r  s   &   r   test_habs_ptxTestCudaIntrinsic.test_habs_ptx  r  r   c                Z   \         \        \        \        \        \
        \        \        \        \        \        \        3p\        \        3p\        P                  \        P                   \        P"                  \        P$                  \        P&                  \        P(                  \        P*                  \        P,                  \        P.                  \        P0                  \        P2                  \4        3p\        P6                  \        P8                  3p^ p\        P:                  P=                  ^4       \        P:                  P?                  ^RVR7      PA                  \        PB                  4      p\        PD                  ! V4      p\G        W4       F  w  rV PI                  V	R7      ;_uu_ 4        \J        PL                  ! R4      ! V4      pV^V3,          ! Wv4       V	! V\        PB                  R7      p
\        PN                  PQ                  Wz4       RRR4       K  	  \        P:                  P?                  ^^
VR7      PA                  \        PB                  4      p\G        W$4       F  w  rV PI                  V	R7      ;_uu_ 4        \J        PL                  ! R4      ! V4      pV^V3,          ! W{4       V	! V\        PB                  R7      p
\        PN                  PQ                  Wz4       RRR4       K  	  R#   + '       g   i     EK~  ; i  + '       g   i     K  ; i)    r  r  fnr  r  N))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  ro  r  )r  kernelsexp_kernelsexpected_functionsexpected_exp_functionsr  r   r   kernelr  ru  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   7A%LA%LLL*c                   \         P                  ! 4       R  4       p^ p\        P                  P	                  ^4       \        P                  P                  V4      P                  \        P                  4      p\        P                  ! V4      pV^V3,          ! WC4       \        P                  P                  V^
V,          4       R# )c                     \         P                  ! ^4      pV\        V 4      8  d*   \         P                  P	                  W,          4      W&   R# R# r   )r   r%   r   rK   hexp10r   s   && r   hexp10_vectors5TestCudaIntrinsic.test_hexp10.<locals>.hexp10_vectors  s7    		!A3q6zyy''- r   N)r   r  r  r	  r
  randr  r  r  ro  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                   \         \        \        \        \        \
        3p\        P                  \        P                  \        P                  \        P                  \        P                  \        P                  3p\        W4       EFc  w  r4V P                  VR 7      ;_uu_ 4        \        P                   ! R4      ! V4      p\"        P$                  ! ^\"        P&                  R7      p\"        P$                  ! ^\"        P&                  R7      p\"        P(                  ! ^4      p\"        P(                  ! ^4      p	\"        P(                  ! ^4      p
VR,          ! WyV	4       V! W4      pV P+                  Wg^ ,          4       VR,          ! WyV
4       V! W4      pV P+                  Wg^ ,          4       VR,          ! WyV4       V! W4      pV P+                  Wg^ ,          4       RRR4       EKf  	  R#   + '       g   i     EK|  ; i))opzvoid(b1[:], f2, f2)r  Nr  )r~   r   r   r   r   r   operatoreqnegegtleltr}  r  r   r  r  r%  r  r  rS  )r  fnsopsr  r-  r   ru  gotr  r  arg4s   &          r   test_fp16_comparison&TestCudaIntrinsic.test_fp16_comparison!  sb    "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;G22Hc           	        \         \        \        \        \        3pV F  pV P                  VR 7      ;_uu_ 4        \        P                  ! R4      ! V4      p\        P                  ! ^\        P                  R7      p\        P                  ! R4      p\        P                  ! R4      p\        P                  ! R4      pVR,          ! WEWg4       V P                  V^ ,          4       RRR4       K  	  R#   + '       g   i     K  ; i)r  zvoid(b1[:], f2, f2, f2)r  r  r  r  Nr  )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-C66Dc                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      p\        P
                  ! R4      pVR,          ! W#V4       \        P                  P                  V^ ,          V4       \        P
                  ! R4      pVR,          ! W#V4       \        P                  P                  V^ ,          V4       R# r  r  r  r  g      @Nr  )r   r  r   r  r%  r  ro  r  r  s   &    r   	test_hmaxTestCudaIntrinsic.test_hmaxR      88123EFhhq

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


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


""3q640r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P
                  ! R4      p\        P
                  ! R4      pVR,          ! W#V4       \        P                  P                  V^ ,          V4       \        P
                  ! R4      pVR,          ! W#V4       \        P                  P                  V^ ,          V4       R# r@  )r   r  r   r  r%  r  ro  r  r  s   &    r   	test_hminTestCudaIntrinsic.test_hmin^  rC  r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pRpVR,          ! W#4       \        P                  P                  V^ ,          VR,          4       R# )zvoid(float32[:], float32)r  r  Nr  UUUUUU?)r   r  r   r  r%  r  ro  r  r  r  r   cbrt_args   &   r   test_cbrt_f32TestCudaIntrinsic.test_cbrt_f32j  X    8878Ehhq

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


""3q68+>?r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pRpVR,          ! W#4       \        P                  P                  V^ ,          VR,          4       R# )zvoid(float64[:], float64)r  g      @Nr  rH  )r   r  r   r  r%  r  ro  r  rI  s   &   r   test_cbrt_f64TestCudaIntrinsic.test_cbrt_f64q  rM  r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          R4       R# )zvoid(uint32[:], uint32)r  i0  i  Nr  )r   r  r   r  r%  uint32rS  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                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          R4       R# )zvoid(uint64[:], uint64)r  l   0  C l       `xNr  )r   r  r   r  r%  rc  rS  r  s   &  r   test_brev_u8TestCudaIntrinsic.test_brev_u8~  sN    8856{Chhq		*s./Q!34r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^4       R# )void(int32[:], int32)r     Nr  r   r  r   r  r%  r  rS  r  s   &  r   test_clz_i4TestCudaIntrinsic.test_clz_i4  sL    8834Z@hhq)sJ'Q$r   c                    \         P                  ! R4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^4       R# )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  rZ  Nr  r[  r  s   &  r   test_clz_u4TestCudaIntrinsic.test_clz_u4  sN     8845jAhhq)sJ'Q$r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^ 4       R# rY  r  l    Nr  r[  r  s   &  r   test_clz_i4_1s TestCudaIntrinsic.test_clz_i4_1s  L    8834Z@hhq)sJ'Q#r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! V^ 4       V P                  V^ ,          ^ R4       R# )rY  r  CUDA semanticsNr  r[  r  s   &  r   test_clz_i4_0s TestCudaIntrinsic.test_clz_i4_0s  sO    8834Z@hhq)sC Q%56r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^/4       R# )void(int32[:], int64)r     Nr  r[  r  s   &  r   test_clz_i8TestCudaIntrinsic.test_clz_i8  sM    8834Z@hhq)s-.Q$r   c                6   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^4       VR,          ! VR4       V P                  V^ ,          ^ 4       R# )rY  r  rZ          Nr  r   r  r   r  r%  r  rS  r  s   &  r   test_ffs_i4TestCudaIntrinsic.test_ffs_i4  sn    8834Z@hhq)sJ'Q$sJ'Q$r   c                6   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^4       VR,          ! VR4       V P                  V^ ,          ^ 4       R# )r  r  rZ  rp  Nr  rq  r  s   &  r   test_ffs_u4TestCudaIntrinsic.test_ffs_u4  sn    8845jAhhq)sJ'Q$sJ'Q$r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^4       R# rb  rq  r  s   &  r   test_ffs_i4_1s TestCudaIntrinsic.test_ffs_i4_1s  re  r   c                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! V^ 4       V P                  V^ ,          ^ 4       R# )rY  r  Nr  rq  r  s   &  r   test_ffs_i4_0s TestCudaIntrinsic.test_ffs_i4_0s  sL    8834Z@hhq)sC Q#r   c                6   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VR4       V P                  V^ ,          ^4       VR,          ! VR4       V P                  V^ ,          ^!4       R# )rk  r  rl  l        Nr  rq  r  s   &  r   test_ffs_i8TestCudaIntrinsic.test_ffs_i8  so    8834Z@hhq)s-.Q$sK(Q$r   c                   \         P                  ! R 4      ! \        4      p^p\        P                  ! V^ ,          \        P
                  R7      p\        P                  ! \        P                  ! ^ \        P
                  R7      V4      pV^V^ ,          3,          ! V4       V P                  \        P                  ! W48H  4      4       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                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! V4       V P                  V^ ,          ^ R4       R# )r  r  rg  Nr  )r   r  r   r  r%  r  rS  r  s   &  r   test_simple_warpsize&TestCudaIntrinsic.test_simple_warpsize  sM    88,-o>hhq)sQ%56r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pR F4  pVR,          ! W#4       V P                  V^ ,          \        V4      4       K6  	  R# )zvoid(int64[:], float32)r  Nr  g      g      g      g      ?g      @g      @g      @r  r   r  r   r  r%  r   rS  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                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pR F4  pVR,          ! W#4       V P                  V^ ,          \        V4      4       K6  	  R# )zvoid(int64[:], float64)r  Nr  r  r  r  s   &   r   test_round_f8TestCudaIntrinsic.test_round_f8  r  r   c           
        \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  P                  ^{4       \        P                  P                  ^ 4      P                  \        P
                  4      p\        P                  ! V\        P                  ! \        P                  \        P                  ) \        P                  .4      34       Rp\        P                  ! W44       F\  w  rVV P                  WVR7      ;_uu_ 4        VR,          ! W%V4       V P!                  V^ ,          \#        WV4      RR7       RRR4       K^  	  R#   + '       g   i     Ks  ; i) void(float32[:], float32, int32)r  valr   singleprecN)r   r   r)   rK  r#  rM     r  )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7s1''Ac0C-5 ( 7 87 <777s   35E66Fz$Overflow behavior differs on CPythonc                D   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! \        P
                  4      P                  pRpVR,          ! W#V4       V P                  V^ ,          V4       R# )r  r  i,  Nr  )	r   r  r   r  r%  r  finfomaxrS  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                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pRp^pVR,          ! W#V4       V P                  V^ ,          \        W44      RR7       R# )r  r  gQ?r  r  Nr  )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           
        \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  P                  ^{4       \        P                  P                  ^ 4      p\        P                  ! V\        P                  ! \        P                  \        P                  ) \        P                  .4      34       Rp\        P                  ! W44       F\  w  rVV P                  WVR7      ;_uu_ 4        VR,          ! W%V4       V P                  V^ ,          \!        WV4      RR7       RRR4       K^  	  R	p^pV P                  WVR7      ;_uu_ 4        VR,          ! W%V4       V P                  V^ ,          \!        WV4      RR7       RRR4       R#   + '       g   i     K  ; i  + '       g   i     R# ; i)
 void(float64[:], float64, int32)r  r  exactr  Ndouble)r  r  r  r  r  r   r   r)   rK  r#  rM  r  g`8p=<)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!  sQ   88>?Phhq

+
		syy#
bhh'@ABC7%--d;LC#77s1''Ac0C-4 ( 6 87 < +\\c\33TN3W-##CFE#,?)1 $ 3 43 877 433s   5F455G4GG	c                D   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      p\        P                  ! \        P
                  4      P                  p^pVR,          ! W#V4       V P                  V^ ,          V4       R# )r  r  Nr  )	r   r  r   r  r%  r  r  r  rS  r  s   &    r   test_round_to_f8_overflow+TestCudaIntrinsic.test_round_to_f8_overflow8  sm     88>?Phhq

+hhrzz"&& s)Q%r   c                   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pRp^pVR,          ! W#V4       V P                  V^ ,          \        W44      RR7       R# )r  r  g\(\?r  r  Nr  )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   )N__name__
__module____qualname____firstlineno__r  r  r   r2  r
   r9  r=  rF  rP  rU  rd  rw  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*  r9  r=  rA  rE  rK  rO  rS  rV  r\  r_  rc  rh  rm  rr  ru  rx  r{  r~  r  r  r  r  r  r  r  r  r  r  __static_attributes____classdictcell____classcell__)r  __classdict__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  r.  rm  numbar   r   
numba.cudar   numba.core.errorsr   numba.core.typesr   numba.cuda.testingr   r	   r
   r   r   r   r"   r&   r*   r.   r0   r=   rA   rH   rN   rQ   rV   rX   r\   r^   rb   rd   rg   rl   rq   rt   ry   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   