
    xKgs                        d dl Z d dlZd dlZd dlZd dlmZmZ d dlm	Z	 d dl
mZ d dlmZ d dlmZmZmZmZ d Zd Zd	 Zd
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z d Z!d Z"d Z#d Z$d Z%d Z&d Z'd Z(d Z)d Z*d Z+d  Z,d! Z-d" Z.d# Z/d$ Z0 ejb                  d%&      d'        Z2 ejb                  d%&      d(        Z3d) Z4d* Z5d+ Z6d, Z7d- Z8d. Z9d/ Z:d0 Z;d1 Z<d2 Z=d3 Z>d4 Z?d5 Z@d6 ZAd7 ZBd8 ZCd9 ZDd: ZEd; ZFd< ZGd= ZHd> ZId? ZJd@ ZKdA ZLdB ZMdC ZNdD ZOdE ZPdF ZQdG ZRdH ZSdI ZTdJ ZU G dK dLe      ZVeWdMk(  r ej                          yy)N    N)cudaint64)compile_ptx)TypingError)f2)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_53c                 B    t         j                  j                  }|| d<   y Nr   r   	threadIdxxaryis     k/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pysimple_threadidxr          ACF    c                 B    t         j                  j                  }|| |<   y Nr   r   s     r   fill_threadidxr      r   r   c                     t         j                  j                  }t         j                  j                  }t         j                  j                  }|dz   |dz   z  |dz   z  | |||f<   y N   )r   r   r   yz)r   r   jks       r   fill3d_threadidxr"      sT    AAAEa!e$A.C1aLr   c                 8    t        j                  d      }|| |<   y r   r   gridr   s     r   simple_grid1dr&      s    		!ACFr   c                 H    t        j                  d      \  }}||z   | ||f<   y N   r$   )r   r   r    s      r   simple_grid2dr*   $   s$    99Q<DAqAC1Ir   c                 n    t        j                  d      }t        j                  d      }|dk(  r|| d<   y y Nr   r   r   r%   gridsize)r   r   r   s      r   simple_gridsize1dr/   )   s2    		!AaAAvA r   c                     t        j                  d      \  }}t        j                  d      \  }}|dk(  r|dk(  r|| d<   || d<   y y y )Nr)   r   r   r-   )r   r   r    r   r   s        r   simple_gridsize2dr1   0   sK    99Q<DAq==DAqAv!q&AA vr   c                    t        j                  d      \  }}t         j                  j                  t         j                  j                  z  }t         j                  j
                  t         j                  j
                  z  }| j                  \  }}t        |||      D ]  }t        |||      D ]  }||z   | ||f<      y r(   )r   r%   gridDimr   blockDimr   shaperange)	cstartXstartYgridXgridYheightwidthr   r   s	            r   intrinsic_forloop_stepr>   8   s    YYq\NFFLLNNT]]__,ELLNNT]]__,EGGMFE65%(vvu-A!eAadG . )r   c                 4    t        j                  |      | d<   y r   )r   popcr   r7   s     r   simple_popcrB   C       YYq\CFr   c                 8    t        j                  |||      | d<   y r   )r   fmar   abr7   s       r   
simple_fmarI   G   s    XXaACFr   c                 V    t         j                  j                  |d   |d         | d<   y r   r   fp16haddr   rG   rH   s      r   simple_haddrO   K   "    YY^^AaD!A$'CFr   c                 J    t         j                  j                  ||      | d<   y r   rK   rN   s      r   simple_hadd_scalarrR   O       YY^^Aq!CFr   c                 ^    t         j                  j                  |d   |d   |d         | d<   y r   r   rL   hfmarF   s       r   simple_hfmarW   S   s(    YY^^AaD!A$!-CFr   c                 L    t         j                  j                  |||      | d<   y r   rU   rF   s       r   simple_hfma_scalarrY   W   s    YY^^Aq!$CFr   c                 V    t         j                  j                  |d   |d         | d<   y r   r   rL   hsubrN   s      r   simple_hsubr]   [   rP   r   c                 J    t         j                  j                  ||      | d<   y r   r[   rN   s      r   simple_hsub_scalarr_   _   rS   r   c                 V    t         j                  j                  |d   |d         | d<   y r   r   rL   hmulrN   s      r   simple_hmulrc   c   rP   r   c                 J    t         j                  j                  ||      | d<   y r   ra   rN   s      r   simple_hmul_scalarre   g   rS   r   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   hdivrN   s      r   simple_hdiv_scalarrh   k   rS   r   c                     t        j                  d      }|| j                  k  r.||   }||   }t         j                  j	                  ||      | |<   y y r   )r   r%   sizerL   rg   )r   array_aarray_br   rG   rH   s         r   simple_hdiv_kernelrm   o   sI    		!A388|AJAJ1%A r   c                 N    t         j                  j                  |d         | d<   y r   r   rL   hnegr   rG   s     r   simple_hnegrr   w       YY^^AaD!CFr   c                 H    t         j                  j                  |      | d<   y r   ro   rq   s     r   simple_hneg_scalarru   {       YY^^ACFr   c                 N    t         j                  j                  |d         | d<   y r   r   rL   habsrq   s     r   simple_habsrz      rs   r   c                 H    t         j                  j                  |      | d<   y r   rx   rq   s     r   simple_habs_scalarr|      rv   r   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   heqrN   s      r   simple_heq_scalarr          YY]]1a CFr   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   hnerN   s      r   simple_hne_scalarr      r   r   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   hgerN   s      r   simple_hge_scalarr      r   r   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   hgtrN   s      r   simple_hgt_scalarr      r   r   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   hlerN   s      r   simple_hle_scalarr      r   r   c                 J    t         j                  j                  ||      | d<   y r   r   rL   hltrN   s      r   simple_hlt_scalarr      r   r   T)devicec                 B    t         j                  j                  | |      S r   r   r   r   s     r   
hlt_func_1r          99==Ar   c                 B    t         j                  j                  | |      S r   r   r   s     r   
hlt_func_2r      r   r   c                 >    t        ||      xr t        ||      | d<   y r   )r   r   rrG   rH   r7   s       r   test_multiple_hcmp_1r      s    a0
1a 0AaDr   c                 f    t        ||      xr  t        j                  j                  ||      | d<   y r   )r   r   rL   r   r   s       r   test_multiple_hcmp_2r      &    a3		a 3AaDr   c                 f    t        ||      xr  t        j                  j                  ||      | d<   y r   )r   r   rL   r   r   s       r   test_multiple_hcmp_3r      r   r   c                     t         j                  j                  ||      xr  t         j                  j                  ||      | d<   y r   r   r   s       r   test_multiple_hcmp_4r      .    99==A6499==A#6AaDr   c                     t         j                  j                  ||      xr  t         j                  j                  ||      | d<   y r   )r   rL   r   r   r   s       r   test_multiple_hcmp_5r      r   r   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   hmaxrN   s      r   simple_hmax_scalarr      rS   r   c                 J    t         j                  j                  ||      | d<   y r   )r   rL   hminrN   s      r   simple_hmin_scalarr      rS   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   lenrL   hsinr   r   r   s      r   simple_hsinr      9    		!A3q6zyy~~ad#! r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hcosr   s      r   simple_hcosr      r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hlogr   s      r   simple_hlogr      r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hlog2r   s      r   simple_hlog2r      9    		!A3q6zyyqt$! r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hlog10r   s      r   simple_hlog10r      ;    		!A3q6zyy!%! r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hexpr   s      r   simple_hexpr      r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hexp2r   s      r   simple_hexp2r      r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hsqrtr   s      r   simple_hsqrtr      r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hrsqrtr   s      r   simple_hrsqrtr     s;    		!A3q6zyy!%! r   c                     | dz  S )Ng      ࿩ )r   dtypes     r   numpy_hrsqrtr   
  s    9r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hceilr   s      r   simple_hceilr     r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hfloorr   s      r   simple_hfloorr     r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hrcpr   s      r   simple_hrcpr     r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   htruncr   s      r   simple_htruncr   #  r   r   c                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hrintr   s      r   simple_hrintr   *  r   r   c                 4    t        j                  |      | d<   y r   )r   cbrtrq   s     r   simple_cbrtr   1  rC   r   c                 4    t        j                  |      | d<   y r   )r   brevrA   s     r   simple_brevr   5  rC   r   c                 4    t        j                  |      | d<   y r   )r   clzrA   s     r   
simple_clzr   9      XXa[CFr   c                 4    t        j                  |      | d<   y r   )r   ffsrA   s     r   
simple_ffsr   =  r   r   c                      t        |      | d<   y r   roundrA   s     r   simple_roundr   A  s    1XCFr   c                 "    t        ||      | d<   y r   r   )r   r7   ndigitss      r   simple_round_tor   E  s    1gCFr   c                 v    t        j                  d      }| |   dkD  r|dz  dk(  r	||   | |<   y d| |<   y d| |<   y )Nr      r)   r         r$   )rG   rH   r7   r   s       r   branching_with_ifsr   I  sD    		!Ataxq5A:Q4AaDAaD!r   c                     t        j                  d      }t        j                  |dz  dk(  ||   d      }t        j                  | |   dkD  |d      | |<   y )Nr   r)   r   r   r   r   )r   r%   selp)rG   rH   r7   r   inners        r   branching_with_selpsr   U  sK    		!AIIa!eqj!A$+E99QqTAXua(AaDr   c                 T    t        j                  d      }t         j                  | |<   y r   )r   r%   laneidr   s     r   simple_laneidr   \  s    		!A[[CFr   c                 *    t         j                  | d<   y r   )r   warpsize)r   s    r   simple_warpsizer  a  s    ]]CFr   c                 .    t        j                  |        y r   r$   r   s    r   nonliteral_gridr  e  s    IIaLr   c                 .    t        j                  |        y r   )r   r.   r  s    r   nonliteral_gridsizer  i  s    MM!r   c                       e Zd Z fdZd Zd Zd Z ed      d        Z ed      d        Z	d Z
d	 Zd
 Z ed      d        Z ed      d        Zd Zd Zd Zd Zd Zd Zd Zd Zed        Zed        Z ed      d        Zed        Zed        Z ed      d        Zed        Zed        Z ed      d         Z ed!        Z!ed"        Z" ed      d#        Z#ed$        Z$ed%        Z%ed&        Z&ed'        Z' ed      d(        Z(ed)        Z)ed*        Z* ed      d+        Z+ed,        Z,ed-        Z-ed.        Z.ed/        Z/ed0        Z0ed1        Z1d2 Z2d3 Z3d4 Z4 ed5      d6        Z5d7 Z6d8 Z7d9 Z8d: Z9 ed5      d;        Z:d< Z;d= Z<d> Z=d? Z> ed5      d@        Z?dA Z@dB ZAdC ZBdD ZCdE ZD edF      dG        ZEdH ZFdI ZG edF      dJ        ZHdK ZI xZJS )LTestCudaIntrinsicc                 `    t         |           t        j                  j	                  d       y r   )supersetUpnprandomseed)self	__class__s    r   r  zTestCudaIntrinsic.setUpn  s    
		qr   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |       | j                  |d   dk(         y )Nvoid(int32[:])r   r   r   r   r   )r   jitr   r  onesint32
assertTruer  compiledr   s      r   test_simple_threadidxz'TestCudaIntrinsic.test_simple_threadidxr  sO    -488,-.>?ggarxx(sA!$r   c                 B    t        j                  d      t              }d}t        j                  |t        j
                        }t        j                  |t        j
                        } |d|f   |       | j                  t        j                  ||k(               y )Nr  
   r  r   )	r   r  r   r  r  r  aranger  all)r  r  Nr   exps        r   test_fill_threadidxz%TestCudaIntrinsic.test_fill_threadidxx  sp    -488,-n=ggarxx(ii*Asscz*+r   c                     d\  fd}fd} |       } |       }| j                  t        j                  ||k(               y )N)r         c                       t        j                  d      t              } t        j                  ft        j
                        } | dff   |       |S )Nzvoid(int32[:,:,::1])r  r   )r   r  r"   r  zerosr  r  r   XYZs     r   c_contigousz<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.c_contigous  sR    7txx 678HIH((Aq!9BHH5C"HQAq	\"3'Jr   c                       t        j                  d      t              } t        j                  t        j
                  ft        j                              } | dff   |       |S )Nzvoid(int32[::1,:,:])r  r   )r   r  r"   r  asfortranarrayr'  r  r(  s     r   f_contigousz<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.f_contigous  s]    7txx 678HIH##BHHaAYbhh$GHC"HQAq	\"3'Jr   )r  r  r  )r  r,  r/  c_resf_resr)  r*  r+  s        @@@r   test_fill3d_threadidxz'TestCudaIntrinsic.test_fill3d_threadidx  sA    1a		 u~./r   zCudasim does not check typesc                     | j                  t        d      5   t        j                  d      t               d d d        y # 1 sw Y   y xY wNRequireLiteralValuezvoid(int32))assertRaisesRegexr   r   r  r  r  s    r   test_nonliteral_grid_errorz,TestCudaIntrinsic.test_nonliteral_grid_error  s3    ##K1FG#DHH]#O4 HGG    A  A	c                     | j                  t        d      5   t        j                  d      t               d d d        y # 1 sw Y   y xY wr4  )r6  r   r   r  r  r7  s    r   test_nonliteral_gridsize_errorz0TestCudaIntrinsic.test_nonliteral_gridsize_error  s4    ##K1FG#DHH]#$78 HGGr9  c                 .    t        j                  d      t              }d\  }}||z  }t        j                  |t        j
                        } |||f   |       | j                  t        j                  |t        j                  |      k(               y )Nvoid(int32[::1])r      r  )	r   r  r&   r  emptyr  r  r  r  )r  r  ntidnctaidnelemr   s         r   test_simple_grid1dz$TestCudaIntrinsic.test_simple_grid1d  st    /488./>fvhhuBHH-s#sbii&6678r   c                     t        j                  d      t              }d}d}|d   |d   z  |d   |d   z  f}t        j                  |t        j
                        }|j                         } |||f   |       t        |j                  d         D ])  }t        |j                  d         D ]  }||z   |||f<    + | j                  t        j                  ||k(               y Nzvoid(int32[:,::1])r   r   r$  r%  r   r   r  )r   r  r*   r  r@  r  copyr6   r5   r  r  )	r  r  rA  rB  r5   r   r!  r   r    s	            r   test_simple_grid2dz$TestCudaIntrinsic.test_simple_grid2d  s    148801-@a6!9$d1gq	&9:hhuBHH-hhjs#syy|$A399Q<(EAqD	 ) % 	scz*+r   c                      t        j                  d      t              }d\  }}t        j                  dt        j
                        } |||f   |       | j                  |d   ||z         y )Nr=  r>  r   r  r   )r   r  r/   r  r'  r  assertEqualr  r  rA  rB  r   s        r   test_simple_gridsize1dz(TestCudaIntrinsic.test_simple_gridsize1d  sa    /488./0ABfhhq)s#Q$/r   zRequires too many threadsc                 0   t         j                  d        }t        j                  dt        j                        }t        j                  dt        j                        } |d   ||       | j                  |d   d       | j                  |d   d       y )Nc                    t        j                  d      }t         j                  j                  t         j                  j                  z  t         j
                  j                  z   }t        j                  d      }t         j                  j                  t         j                  j                  z  }||k7  rd| d<   ||k7  rd|d<   y y r,   )r   r%   blockIdxr   r4   r   r.   r3   )
grid_errorgridsize_errori1i2gs1gs2s         r   fz,TestCudaIntrinsic.test_issue_9229.<locals>.f  s    1B4==??2T^^5E5EEB--"C--//DLLNN2CRx !
1cz$%q! r   r   r  )i Q   r   )r   r  r  r'  uint64rL  )r  rX  rR  rS  s       r   test_issue_9229z!TestCudaIntrinsic.test_issue_9229  s     
	& 
	& XXaryy1
!2995 	-^4A**A.r   zTests PTX emissionc           	         t         d d  t         t         d d  f} t        j                  |      t              } t        j                  |      t              }d}d}t        j                  ddt
        j                         }|j                         }d|d d t        j                  |t
        j                         } ||df   |||       |j                  |      }	| j                  d	t        t        j                  d
|	                   t
        j                  j                  ||d       t        j                  |t
        j                         } ||df   |||       |j                  |      }	| j                  dt        t        j                  d
|	                   t
        j                  j                  ||d       y )N    r%     )r5   
fill_valuer   r   r$  r  r   r)   z	\s+bra\s+	branching)err_msgr   r   )r   r   r  r   r   r  fullrI  r  inspect_asmrL  r   refindalltestingassert_array_equal)
r  sigcu_branching_with_ifscu_branching_with_selpsnrH   r7   expectedrG   ptxs
             r   	test_selpzTestCudaIntrinsic.test_selp  sb   Qxa) -.@ A"/$((3-0D"EGG"288<668!IIarxx(#ad#Aq!,#//4C

< =>?


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

< =>?


%%a6%Br   c                 &    t        j                  d      t              }d}d}t        j                  dt        j
                        } |||f   |       | j                  |d   |d   |d   z         | j                  |d   |d   |d   z         y )Nr=  rG  rH  r)   r  r   r   )r   r  r1   r  r'  r  rL  rM  s        r   test_simple_gridsize2dz(TestCudaIntrinsic.test_simple_gridsize2d  s    /488./0ABhhq)s#QT!W!45QT!W!45r   c           	          t        j                  d      t              }d}d}|d   |d   z  |d   |d   z  f}t        j                  |t        j
                        } |||f   |       |\  }}|j                  \  }}	t        t        |d         t        |d               D ]Z  \  }
}||
z   ||z   }}t        ||	|      D ];  }t        |||      D ])  }| j                  |||f   ||z   k(  |||f   ||z   f       + = \ y rF  )
r   r  r>   r  r@  r  r5   zipr6   r  )r  r  rA  rB  r5   r   r:   r;   r<   r=   r   r    r8   r9   r   r   s                   r   test_intrinsic_forloop_stepz-TestCudaIntrinsic.test_intrinsic_forloop_step  s   1488012HIa6!9$d1gq	&9:hhuBHH-s#u		d1gd1g7DAq"QY	FF65%0vvu5AOOC1IQ$6QTAE8JK 6 1 8r   c                     t         j                  d        }t        j                  dt        j                        j                  ddd      } |d   |       t        j                  j                  |d       y )Nc                     t        j                  d      \  }}}t        j                  d      \  }}}||z  |z  | |||f<   y Nr   r-   )outr   r   r   rG   rH   r7   s          r   fooz*TestCudaIntrinsic.test_3dgrid.<locals>.foo	  s@    iilGAq!mmA&GAq!q519C1aLr   i  r  	   )r   r   r   rz  )r   r  r  r'  r  reshaperf  assert_equal)r  rx  arrs      r   test_3dgridzTestCudaIntrinsic.test_3dgrid  sa    		% 
	%
 hhvRXX.66q!Q?! !#&


V,r   c                    t         j                  d        }d\  }}}t        j                  ||z  |z  t        j                        j                  |||      } |d   |       | j                  t        j                  |             y )Nc                    t        j                  d      \  }}}t        j                  d      \  }}}|t         j                  j                  t         j
                  j                  t         j                  j                  z  z   k(  xr |t         j                  j                  t         j
                  j                  t         j                  j                  z  z   k(  xrS |t         j                  j                  t         j
                  j                  t         j                  j                  z  z   k(  }|t         j                  j                  t         j                  j                  z  k(  xrr |t         j                  j                  t         j                  j                  z  k(  xr8 |t         j                  j                  t         j                  j                  z  k(  }|xr || |||f<   y rv  )
r   r%   r.   r   r   rQ  r4   r   r   r3   )	rw  r   r   r   rG   rH   r7   grid_is_rightgridsize_is_rights	            r   rx  z,TestCudaIntrinsic.test_3dgrid_2.<locals>.foo  s?   iilGAq!mmA&GAq!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  )>->C1aLr   )   r%     r  ))r   r   r)   )r   r)   r   )r   r  r  r'  bool_r{  r  r  )r  rx  r   r   r   r}  s         r   test_3dgrid_2zTestCudaIntrinsic.test_3dgrid_2  st    		? 
	? (1ahhA	"((3;;Aq!D! !#&s$r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y )Nvoid(int32[:], uint32)r   r  r     r   r   r   r  rB   r  r'  r  rL  r  s      r   test_popc_u4zTestCudaIntrinsic.test_popc_u4)  sP    548845kBhhq)sD!Q#r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y )Nzvoid(int32[:], uint64)r   r  r  l        @ r   r   r  r  s      r   test_popc_u8zTestCudaIntrinsic.test_popc_u8/  sP    548845kBhhq)sN+Q#r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |ddd       t        j                  j                  |d   d	       y )
Nzvoid(f4[:], f4, f4, f4)r   r  r         @      @      @r   r  )r   r  rI   r  r'  float32rf  assert_allcloser  s      r   test_fma_f4zTestCudaIntrinsic.test_fma_f45  X    648856zBhhq

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


""3q695r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |ddd       t        j                  j                  |d   d	       y )
Nzvoid(f8[:], f8, f8, f8)r   r  r  r  r  r  r   r  )r   r  rI   r  r'  float64rf  r  r  s      r   test_fma_f8zTestCudaIntrinsic.test_fma_f8;  r  r   c                     t        j                  d      t              }t        j                  dt        j
                        }t        j                  dgt        j
                        }t        j                  dgt        j
                        } |d   |||       t        j                  j                  |d   ||z          y Nvoid(f2[:], f2[:], f2[:])r   r  r  r  r  r   )	r   r  rO   r  r'  float16arrayrf  r  r  r  r   arg1arg2s        r   	test_haddzTestCudaIntrinsic.test_haddA      848878Ehhq

+xxBJJ/xxBJJ/sD$'


""3q64$;7r   c                 J    t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      }t        j
                  d      } |d   |||       ||z   }t        j                  j                  |d   |       y )Nvoid(f2[:], f2, f2)r   r  JM!	@r  r  r   )r   r  rR   r  r'  r  rf  r  r  r  r   r  r  refs         r   test_hadd_scalarz"TestCudaIntrinsic.test_hadd_scalarJ  s{    2488123EFhhq

+zz)$zz"~sD$'Tk


""3q63/r   z(Compilation unsupported in the simulatorc                 z    t         d d  t         t         f}t        t        |d      \  }}| j                  d|       y )Nr$  r   cczadd.f16)r   r   rR   assertInr  argsrm  _s       r   test_hadd_ptxzTestCudaIntrinsic.test_hadd_ptxT  2    1r2/&AQi%r   c                     t        j                  d      t              }t        j                  dt        j
                        }t        j                  dgt        j
                        }t        j                  dgt        j
                        }t        j                  dgt        j
                        } |d   ||||       t        j                  j                  |d   ||z  |z          y )	Nz void(f2[:], f2[:], f2[:], f2[:])r   r  r  r  r  r  r   )	r   r  rW   r  r'  r  r  rf  r  )r  r  r   r  r  arg3s         r   	test_hfmazTestCudaIntrinsic.test_hfmaZ  s    ?488>?Lhhq

+xxBJJ/xxBJJ/xxBJJ/sD$-


""3q64$;+=>r   c                 |    t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      }t        j
                  d      }t        j
                  d      } |d   ||||       ||z  |z   }t        j                  j                  |d   |       y )	Nzvoid(f2[:], f2, f2, f2)r   r  r  r  r  r  r   )r   r  rY   r  r'  r  rf  r  )r  r  r   r  r  r  r  s          r   test_hfma_scalarz"TestCudaIntrinsic.test_hfma_scalard  s    6488567IJhhq

+zz"~zz"~zz"~sD$-TkD 


""3q63/r   c                     t         d d  t         t         t         f}t        t        |d      \  }}| j                  d|       y )Nr  r  z
fma.rn.f16)r   r   rY   r  r  s       r   test_hfma_ptxzTestCudaIntrinsic.test_hfma_ptxo  s5    1r2r"/&AQlC(r   c                     t        j                  d      t              }t        j                  dt        j
                        }t        j                  dgt        j
                        }t        j                  dgt        j
                        } |d   |||       t        j                  j                  |d   ||z
         y r  )	r   r  r]   r  r'  r  r  rf  r  r  s        r   	test_hsubzTestCudaIntrinsic.test_hsubu  r  r   c                 J    t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      }t        j
                  d      } |d   |||       ||z
  }t        j                  j                  |d   |       y Nr  r   r  r  gQ?r  r   )r   r  r_   r  r'  r  rf  r  r  s         r   test_hsub_scalarz"TestCudaIntrinsic.test_hsub_scalar~  |    2488123EFhhq

+zz)$zz$sD$'Tk


""3q63/r   c                 z    t         d d  t         t         f}t        t        |d      \  }}| j                  d|       y )Nr  r  zsub.f16)r   r   r_   r  r  s       r   test_hsub_ptxzTestCudaIntrinsic.test_hsub_ptx  r  r   c                     t        j                         t              }t        j                  dt        j
                        }t        j                  dgt        j
                        }t        j                  dgt        j
                        } |d   |||       t        j                  j                  |d   ||z         y )Nr   r  r  r  r  r   )	r   r  rc   r  r'  r  r  rf  r  r  s        r   	test_hmulzTestCudaIntrinsic.test_hmul  s    488:k*hhq

+xxBJJ/xxBJJ/sD$'


""3q64$;7r   c                 J    t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      }t        j
                  d      } |d   |||       ||z  }t        j                  j                  |d   |       y r  )r   r  re   r  r'  r  rf  r  r  s         r   test_hmul_scalarz"TestCudaIntrinsic.test_hmul_scalar  r  r   c                 z    t         d d  t         t         f}t        t        |d      \  }}| j                  d|       y )Nr  r  zmul.f16)r   r   re   r  r  s       r   test_hmul_ptxzTestCudaIntrinsic.test_hmul_ptx  r  r   c                 J    t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      }t        j
                  d      } |d   |||       ||z  }t        j                  j                  |d   |       y r  )r   r  rh   r  r'  r  rf  r  r  s         r   test_hdiv_scalarz"TestCudaIntrinsic.test_hdiv_scalar  s|    2488123EFhhq

+zz)$zz$sD$'Tk


""3q63/r   c                     t        j                  d      t              }t        j                  j                  ddd      j                  t        j                        }t        j                  j                  ddd      j                  t        j                        }t        j                  |t        j                        } |j                  |j                        |||       ||z  }t        j                  j                  ||       y )Nr  i    i  rj   r  )r   r  rm   r  r  randintastyper  
zeros_likeforallrj   rf  r  )r  r  arry1arry2r   r  s         r   	test_hdivzTestCudaIntrinsic.test_hdiv  s    8488789KL		!!&%c!:AA"**M		!!&%c!:AA"**MmmE4!!#ue4em


""3,r   c                 8    t        j                  d      t              }t        j                  dt        j
                        }t        j                  dgt        j
                        } |d   ||       t        j                  j                  |d   |        y )Nvoid(f2[:], f2[:])r   r  r  r  r   )	r   r  rr   r  r'  r  r  rf  r  r  r  r   r  s       r   	test_hnegzTestCudaIntrinsic.test_hneg  sl    148801+>hhq

+xxBJJ/sD!


""3q6D51r   c                     t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      } |d   ||       | }t        j                  j                  |d   |       y )Nvoid(f2[:], f2)r   r  r  r  r   )r   r  ru   r  r'  r  rf  r  r  r  r   r  r  s        r   test_hneg_scalarz"TestCudaIntrinsic.test_hneg_scalar  sj    .488-./ABhhq

+zz)$sD!e


""3q63/r   c                 p    t         d d  t         f}t        t        |d      \  }}| j                  d|       y )Nr  r  zneg.f16)r   r   ru   r  r  s       r   test_hneg_ptxzTestCudaIntrinsic.test_hneg_ptx  0    1r{/&AQi%r   c                 F    t        j                         t              }t        j                  dt        j
                        }t        j                  dgt        j
                        } |d   ||       t        j                  j                  |d   t        |             y )Nr   r        r  r   )
r   r  rz   r  r'  r  r  rf  r  absr  s       r   	test_habszTestCudaIntrinsic.test_habs  sj    488:k*hhq

+xxRZZ0sD!


""3q63t95r   c                 *    t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      } |d   ||       t        |      }t        j                  j                  |d   |       y )Nr  r   r  gJM!	r  r   )	r   r  r|   r  r'  r  r  rf  r  r  s        r   test_habs_scalarz"TestCudaIntrinsic.test_habs_scalar  sl    .488-./ABhhq

+zz*%sD!$i


""3q63/r   c                 p    t         d d  t         f}t        t        |d      \  }}| j                  d|       y )Nr  r  zabs.f16)r   r   r|   r  r  s       r   test_habs_ptxzTestCudaIntrinsic.test_habs_ptx  r  r   c                 
   t         t        t        t        t        t
        t        t        t        t        t        t        f}t        t        f}t        j                  t        j                   t        j"                  t        j$                  t        j&                  t        j(                  t        j*                  t        j,                  t        j.                  t        j0                  t        j2                  t4        f}t        j6                  t        j8                  f}d}t        j:                  j=                  d       t        j:                  j?                  dd|      jA                  t        jB                        }t        jD                  |      }tG        ||      D ]  \  }}	| jI                  |	      5   tK        jL                  d      |      } |d|f   ||        |	|t        jB                        }
t        jN                  jQ                  ||
       d d d         t        j:                  j?                  dd|      jA                  t        jB                        }tG        ||      D ]  \  }}	| jI                  |	      5   tK        jL                  d      |      } |d|f   ||        |	|t        jB                        }
t        jN                  jQ                  ||
       d d d         y # 1 sw Y   ^xY w# 1 sw Y   xY w)	Nr]  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  rr  subTestr   r  rf  r  )r  kernelsexp_kernelsexpected_functionsexpected_exp_functionsr   r   r   kernelr  rl  x2s               r   test_fp16_intrinsics_commonz-TestCudaIntrinsic.test_fp16_intrinsics_common  s   m}| 	"
 #L1 ffbff ffbggrxx ggrww mmRXXrww*	,
 #%&&"''!2 
		qIIaQ/66rzzBMM!g'9:JFB$7"67?qsAq!arzz2

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

3

**1h7	 %$ C %$ %$s   0A"K,?A"K9,K6	9L	c                 ~   t        j                         d        }d}t        j                  j	                  d       t        j                  j                  |      j                  t        j                        }t        j                  |      } |d|f   ||       t        j                  j                  |d|z         y )Nc                     t        j                  d      }|t        |       k  r&t         j                  j	                  ||         | |<   y y r   )r   r%   r   rL   hexp10r   s      r   hexp10_vectorsz5TestCudaIntrinsic.test_hexp10.<locals>.hexp10_vectors  s;    		!A3q6zyy''!-! r   r]  r   r  )r   r  r  r  r  randr  r  r  rf  r  )r  r  r   r   r   s        r   test_hexp10zTestCudaIntrinsic.test_hexp10  s    		. 
	. 
		qIINN1$$RZZ0MM! 	q!tQ"


""1bAg.r   c                    t         t        t        t        t        t
        f}t        j                  t        j                  t        j                  t        j                  t        j                  t        j                  f}t        ||      D ]F  \  }}| j                  |      5   t        j                   d      |      }t#        j$                  dt"        j&                        }t#        j$                  dt"        j&                        }t#        j(                  d      }t#        j(                  d      }	t#        j(                  d      }
 |d   ||	|	        ||	|	      }| j+                  ||d	           |d   ||	|
        ||	|
      }| j+                  ||d	           |d   ||	|        ||	|      }| j+                  ||d	          d d d        I y # 1 sw Y   UxY w)
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ltrr  r  r   r  r  r'  r  r  rL  )r  fnsopsr  r  r  rl  gotr  r  arg4s              r   test_fp16_comparisonz&TestCudaIntrinsic.test_fp16_comparison!  sz    "35F "35FH{{HKK{{HKK6 #smFB$8"78<88ARXX6hhq1zz!}zz!}zz!} tS$-dD>  3q62 tS$-dD>  3q62 tS$-dD>  3q62- %$ $$$s   "D&GG	c                    t         t        t        t        t        f}|D ]  }| j                  |      5   t        j                  d      |      }t        j                  dt        j                        }t        j                  d      }t        j                  d      }t        j                  d      } |d   ||||       | j                  |d	          d d d         y # 1 sw Y   xY w)
Nr  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_comparisonsz3TestCudaIntrinsic.test_multiple_float16_comparisonsA  s    )))))	+	
 B$>488$=>rBhhq1zz"~zz"~zz"~sD$5A' %$ $$s   B"C  C)	c                     t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      }t        j
                  d      } |d   |||       t        j                  j                  |d   |       t        j
                  d      } |d   |||       t        j                  j                  |d   |       y 	Nr  r   r  r  r  r  r   g      @)r   r  r   r  r'  r  rf  r  r  s        r   	test_hmaxzTestCudaIntrinsic.test_hmaxR      2488123EFhhq

+zz"~zz"~sD$'


""3q640zz"~sD$'


""3q640r   c                     t        j                  d      t              }t        j                  dt        j
                        }t        j
                  d      }t        j
                  d      } |d   |||       t        j                  j                  |d   |       t        j
                  d      } |d   |||       t        j                  j                  |d   |       y r  )r   r  r   r  r'  r  rf  r  r  s        r   	test_hminzTestCudaIntrinsic.test_hmin^  r  r   c                      t        j                  d      t              }t        j                  dt        j
                        }d} |d   ||       t        j                  j                  |d   |dz         y )Nzvoid(float32[:], float32)r   r  r  r  r   UUUUUU?)r   r  r   r  r'  r  rf  r  r  r  r   cbrt_args       r   test_cbrt_f32zTestCudaIntrinsic.test_cbrt_f32j  ^    848878Ehhq

+sH%


""3q68+>?r   c                      t        j                  d      t              }t        j                  dt        j
                        }d} |d   ||       t        j                  j                  |d   |dz         y )Nzvoid(float64[:], float64)r   r  g      @r  r   r  )r   r  r   r  r'  r  rf  r  r  s       r   test_cbrt_f64zTestCudaIntrinsic.test_cbrt_f64q  r   r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y )Nzvoid(uint32[:], uint32)r   r  r  i0  r   i  )r   r  r   r  r'  uint32rL  r  s      r   test_brev_u4zTestCudaIntrinsic.test_brev_u4x  sP    648856{Chhq		*sJ'Q,r   z.only get given a Python "int", assumes 32 bitsc                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y )Nzvoid(uint64[:], uint64)r   r  r  l   0  C r   l       `x)r   r  r   r  r'  rZ  rL  r  s      r   test_brev_u8zTestCudaIntrinsic.test_brev_u8~  sR    648856{Chhq		*s./Q!34r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y )Nvoid(int32[:], int32)r   r  r     r      r   r  r   r  r'  r  rL  r  s      r   test_clz_i4zTestCudaIntrinsic.test_clz_i4  sP    448834Z@hhq)sJ'Q$r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y)	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  r*  r   r+  Nr,  r  s      r   test_clz_u4zTestCudaIntrinsic.test_clz_u4  sR     648845jAhhq)sJ'Q$r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y Nr)  r   r  r  l    r   r,  r  s      r   test_clz_i4_1sz TestCudaIntrinsic.test_clz_i4_1s  P    448834Z@hhq)sJ'Q#r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   dd       y )Nr)  r   r  r  r   r]  CUDA semanticsr,  r  s      r   test_clz_i4_0sz TestCudaIntrinsic.test_clz_i4_0s  sS    448834Z@hhq)sC Q%56r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y )Nvoid(int32[:], int64)r   r  r     r   /   r,  r  s      r   test_clz_i8zTestCudaIntrinsic.test_clz_i8  sQ    448834Z@hhq)s-.Q$r   c                     t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d        |d   |d       | j                  |d   d	       y )
Nr)  r   r  r  r*  r              r]  r   r  r   r  r'  r  rL  r  s      r   test_ffs_i4zTestCudaIntrinsic.test_ffs_i4  st    448834Z@hhq)sJ'Q$sJ'Q$r   c                     t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d        |d   |d       | j                  |d   d	       y )
Nr  r   r  r  r*  r   r=  r>  r]  r?  r  s      r   test_ffs_u4zTestCudaIntrinsic.test_ffs_u4  st    548845jAhhq)sJ'Q$sJ'Q$r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y r1  r?  r  s      r   test_ffs_i4_1sz TestCudaIntrinsic.test_ffs_i4_1s  r3  r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d       y )Nr)  r   r  r  r   r?  r  s      r   test_ffs_i4_0sz TestCudaIntrinsic.test_ffs_i4_0s  sP    448834Z@hhq)sC Q#r   c                     t        j                  d      t              }t        j                  dt        j
                        } |d   |d       | j                  |d   d        |d   |d       | j                  |d   d	       y )
Nr8  r   r  r  r9  r   r^  l        !   r?  r  s      r   test_ffs_i8zTestCudaIntrinsic.test_ffs_i8  su    448834Z@hhq)s-.Q$sK(Q$r   c                 v    t        j                  d      t              }d}t        j                  |dz  t        j
                        }t        j                  t        j                  dt        j
                        |      } |d|dz  f   |       | j                  t        j                  ||k(               y )Nr  r)   r]  r  r   )
r   r  r   r  r'  r  tiler  r  r  )r  r  countr   r!  s        r   test_simple_laneidz$TestCudaIntrinsic.test_simple_laneid  s    -488,-m<hhurz2ggbii"((3U;EBJ$scz*+r   c                      t        j                  d      t              }t        j                  dt        j
                        } |d   |       | j                  |d   dd       y )Nr  r   r  r  r   r]  r5  )r   r  r  r  r'  r  rL  r  s      r   test_simple_warpsizez&TestCudaIntrinsic.test_simple_warpsize  sQ    -488,-o>hhq)sQ%56r   c                      t        j                  d      t              }t        j                  dt        j
                        }dD ],  } |d   ||       | j                  |d   t        |             . y )Nzvoid(int64[:], float32)r   r  r  g      g      g      g      ?g      @g      @g      @r  r   r   r  r   r  r'  r   rL  r   r  r  r   r   s       r   test_round_f4zTestCudaIntrinsic.test_round_f4  _    648856|Dhhq)@AHTN3"SVU1X. Ar   c                      t        j                  d      t              }t        j                  dt        j
                        }dD ],  } |d   ||       | j                  |d   t        |             . y )Nzvoid(int64[:], float64)r   r  rQ  r  r   rR  rS  s       r   test_round_f8zTestCudaIntrinsic.test_round_f8  rU  r   c           	          t        j                  d      t              }t        j                  dt        j
                        }t        j                  j                  d       t        j                  j                  d      j                  t        j
                        }t        j                  |t        j                  t        j                  t        j                   t        j                  g      f       d}t        j                  ||      D ]O  \  }}| j                  ||      5   |d   |||       | j!                  |d	   t#        ||      d
       d d d        Q y # 1 sw Y   \xY w)N void(float32[:], float32, int32)r   r  {   r]  )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digitsra  r   s          r   test_round_to_f4z"TestCudaIntrinsic.test_round_to_f4  s   ?488>?Phhq

+
		syy#**2::6
bhh'@ABC
 &--dF;LC#w7sC1''Ac70C-5 ( 7 87 <77s   )/E##E,	z$Overflow behavior differs on CPythonc                 .    t        j                  d      t              }t        j                  dt        j
                        }t        j                  t        j
                        j                  }d} |d   |||       | j                  |d   |       y )NrY  r   r  i,  r  r   )	r   r  r   r  r'  r  finfomaxrL  r  r  r   ra  r   s        r   test_round_to_f4_overflowz+TestCudaIntrinsic.test_round_to_f4_overflow	  st     @488>?Phhq

+hhrzz"&& sC)Q%r   c                      t        j                  d      t              }t        j                  dt        j
                        }d}d} |d   |||       | j                  |d   t        ||      d	       y )
NrY  r   r  gQ?r   r  r   rb  rc  )r   r  r   r  r'  r  rj  r   rr  s        r   test_round_to_f4_halfwayz*TestCudaIntrinsic.test_round_to_f4_halfway  sj    ?488>?Phhq

+ sC)Ac7(;(Kr   c           	      X    t        j                  d      t              }t        j                  dt        j
                        }t        j                  j                  d       t        j                  j                  d      }t        j                  |t        j                  t        j                  t        j                   t        j                  g      f       d}t        j                  ||      D ]O  \  }}| j                  ||      5   |d   |||       | j                  |d	   t!        ||      d
       d d d        Q d}d}| j                  ||      5   |d   |||       | j                  |d	   t!        ||      d       d d d        y # 1 sw Y   xY w# 1 sw Y   y xY w)N void(float64[:], float64, int32)r   r  rZ  r]  )r[  r\  r]  r^  r_  r   r   r)   r   r   r$  r`  r  r   exactrc  g`8p=<   double)r   r  r   r  r'  r  r  r  re  r  rf  rg  rh  ri  r  rj  r   rk  s          r   test_round_to_f8z"TestCudaIntrinsic.test_round_to_f8!  s[   ?488>?Phhq

+
		syy#
bhh'@ABC7%--dF;LC#w7sC1''Ac70C-4 ( 6 87 < +\\c7\3HTN3W-##CFE#w,?)1 $ 3 43 87 43s   /F/F F	 F)c                 .    t        j                  d      t              }t        j                  dt        j
                        }t        j                  t        j
                        j                  }d} |d   |||       | j                  |d   |       y )Nrw  r   r  r  r  r   )	r   r  r   r  r'  r  rp  rq  rL  rr  s        r   test_round_to_f8_overflowz+TestCudaIntrinsic.test_round_to_f8_overflow8  st     @488>?Phhq

+hhrzz"&& sC)Q%r   c                      t        j                  d      t              }t        j                  dt        j
                        }d}d} |d   |||       | j                  |d   t        ||      d	       y )
Nrw  r   r  g\(\?r   r  r   rz  rc  )r   r  r   r  r'  r  rj  r   rr  s        r   test_round_to_f8_halfwayz*TestCudaIntrinsic.test_round_to_f8_halfwayE  sj    ?488>?Phhq

+ sC)Ac7(;(Kr   )K__name__
__module____qualname__r  r  r"  r2  r
   r8  r;  rD  rJ  rN  r[  rn  rp  rs  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/  r2  r6  r;  r@  rB  rD  rF  rI  rM  rO  rT  rW  rn  rs  ru  r{  r}  r  __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r   r  __main__)Yrh  numpyr  r  rd  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HMMO r   