U
    hâËd�ã  ã                   @   s‚  d dl Zd dlmZ d dlmZmZmZmZm	Z	 d dl
mZmZmZmZ d dlmZ ejdd�dd	„ ƒZejdd�d
d„ ƒZejdd�dd„ ƒZejdd�dd„ ƒZejdd�dd„ ƒZejdd�dd„ ƒZejdd�dd„ ƒZejdd�dd„ ƒZdd„ Zdd„ Zdd„ Zdd„ Zd d!„ Zd"d#„ Zd$d%„ Z d&d'„ Z!d(d)„ Z"d*d+„ Z#d,d-„ Z$d.d/„ Z%d0d1„ Z&d2d3„ Z'd4d5„ Z(d6d7„ Z)d8d9„ Z*d:d;„ Z+d<d=„ Z,d>d?„ Z-d@dA„ Z.dBdC„ Z/dDdE„ Z0dFdG„ Z1dHdI„ Z2dJdK„ Z3dLdM„ Z4dNdO„ Z5dPdQ„ Z6dRdS„ Z7dTdU„ Z8dVdW„ Z9dXdY„ Z:dZd[„ Z;d\d]„ Z<d^d_„ Z=d`da„ Z>dbdc„ Z?ddde„ Z@dfdg„ ZAdhdi„ ZBdjdk„ ZCdldm„ ZDdndo„ ZEdpdq„ ZFdrds„ ZGdtdu„ ZHdvdw„ ZIdxdy„ ZJdzd{„ ZKd|d}„ ZLd~d„ ZMd€d�„ ZNd‚dƒ„ ZOd„d…„ ZPd†d‡„ ZQdˆd‰„ ZRdŠd‹„ ZSdŒd�„ ZTdŽd�„ ZUd�d‘„ ZVd’d“„ ZWd”d•„ ZXd–d—„ ZYd˜d™„ ZZdšd›„ Z[e[dœƒ\Z\Z]Z^Z_e[d�ƒ\Z`ZaZbZce[džƒ\ZdZeZfZge[dŸƒ\ZhZiZjZkd d¡„ Zld¢d£„ Zmd¤d¥„ ZnG d¦d§„ d§eƒZoepd¨k�r~e q¡  dS )©é    N)Údedent)ÚcudaÚuint32Úuint64Úfloat32Úfloat64)ÚunittestÚCUDATestCaseÚskip_unless_cc_50Úcc_X_or_above)ÚconfigT)Zdevicec                 C   s   t | ƒS ©N)r   ©Únum© r   ú]/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_atomics.pyÚatomic_cast_to_uint64
   s    r   c                 C   s   t | ƒS r   )Úintr   r   r   r   Úatomic_cast_to_int   s    r   c                 C   s   | S r   r   r   r   r   r   Úatomic_cast_none   s    r   c	                 C   sf   t jj}	t j ||¡}
||
|	< t  ¡  |||	 | ƒ}|rB|| }||
||ƒ t  ¡  |
|	 | |	< d S r   ©r   Ú	threadIdxÚxÚsharedÚarrayÚsyncthreads)ÚaryÚidxÚop2Ú	ary_dtypeÚary_nelementsÚ
binop_funcÚ	cast_funcZinitializerÚneg_idxÚtidÚsmÚbinr   r   r   Úatomic_binary_1dim_shared   s    r'   c           
      C   s^   t jj}t j ||¡}| | ||< t  ¡  ||| | ƒ}	|||	|ƒ t  ¡  || | |< d S r   r   )
r   r   r   r   r    r!   r"   r$   r%   r&   r   r   r   Úatomic_binary_1dim_shared2)   s    r(   c                 C   s–   t jj}t jj}t j ||¡}	| ||f |	||f< t  ¡  |||ƒf}
|rj|
d |d  |
d |d  f}
||	|
|ƒ t  ¡  |	||f | ||f< d S ©Nr   é   )r   r   r   Úyr   r   r   )r   r   r   Z	ary_shaper!   Úy_cast_funcr#   ÚtxÚtyr%   r&   r   r   r   Úatomic_binary_2dim_shared6   s     r/   c                 C   sT   t jj}t jj}|||ƒf}|rD|d | jd  |d | jd  f}|| ||ƒ d S r)   )r   r   r   r+   Úshape)r   r   r!   r,   r#   r-   r.   r&   r   r   r   Úatomic_binary_2dim_globalF   s    $r1   c                 C   s4   t jj}t|| | ƒ}|r$|| }|| ||ƒ d S r   )r   r   r   r   )r   r   r    r   r!   r#   r$   r&   r   r   r   Úatomic_binary_1dim_globalP   s
    r2   c              
   C   s    t | | dtdtjjtddƒ	 d S ©Nr*   é    r   F©r'   r   r   ÚatomicÚaddr   ©r   r   r   r   Ú
atomic_addZ   s       ÿr9   c              
   C   s    t | | dtdtjjtddƒ	 d S )Nr*   r4   r   Tr5   r8   r   r   r   Úatomic_add_wrap_   s       ÿr:   c                 C   s   t | dtdtjjtdƒ d S ©Nr*   ©é   é   F©r/   r   r   r6   r7   r   r8   r   r   r   Úatomic_add2d   s
    
  ÿr@   c                 C   s   t | dtdtjjtdƒ d S )Nr*   r<   Tr?   r8   r   r   r   Úatomic_add2_wrapi   s
    
  ÿrA   c                 C   s   t | dtdtjjtdƒ d S r;   )r/   r   r   r6   r7   r   r8   r   r   r   Úatomic_add3n   s
    
  ÿrB   c              
   C   s    t | | dtdtjjtddƒ	 d S ©Nç      ð?r4   ç        F©r'   r   r   r6   r7   r   r8   r   r   r   Úatomic_add_floats   s       ÿrG   c              
   C   s    t | | dtdtjjtddƒ	 d S ©NrD   r4   rE   TrF   r8   r   r   r   Úatomic_add_float_wrapx   s       ÿrI   c                 C   s   t | dtdtjjtdƒ d S ©NrD   r<   F©r/   r   r   r6   r7   r   r8   r   r   r   Úatomic_add_float_2}   s
    
  ÿrL   c                 C   s   t | dtdtjjtdƒ d S ©NrD   r<   TrK   r8   r   r   r   Úatomic_add_float_2_wrap‚   s
    
  ÿrN   c                 C   s   t | dtdtjjtdƒ d S rJ   )r/   r   r   r6   r7   r   r8   r   r   r   Úatomic_add_float_3‡   s
    
  ÿrO   c                 C   s   t || ddtjjdƒ d S ©Nr4   rD   F©r2   r   r6   r7   ©r   r   r   r   r   Úatomic_add_double_globalŒ   s    rS   c                 C   s   t || ddtjjdƒ d S )Nr4   rD   TrQ   rR   r   r   r   Úatomic_add_double_global_wrap�   s    rT   c                 C   s   t | dtjjtdƒ d S ©Nr*   F©r1   r   r6   r7   r   r8   r   r   r   Úatomic_add_double_global_2”   s    rW   c                 C   s   t | dtjjtdƒ d S )Nr*   TrV   r8   r   r   r   Úatomic_add_double_global_2_wrap˜   s    rX   c                 C   s   t | dtjjtdƒ d S rU   )r1   r   r6   r7   r   r8   r   r   r   Úatomic_add_double_global_3œ   s    ÿrY   c              
   C   s    t || dtdtjjtddƒ	 d S rC   ©r'   r   r   r6   r7   r   rR   r   r   r   Úatomic_add_double¡   s       ÿr[   c              
   C   s    t || dtdtjjtddƒ	 d S rH   rZ   rR   r   r   r   Úatomic_add_double_wrap¦   s       ÿr\   c                 C   s   t | dtdtjjtdƒ d S rJ   ©r/   r   r   r6   r7   r   r8   r   r   r   Úatomic_add_double_2«   s
    
  ÿr^   c                 C   s   t | dtdtjjtdƒ d S rM   r]   r8   r   r   r   Úatomic_add_double_2_wrap°   s
    
  ÿr_   c                 C   s   t | dtdtjjtdƒ d S rJ   )r/   r   r   r6   r7   r   r8   r   r   r   Úatomic_add_double_3µ   s
    
  ÿr`   c              
   C   s    t | | dtdtjjtddƒ	 d S r3   )r'   r   r   r6   Úsubr   r8   r   r   r   Ú
atomic_subº   s       ÿrb   c                 C   s   t | dtdtjjtdƒ d S r;   )r/   r   r   r6   ra   r   r8   r   r   r   Úatomic_sub2¿   s
    
  ÿrc   c                 C   s   t | dtdtjjtdƒ d S r;   )r/   r   r   r6   ra   r   r8   r   r   r   Úatomic_sub3Ä   s
    
  ÿrd   c              
   C   s    t | | dtdtjjtddƒ	 d S rC   )r'   r   r   r6   ra   r   r8   r   r   r   Úatomic_sub_floatÉ   s       ÿre   c                 C   s   t | dtdtjjtdƒ d S rJ   )r/   r   r   r6   ra   r   r8   r   r   r   Úatomic_sub_float_2Î   s
    
  ÿrf   c                 C   s   t | dtdtjjtdƒ d S rJ   )r/   r   r   r6   ra   r   r8   r   r   r   Úatomic_sub_float_3Ó   s
    
  ÿrg   c              
   C   s    t || dtdtjjtddƒ	 d S rC   )r'   r   r   r6   ra   r   rR   r   r   r   Úatomic_sub_doubleØ   s       ÿrh   c                 C   s   t | dtdtjjtdƒ d S rJ   )r/   r   r   r6   ra   r   r8   r   r   r   Úatomic_sub_double_2Ý   s
    
  ÿri   c                 C   s   t | dtdtjjtdƒ d S rJ   ©r/   r   r   r6   ra   r   r8   r   r   r   Úatomic_sub_double_3â   s
    
  ÿrk   c                 C   s   t || ddtjjdƒ d S rP   )r2   r   r6   ra   rR   r   r   r   Úatomic_sub_double_globalç   s    rl   c                 C   s   t | dtjjtdƒ d S )NrD   F)r1   r   r6   ra   r   r8   r   r   r   Úatomic_sub_double_global_2ë   s    ÿrm   c                 C   s   t | dtdtjjtdƒ d S rJ   rj   r8   r   r   r   Úatomic_sub_double_global_3ð   s
    
  ÿrn   c              
   C   s    t | | |tdtjjtddƒ	 d S )Nr4   r*   F)r'   r   r   r6   Úand_r   ©r   r   r   r   r   Ú
atomic_andõ   s       ÿrq   c                 C   s   t | |tdtjjtdƒ d S ©Nr<   F)r/   r   r   r6   ro   r   rp   r   r   r   Úatomic_and2ú   s
    
  ÿrs   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   ro   r   rp   r   r   r   Úatomic_and3ÿ   s
    
  ÿrt   c                 C   s   t || d|tjjdƒ d S ©Nr4   F)r2   r   r6   ro   ©r   r   r   r   r   r   Úatomic_and_global  s    rw   c                 C   s   t | |tjjtdƒ d S ©NF)r1   r   r6   ro   r   rp   r   r   r   Úatomic_and_global_2  s     ÿry   c              
   C   s    t | | |tdtjjtddƒ	 d S ©Nr4   r   F)r'   r   r   r6   Úor_r   rp   r   r   r   Ú	atomic_or  s       ÿr|   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r{   r   rp   r   r   r   Ú
atomic_or2  s
    
  ÿr}   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r{   r   rp   r   r   r   Ú
atomic_or3  s
    
  ÿr~   c                 C   s   t || d|tjjdƒ d S ru   )r2   r   r6   r{   rv   r   r   r   Úatomic_or_global  s    r   c                 C   s   t | |tjjtdƒ d S rx   )r1   r   r6   r{   r   rp   r   r   r   Úatomic_or_global_2   s     ÿr€   c              
   C   s    t | | |tdtjjtddƒ	 d S rz   )r'   r   r   r6   Úxorr   rp   r   r   r   Ú
atomic_xor%  s       ÿr‚   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r�   r   rp   r   r   r   Úatomic_xor2*  s
    
  ÿrƒ   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r�   r   rp   r   r   r   Úatomic_xor3/  s
    
  ÿr„   c                 C   s   t || d|tjjdƒ d S ru   )r2   r   r6   r�   rv   r   r   r   Úatomic_xor_global4  s    r…   c                 C   s   t | |tjjtdƒ d S rx   )r1   r   r6   r�   r   rp   r   r   r   Úatomic_xor_global_28  s     ÿr†   c                 C   s   t | ||tdtjjtƒ d S ©Nr4   )r(   r   r   r6   Úincr   ©r   r   r   r   r   r   Úatomic_inc32=  s     ÿrŠ   c                 C   s   t | ||tdtjjtƒ d S r‡   )r(   r   r   r6   rˆ   r   r‰   r   r   r   Úatomic_inc64B  s     ÿr‹   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   rˆ   r   rp   r   r   r   Úatomic_inc2_32G  s
    
  ÿrŒ   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   rˆ   r   rp   r   r   r   Úatomic_inc2_64L  s
    
  ÿr�   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   rˆ   r   rp   r   r   r   Úatomic_inc3Q  s
    
  ÿrŽ   c                 C   s   t || d|tjjdƒ d S ru   )r2   r   r6   rˆ   rv   r   r   r   Úatomic_inc_globalV  s    r�   c                 C   s   t | |tjjtdƒ d S rx   )r1   r   r6   rˆ   r   rp   r   r   r   Úatomic_inc_global_2Z  s     ÿr�   c                 C   s   t | ||tdtjjtƒ d S r‡   )r(   r   r   r6   Údecr   r‰   r   r   r   Úatomic_dec32_  s     ÿr’   c                 C   s   t | ||tdtjjtƒ d S r‡   )r(   r   r   r6   r‘   r   r‰   r   r   r   Úatomic_dec64d  s     ÿr“   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r‘   r   rp   r   r   r   Úatomic_dec2_32i  s
    
  ÿr”   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r‘   r   rp   r   r   r   Úatomic_dec2_64n  s
    
  ÿr•   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r‘   r   rp   r   r   r   Úatomic_dec3s  s
    
  ÿr–   c                 C   s   t || d|tjjdƒ d S ru   )r2   r   r6   r‘   rv   r   r   r   Úatomic_dec_globalx  s    r—   c                 C   s   t | |tjjtdƒ d S rx   )r1   r   r6   r‘   r   rp   r   r   r   Úatomic_dec_global_2|  s     ÿr˜   c                 C   s   t | ||tdtjjtƒ d S r‡   )r(   r   r   r6   Úexchr   r‰   r   r   r   Úatomic_exch�  s     ÿrš   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r™   r   rp   r   r   r   Úatomic_exch2†  s
    
  ÿr›   c                 C   s   t | |tdtjjtdƒ d S rr   )r/   r   r   r6   r™   r   rp   r   r   r   Úatomic_exch3‹  s
    
  ÿrœ   c                 C   s   t || d|tjjdƒ d S ru   )r2   r   r6   r™   rv   r   r   r   Úatomic_exch_global�  s    r�   c                 C   sD   t dƒj| d�}i }t|tttdœ|ƒ |d |d |d |d fS )Na  
    def atomic(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, bx])

    def atomic_double_normalizedindex(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, uint64(bx)])

    def atomic_double_oneindex(res, ary):
        tx = cuda.threadIdx.x
        {func}(res, 0, ary[tx])

    def atomic_double_shared(res, ary):
        tid = cuda.threadIdx.x
        smary = cuda.shared.array(32, float64)
        smary[tid] = ary[tid]
        smres = cuda.shared.array(1, float64)
        if tid == 0:
            smres[0] = res[0]
        cuda.syncthreads()
        {func}(smres, 0, smary[tid])
        cuda.syncthreads()
        if tid == 0:
            res[0] = smres[0]
    )Úfunc)r   r   r   r6   Zatomic_double_normalizedindexZatomic_double_oneindexZatomic_double_shared)r   ÚformatÚexecr   r   r   )rž   ÚfnsÚldr   r   r   Úgen_atomic_extreme_funcs”  s    å ÿr£   zcuda.atomic.maxzcuda.atomic.minzcuda.atomic.nanmaxzcuda.atomic.nanminc                 C   s8   t  d¡}|| jk r4t j | |d … ||| ¡||< d S ©Nr*   )r   ÚgridÚsizer6   Zcompare_and_swap©ÚresÚoldr   Zfill_valÚgidr   r   r   Úatomic_compare_and_swapÄ  s    

r«   c                 C   s2   t  d¡}|| jk r.t j | |||| ¡||< d S r¤   )r   r¥   r¦   r6   Úcasr§   r   r   r   Úatomic_cas_1dimÊ  s    

r­   c                 C   sL   t  d¡}|d | jd k rH|d | jd k rHt j | |||| ¡||< d S )Né   r   r*   )r   r¥   r0   r6   r¬   r§   r   r   r   Úatomic_cas_2dimÐ  s    
$r¯   c                       s  e Zd Z‡ fdd„Zdd„ Zdd„ Zdd„ Zd	d
„ Zdd„ Zdd„ Z	�d)dd„Z
edd„ ƒZdd„ Zdd„ Zedd„ ƒZdd„ Zdd„ Zdd„ Zd d!„ Zd"d#„ Zd$d%„ Zd&d'„ Zd(d)„ Zd*d+„ Zd,d-„ Zd.d/„ Zd0d1„ Zd2d3„ Zd4d5„ Zd6d7„ Zd8d9„ Zd:d;„ Z d<d=„ Z!d>d?„ Z"d@dA„ Z#dBdC„ Z$dDdE„ Z%dFdG„ Z&dHdI„ Z'dJdK„ Z(dLdM„ Z)dNdO„ Z*dPdQ„ Z+dRdS„ Z,dTdU„ Z-dVdW„ Z.dXdY„ Z/dZd[„ Z0d\d]„ Z1d^d_„ Z2d`da„ Z3dbdc„ Z4ddde„ Z5dfdg„ Z6dhdi„ Z7djdk„ Z8dldm„ Z9dndo„ Z:dpdq„ Z;drds„ Z<dtdu„ Z=dvdw„ Z>dxdy„ Z?dzd{„ Z@d|d}„ ZAd~d„ ZBd€d�„ ZCd‚dƒ„ ZDd„d…„ ZEd†d‡„ ZFdˆd‰„ ZGdŠd‹„ ZHdŒd�„ ZIdŽd�„ ZJd�d‘„ ZKd’d“„ ZLd”d•„ ZMd–d—„ ZNd˜d™„ ZOdšd›„ ZPdœd�„ ZQdždŸ„ ZRd d¡„ ZSd¢d£„ ZTd¤d¥„ ZUd¦d§„ ZVd¨d©„ ZWdªd«„ ZXd¬d­„ ZYd®d¯„ ZZd°d±„ Z[d²d³„ Z\d´dµ„ Z]d¶d·„ Z^d¸d¹„ Z_dºd»„ Z`d¼d½„ Zad¾d¿„ ZbdÀdÁ„ ZcdÂdÃ„ Zd�d*dÅdÆ„ZedÇdÈ„ ZfdÉdÊ„ ZgdËdÌ„ ZhdÍdÎ„ ZidÏdÐ„ ZjdÑdÒ„ ZkdÓdÔ„ ZldÕdÖ„ Zmd×dØ„ ZndÙdÚ„ ZodÛdÜ„ ZpdÝdÞ„ Zqdßdà„ Zrdádâ„ Zsdãdä„ Ztdådæ„ Zudçdè„ Zvdédê„ Zwdëdì„ Zxdídî„ Zydïdð„ Zzdñdò„ Z{dódô„ Z|dõdö„ Z}d÷dø„ Z~dùdú„ Zdûdü„ Z€dýdþ„ Z�dÿ�d „ Z‚�d�d„ Zƒ�d�d„ Z„�d�d„ Z…�d�d„ Z†�d	�d
„ Z‡�d�d„ Zˆ�d�d„ Z‰�d�d„ ZŠ�d�d„ Z‹�d�d„ ZŒ�d�d„ Z��d�d„ ZŽ�d�d„ Z��d�d„ Z��d�d„ Z‘�d�d „ Z’�d!�d"„ Z“�d#�d$„ Z”�d%�d&„ Z•�d'�d(„ Z–‡  Z—S (+  ÚTestCudaAtomicsc                    s   t ƒ  ¡  tj d¡ d S )Nr   )ÚsuperÚsetUpÚnpÚrandomÚseed©Úself©Ú	__class__r   r   r²   ×  s    
zTestCudaAtomics.setUpc                 C   s¾   t jjdddd� t j¡}| ¡ }| ¡ }t d¡tƒ}|d |ƒ t d¡t	ƒ}|d |ƒ t j
dt jd�}t|jƒD ]}|||   d7  < qx|  t  ||k¡¡ |  t  ||k¡¡ d S ©Nr   r4   ©r¦   zvoid(uint32[:])©r*   r4   ©Údtyper*   )r³   r´   ÚrandintÚastyper   Úcopyr   Újitr9   r:   ÚzerosÚranger¦   Ú
assertTrueÚall)r·   r   Úary_wrapÚorigZcuda_atomic_addZcuda_atomic_add_wrapÚgoldÚir   r   r   Útest_atomic_addÛ  s    zTestCudaAtomics.test_atomic_addc                 C   sš   t jjdddd� t j¡ dd¡}| ¡ }| ¡ }t d¡t	ƒ}|d |ƒ t d¡t
ƒ}|d |ƒ |  t  ||d k¡¡ |  t  ||d k¡¡ d S ©	Nr   r4   r»   r=   r>   zvoid(uint32[:,:])©r*   r<   r*   )r³   r´   r¿   rÀ   r   ÚreshaperÁ   r   rÂ   r@   rA   rÅ   rÆ   )r·   r   rÇ   rÈ   Úcuda_atomic_add2Zcuda_atomic_add2_wrapr   r   r   Útest_atomic_add2í  s    "z TestCudaAtomics.test_atomic_add2c                 C   s`   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ |  
t  ||d k¡¡ d S rÌ   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rB   rÅ   rÆ   ©r·   r   rÈ   Zcuda_atomic_add3r   r   r   Útest_atomic_add3û  s
    "z TestCudaAtomics.test_atomic_add3c                 C   sÆ   t jjdddd� t j¡}| ¡ }| ¡  t j¡}t d¡t	ƒ}|d |ƒ t d¡t
ƒ}|d |ƒ t jdt jd�}t|jƒD ]}|||   d7  < q€|  t  ||k¡¡ |  t  ||k¡¡ d S ©Nr   r4   r»   zvoid(float32[:])r¼   r½   rD   )r³   r´   r¿   rÀ   r   rÁ   Úintpr   rÂ   rG   rI   rÃ   r   rÄ   r¦   rÅ   rÆ   )r·   r   rÇ   rÈ   Zcuda_atomic_add_floatZadd_float_wraprÉ   rÊ   r   r   r   Útest_atomic_add_float  s    z%TestCudaAtomics.test_atomic_add_floatc                 C   sš   t jjdddd� t j¡ dd¡}| ¡ }| ¡ }t d¡t	ƒ}|d |ƒ t d¡t
ƒ}|d |ƒ |  t  ||d k¡¡ |  t  ||d k¡¡ d S ©	Nr   r4   r»   r=   r>   zvoid(float32[:,:])rÍ   r*   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rL   rN   rÅ   rÆ   )r·   r   rÇ   rÈ   rÏ   Zcuda_func_wrapr   r   r   Útest_atomic_add_float_2  s    "z'TestCudaAtomics.test_atomic_add_float_2c                 C   s`   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ |  
t  ||d k¡¡ d S rÖ   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rO   rÅ   rÆ   rÑ   r   r   r   Útest_atomic_add_float_3#  s
    "z'TestCudaAtomics.test_atomic_add_float_3Tc                 C   sj   t jr
d S tt| ¡  ¡ ƒƒ}tddƒrH|r:|  d|¡ qf|  d|¡ n|rZ|  d|¡ n|  d|¡ d S )Né   r   zatom.shared.add.f64zatom.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ZENABLE_CUDASIMÚnextÚiterZinspect_asmÚvaluesr   ZassertIn)r·   Úkernelr   Úasmr   r   r   ÚassertCorrectFloat64Atomics+  s    
z+TestCudaAtomics.assertCorrectFloat64Atomicsc                 C   sÌ   t jjdddt jd�}t  dt j¡}| ¡ }t d¡t	ƒ}|d ||ƒ t d¡t
ƒ}|d ||ƒ t jdt jd�}t|jƒD ]}|||   d7  < q~t j ||¡ t j ||¡ |  |¡ |  |¡ d S ©Nr   r4   ©r¦   r¾   úvoid(int64[:], float64[:])r¼   r½   rD   )r³   r´   r¿   Úint64rÃ   r   rÁ   r   rÂ   r[   r\   r   rÄ   r¦   ÚtestingÚassert_equalrß   )r·   r   r   rÇ   Úcuda_fnZwrap_fnrÉ   rÊ   r   r   r   Útest_atomic_add_double<  s    
z&TestCudaAtomics.test_atomic_add_doublec                 C   s¢   t jjdddd� t j¡ dd¡}| ¡ }| ¡ }t d¡t	ƒ}|d |ƒ t d¡t
ƒ}|d |ƒ t j ||d ¡ t j ||d ¡ |  |¡ |  |¡ d S ©	Nr   r4   r»   r=   r>   úvoid(float64[:,:])rÍ   r*   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   r^   r_   rä   rå   rß   )r·   r   rÇ   rÈ   ræ   Zcuda_fn_wrapr   r   r   Útest_atomic_add_double_2Q  s    "
z(TestCudaAtomics.test_atomic_add_double_2c                 C   sd   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ t j
 ||d ¡ |  |¡ d S rè   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   r`   rä   rå   rß   ©r·   r   rÈ   Ú	cuda_funcr   r   r   Útest_atomic_add_double_3a  s    "z(TestCudaAtomics.test_atomic_add_double_3c           	      C   sØ   t jjdddt jd�}t  dt j¡}| ¡ }d}t |¡t	ƒ}t |¡t
ƒ}|d ||ƒ |d ||ƒ t jdt jd�}t|jƒD ]}|||   d7  < q‚t j ||¡ t j ||¡ | j|dd	� | j|dd	� d S )
Nr   r4   rá   râ   r¼   r½   rD   F©r   )r³   r´   r¿   rã   rÃ   r   rÁ   r   rÂ   rS   rT   r   rÄ   r¦   rä   rå   rß   )	r·   r   r   rÇ   Úsigrì   Úwrap_cuda_funcrÉ   rÊ   r   r   r   Útest_atomic_add_double_globalj  s    z-TestCudaAtomics.test_atomic_add_double_globalc                 C   s®   t jjdddd� t j¡ dd¡}| ¡ }| ¡ }d}t |¡t	ƒ}t |¡t
ƒ}|d |ƒ |d |ƒ t j ||d ¡ t j ||d ¡ | j|d	d
� | j|d	d
� d S ©Nr   r4   r»   r=   r>   ré   rÍ   r*   Frî   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rW   rX   rä   rå   rß   )r·   r   rÇ   rÈ   rï   rì   rð   r   r   r   Útest_atomic_add_double_global_2€  s    "z/TestCudaAtomics.test_atomic_add_double_global_2c                 C   sh   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ t j
 ||d ¡ | j|d	d
� d S rò   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rY   rä   rå   rß   rë   r   r   r   Útest_atomic_add_double_global_3‘  s    "z/TestCudaAtomics.test_atomic_add_double_global_3c                 C   sˆ   t jjdddd� t j¡}| ¡ }t d¡tƒ}|d |ƒ t j	dt jd�}t
|jƒD ]}|||   d8  < qV|  t  ||k¡¡ d S rº   )r³   r´   r¿   rÀ   r   rÁ   r   rÂ   rb   rÃ   rÄ   r¦   rÅ   rÆ   )r·   r   rÈ   Zcuda_atomic_subrÉ   rÊ   r   r   r   Útest_atomic_subš  s    zTestCudaAtomics.test_atomic_subc                 C   s`   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ |  
t  ||d k¡¡ d S rÌ   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rc   rÅ   rÆ   ©r·   r   rÈ   Zcuda_atomic_sub2r   r   r   Útest_atomic_sub2¦  s
    "z TestCudaAtomics.test_atomic_sub2c                 C   s`   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ |  
t  ||d k¡¡ d S rÌ   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rd   rÅ   rÆ   ©r·   r   rÈ   Zcuda_atomic_sub3r   r   r   Útest_atomic_sub3­  s
    "z TestCudaAtomics.test_atomic_sub3c                 C   s�   t jjdddd� t j¡}| ¡  t j¡}t d¡t	ƒ}|d |ƒ t j
dt jd�}t|jƒD ]}|||   d8  < q^|  t  ||k¡¡ d S rÓ   )r³   r´   r¿   rÀ   r   rÁ   rÔ   r   rÂ   re   rÃ   rÄ   r¦   rÅ   rÆ   )r·   r   rÈ   Zcuda_atomic_sub_floatrÉ   rÊ   r   r   r   Útest_atomic_sub_float´  s    z%TestCudaAtomics.test_atomic_sub_floatc                 C   s`   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ |  
t  ||d k¡¡ d S rÖ   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rf   rÅ   rÆ   rö   r   r   r   Útest_atomic_sub_float_2À  s
    "z'TestCudaAtomics.test_atomic_sub_float_2c                 C   s`   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ |  
t  ||d k¡¡ d S rÖ   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rg   rÅ   rÆ   rø   r   r   r   Útest_atomic_sub_float_3Ç  s
    "z'TestCudaAtomics.test_atomic_sub_float_3c                 C   s†   t jjdddt jd�}t  dt j¡}t d¡tƒ}|d ||ƒ t jdt jd�}t	|j
ƒD ]}|||   d8  < qZt j ||¡ d S rà   )r³   r´   r¿   rã   rÃ   r   r   rÂ   rh   rÄ   r¦   rä   rå   )r·   r   r   rì   rÉ   rÊ   r   r   r   Útest_atomic_sub_doubleÎ  s    z&TestCudaAtomics.test_atomic_sub_doublec                 C   sZ   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ t j
 ||d ¡ d S rè   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   ri   rä   rå   rë   r   r   r   Útest_atomic_sub_double_2Ú  s
    "z(TestCudaAtomics.test_atomic_sub_double_2c                 C   sZ   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ t j
 ||d ¡ d S rè   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rk   rä   rå   rë   r   r   r   Útest_atomic_sub_double_3á  s
    "z(TestCudaAtomics.test_atomic_sub_double_3c                 C   sŠ   t jjdddt jd�}t  dt j¡}d}t |¡tƒ}|d ||ƒ t jdt jd�}t	|j
ƒD ]}|||   d8  < q^t j ||¡ d S rà   )r³   r´   r¿   rã   rÃ   r   r   rÂ   rl   rÄ   r¦   rä   rå   )r·   r   r   rï   rì   rÉ   rÊ   r   r   r   Útest_atomic_sub_double_globalè  s    z-TestCudaAtomics.test_atomic_sub_double_globalc                 C   sZ   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ t j
 ||d ¡ d S rè   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rm   rä   rå   rë   r   r   r   Útest_atomic_sub_double_global_2õ  s
    "z/TestCudaAtomics.test_atomic_sub_double_global_2c                 C   sZ   t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d |ƒ t j
 ||d ¡ d S rè   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rn   rä   rå   rë   r   r   r   Útest_atomic_sub_double_global_3ü  s
    "z/TestCudaAtomics.test_atomic_sub_double_global_3c                 C   sŽ   t j d¡}t jjdddd� t j¡}| ¡ }t d¡tƒ}|d ||ƒ | ¡ }t	|j
ƒD ]}|||   |M  < q\|  t  ||k¡¡ d S )Néô  r   r4   r»   úvoid(uint32[:], uint32)r¼   )r³   r´   r¿   rÀ   r   rÁ   r   rÂ   rq   rÄ   r¦   rÅ   rÆ   ©r·   Ú
rand_constr   rÈ   rì   rÉ   rÊ   r   r   r   Útest_atomic_and  s    zTestCudaAtomics.test_atomic_andc                 C   sn   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ |  
t  |||@ k¡¡ d S ©	Nr  r   r4   r»   r=   r>   úvoid(uint32[:,:], uint32)rÍ   )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rs   rÅ   rÆ   ©r·   r  r   rÈ   Zcuda_atomic_and2r   r   r   Útest_atomic_and2  s    "z TestCudaAtomics.test_atomic_and2c                 C   sn   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ |  
t  |||@ k¡¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rt   rÅ   rÆ   ©r·   r  r   rÈ   Zcuda_atomic_and3r   r   r   Útest_atomic_and3  s    "z TestCudaAtomics.test_atomic_and3c                 C   s˜   t j d¡}t jjdddt jd�}t jjdddt jd�}d}t |¡tƒ}|d |||ƒ | ¡ }t|j	ƒD ]}|||   |M  < qlt j
 ||¡ d S ©Nr  r   r4   rá   zvoid(int32[:], int32[:], int32)r¼   )r³   r´   r¿   Úint32r   rÂ   rw   rÁ   rÄ   r¦   rä   rå   ©r·   r  r   r   rï   rì   rÉ   rÊ   r   r   r   Útest_atomic_and_global   s    z&TestCudaAtomics.test_atomic_and_globalc                 C   sh   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ t j
 |||@ ¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   ry   rä   rå   ©r·   r  r   rÈ   rì   r   r   r   Útest_atomic_and_global_2.  s    "z(TestCudaAtomics.test_atomic_and_global_2c                 C   s–   t j d¡}t jjdddd� t j¡}| ¡ }t d¡tƒ}|d ||ƒ t j	dt jd�}t
|jƒD ]}|||   |O  < qd|  t  ||k¡¡ d S ©Nr  r   r4   r»   r  r¼   r½   )r³   r´   r¿   rÀ   r   rÁ   r   rÂ   r|   rÃ   rÄ   r¦   rÅ   rÆ   r  r   r   r   Útest_atomic_or6  s    zTestCudaAtomics.test_atomic_orc                 C   sn   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ |  
t  |||B k¡¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   r}   rÅ   rÆ   r
  r   r   r   Útest_atomic_or2C  s    "zTestCudaAtomics.test_atomic_or2c                 C   sn   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ |  
t  |||B k¡¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   r~   rÅ   rÆ   r  r   r   r   Útest_atomic_or3K  s    "zTestCudaAtomics.test_atomic_or3c                 C   s˜   t j d¡}t jjdddt jd�}t jjdddt jd�}d}t |¡tƒ}|d |||ƒ | ¡ }t|j	ƒD ]}|||   |O  < qlt j
 ||¡ d S r  )r³   r´   r¿   r  r   rÂ   r   rÁ   rÄ   r¦   rä   rå   r  r   r   r   Útest_atomic_or_globalS  s    z%TestCudaAtomics.test_atomic_or_globalc                 C   sh   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ t j
 |||B ¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   r€   rä   rå   r  r   r   r   Útest_atomic_or_global_2a  s    "z'TestCudaAtomics.test_atomic_or_global_2c                 C   s–   t j d¡}t jjdddd� t j¡}| ¡ }t d¡tƒ}|d ||ƒ t j	dt jd�}t
|jƒD ]}|||   |N  < qd|  t  ||k¡¡ d S r  )r³   r´   r¿   rÀ   r   rÁ   r   rÂ   r‚   rÃ   rÄ   r¦   rÅ   rÆ   r  r   r   r   Útest_atomic_xori  s    zTestCudaAtomics.test_atomic_xorc                 C   sn   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ |  
t  |||A k¡¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   rƒ   rÅ   rÆ   )r·   r  r   rÈ   Zcuda_atomic_xor2r   r   r   Útest_atomic_xor2v  s    "z TestCudaAtomics.test_atomic_xor2c                 C   sn   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ |  
t  |||A k¡¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   r„   rÅ   rÆ   )r·   r  r   rÈ   Zcuda_atomic_xor3r   r   r   Útest_atomic_xor3~  s    "z TestCudaAtomics.test_atomic_xor3c                 C   s˜   t j d¡}t jjdddt jd�}t jjdddt jd�}| ¡ }d}t |¡tƒ}|d |||ƒ t|j	ƒD ]}|||   |N  < qlt j
 ||¡ d S r  )r³   r´   r¿   r  rÁ   r   rÂ   r…   rÄ   r¦   rä   rå   )r·   r  r   r   rÉ   rï   rì   rÊ   r   r   r   Útest_atomic_xor_global†  s    z&TestCudaAtomics.test_atomic_xor_globalc                 C   sh   t j d¡}t jjdddd� t j¡ dd¡}| ¡ }t d¡t	ƒ}|d ||ƒ t j
 |||A ¡ d S r  )r³   r´   r¿   rÀ   r   rÎ   rÁ   r   rÂ   r†   rä   rå   r  r   r   r   Útest_atomic_xor_global_2”  s    "z(TestCudaAtomics.test_atomic_xor_global_2c                 C   s@   t jjd|d�}t jjdddd� |¡}t jd|d�}|||fS )Nr4   r½   r   r»   )r³   r´   r¿   rÀ   Úarange)r·   r¾   ÚrconstÚraryZary_idxr   r   r   Úinc_dec_1dim_setupœ  s    z"TestCudaAtomics.inc_dec_1dim_setupc                 C   s8   t jjd|d�}t jjdddd� |¡ dd¡}||fS )Nr4   r½   r   r»   r=   r>   )r³   r´   r¿   rÀ   rÎ   )r·   r¾   r   r!  r   r   r   Úinc_dec_2dim_setup¢  s     z"TestCudaAtomics.inc_dec_2dim_setupc           
   	   C   sN   |  ¡ }t |¡|ƒ}	|	||f |||ƒ tj |t ||kd|d ¡¡ d S r)   ©rÁ   r   rÂ   r³   rä   rå   Úwhere©
r·   r   r   r   rï   ÚnblocksÚblksizerž   rÈ   rì   r   r   r   Úcheck_inc_index§  s    zTestCudaAtomics.check_inc_indexc           
   	   C   sN   |  ¡ }t |¡|ƒ}	|	||f |||ƒ tj |t ||kd|d ¡¡ d S r)   r$  r&  r   r   r   Úcheck_inc_index2­  s    z TestCudaAtomics.check_inc_index2c           	   	   C   sL   |  ¡ }t |¡|ƒ}|||f ||ƒ tj |t ||kd|d ¡¡ d S r)   r$  ©	r·   r   r   rï   r'  r(  rž   rÈ   rì   r   r   r   Ú	check_inc³  s    zTestCudaAtomics.check_incc              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S ©Nr½   ú"void(uint32[:], uint32[:], uint32)r*   r4   )r"  r³   r   r)  rŠ   ©r·   r  r   r   rï   r   r   r   Útest_atomic_inc_32¹  s    z"TestCudaAtomics.test_atomic_inc_32c              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S ©Nr½   z"void(uint64[:], uint64[:], uint64)r*   r4   )r"  r³   r   r)  r‹   r/  r   r   r   Útest_atomic_inc_64¾  s    z"TestCudaAtomics.test_atomic_inc_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S ©Nr	  r*   r<   )r#  r³   r   r,  rŒ   ©r·   r  r   rï   r   r   r   Útest_atomic_inc2_32Ã  s    z#TestCudaAtomics.test_atomic_inc2_32c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S ©Núvoid(uint64[:,:], uint64)r*   r<   )r#  r³   r   r,  r�   r4  r   r   r   Útest_atomic_inc2_64È  s    z#TestCudaAtomics.test_atomic_inc2_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r3  )r#  r³   r   r,  rŽ   r4  r   r   r   Útest_atomic_inc3Í  s    z TestCudaAtomics.test_atomic_inc3c              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S r-  )r"  r³   r   r*  r�   r/  r   r   r   Útest_atomic_inc_global_32Ò  s
    ÿz)TestCudaAtomics.test_atomic_inc_global_32c              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S r1  )r"  r³   r   r*  r�   r/  r   r   r   Útest_atomic_inc_global_64Ø  s
    ÿz)TestCudaAtomics.test_atomic_inc_global_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r3  )r#  r³   r   r,  r�   r4  r   r   r   Útest_atomic_inc_global_2_32Þ  s    z+TestCudaAtomics.test_atomic_inc_global_2_32c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r6  )r#  r³   r   r,  r�   r4  r   r   r   Útest_atomic_inc_global_2_64ã  s    z+TestCudaAtomics.test_atomic_inc_global_2_64c           
      C   s\   |  ¡ }t |¡|ƒ}	|	||f |||ƒ tj |t |dk|t ||k||d ¡¡¡ d S r)   r$  r&  r   r   r   Úcheck_dec_indexè  s    
þÿzTestCudaAtomics.check_dec_indexc           
      C   s\   |  ¡ }t |¡|ƒ}	|	||f |||ƒ tj |t |dk|t ||k||d ¡¡¡ d S r)   r$  r&  r   r   r   Úcheck_dec_index2ñ  s    
þÿz TestCudaAtomics.check_dec_index2c           	      C   sZ   |  ¡ }t |¡|ƒ}|||f ||ƒ tj |t |dk|t ||k||d ¡¡¡ d S r)   r$  r+  r   r   r   Ú	check_decú  s    
þÿzTestCudaAtomics.check_decc              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S r-  )r"  r³   r   r>  r’   r/  r   r   r   Útest_atomic_dec_32  s    z"TestCudaAtomics.test_atomic_dec_32c              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S r1  )r"  r³   r   r>  r“   r/  r   r   r   Útest_atomic_dec_64  s    z"TestCudaAtomics.test_atomic_dec_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r3  )r#  r³   r   r@  r”   r4  r   r   r   Útest_atomic_dec2_32  s    z#TestCudaAtomics.test_atomic_dec2_32c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r6  )r#  r³   r   r@  r•   r4  r   r   r   Útest_atomic_dec2_64  s    z#TestCudaAtomics.test_atomic_dec2_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r3  )r#  r³   r   r@  r–   r4  r   r   r   Útest_atomic_dec3_new  s    z$TestCudaAtomics.test_atomic_dec3_newc              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S r-  )r"  r³   r   r?  r—   r/  r   r   r   Útest_atomic_dec_global_32  s
    ÿz)TestCudaAtomics.test_atomic_dec_global_32c              	   C   s2   | j tjd�\}}}d}|  ||||ddt¡ d S r1  )r"  r³   r   r?  r—   r/  r   r   r   Útest_atomic_dec_global_64"  s
    ÿz)TestCudaAtomics.test_atomic_dec_global_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r3  )r#  r³   r   r@  r˜   r4  r   r   r   Útest_atomic_dec_global2_32(  s    z*TestCudaAtomics.test_atomic_dec_global2_32c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r6  )r#  r³   r   r@  r˜   r4  r   r   r   Útest_atomic_dec_global2_64-  s    z*TestCudaAtomics.test_atomic_dec_global2_64c                 C   sn   t jjddt jd�}t jjdddd� t j¡}t jdt jd�}t d¡tƒ}|d |||ƒ t j	 
||¡ d S )	Né2   éd   r½   r   r4   r»   r.  r¼   )r³   r´   r¿   r   rÀ   r  r   rÂ   rš   rä   rå   )r·   r  r   r   rì   r   r   r   Útest_atomic_exch2  s    z TestCudaAtomics.test_atomic_exchc                 C   sd   t jjddt jd�}t jjdddd� t j¡ dd¡}t d	¡tƒ}|d
 ||ƒ t j	 
||¡ d S )NrJ  rK  r½   r   r4   r»   r=   r>   r	  rÍ   )r³   r´   r¿   r   rÀ   rÎ   r   rÂ   r›   rä   rå   ©r·   r  r   rì   r   r   r   Útest_atomic_exch2<  s
    "z!TestCudaAtomics.test_atomic_exch2c                 C   sd   t jjddt jd�}t jjdddd� t j¡ dd¡}t d	¡tƒ}|d
 ||ƒ t j	 
||¡ d S )NrJ  rK  r½   r   r4   r»   r=   r>   r7  rÍ   )r³   r´   r¿   r   rÀ   rÎ   r   rÂ   rœ   rä   rå   rM  r   r   r   Útest_atomic_exch3D  s
    "z!TestCudaAtomics.test_atomic_exch3c                 C   sn   t jjddt jd�}t jdt jd�}t jjdddt jd�}d}t |¡tƒ}|d |||ƒ t j 	||¡ d S )	NrJ  rK  r½   r4   r   rá   r.  r¼   )
r³   r´   r¿   r   r  r   rÂ   r�   rä   rå   )r·   r  r   r   rï   rì   r   r   r   Útest_atomic_exch_globalL  s    z'TestCudaAtomics.test_atomic_exch_globalc                 C   s\   t jj||dd� |¡}t jd|jd�}t t¡}|d ||ƒ t  	|¡}t j
 ||¡ d S )N©r4   r4   r»   r*   r½   )r³   r´   r¿   rÀ   rÃ   r¾   r   rÂ   Ú
atomic_maxÚmaxrä   rå   ©r·   r¾   ÚloÚhiÚvalsr¨   rì   rÉ   r   r   r   Úcheck_atomic_maxV  s    

z TestCudaAtomics.check_atomic_maxc                 C   s   | j tjddd� d S ©Né ÿÿéÿÿ  ©r¾   rU  rV  )rX  r³   r  r¶   r   r   r   Útest_atomic_max_int32^  s    z%TestCudaAtomics.test_atomic_max_int32c                 C   s   | j tjddd� d S ©Nr   r[  r\  )rX  r³   r   r¶   r   r   r   Útest_atomic_max_uint32a  s    z&TestCudaAtomics.test_atomic_max_uint32c                 C   s   | j tjddd� d S rY  )rX  r³   rã   r¶   r   r   r   Útest_atomic_max_int64d  s    z%TestCudaAtomics.test_atomic_max_int64c                 C   s   | j tjddd� d S r^  )rX  r³   r   r¶   r   r   r   Útest_atomic_max_uint64g  s    z&TestCudaAtomics.test_atomic_max_uint64c                 C   s   | j tjddd� d S rY  )rX  r³   r   r¶   r   r   r   Útest_atomic_max_float32j  s    z'TestCudaAtomics.test_atomic_max_float32c                 C   s   | j tjddd� d S rY  )rX  r³   r   r¶   r   r   r   Útest_atomic_max_doublem  s    z&TestCudaAtomics.test_atomic_max_doublec                 C   s`   t jjdddd� t j¡}t  dt j¡}t d¡tƒ}|d ||ƒ t  	|¡}t j
 ||¡ d S ©Nr   r[  rQ  r»   r*   úvoid(float64[:], float64[:,:]))r³   r´   r¿   rÀ   r   rÃ   r   rÂ   Ú!atomic_max_double_normalizedindexrS  rä   rå   ©r·   rW  r¨   rì   rÉ   r   r   r   Ú&test_atomic_max_double_normalizedindexp  s    ÿ
z6TestCudaAtomics.test_atomic_max_double_normalizedindexc                 C   s`   t jjdddd� t j¡}t  dt j¡}t d¡tƒ}|d ||ƒ t  	|¡}t j
 ||¡ d S ©Nr   é€   r4   r»   r*   úvoid(float64[:], float64[:])r¼   )r³   r´   r¿   rÀ   r   rÃ   r   rÂ   Úatomic_max_double_oneindexrS  rä   rå   rg  r   r   r   Útest_atomic_max_double_oneindexz  s    ÿ
z/TestCudaAtomics.test_atomic_max_double_oneindexc                 C   s^   t jj||dd� |¡}t jdg|jd�}t t¡}|d ||ƒ t  	|¡}t j
 ||¡ d S )NrQ  r»   r[  r½   )r³   r´   r¿   rÀ   r   r¾   r   rÂ   Ú
atomic_minÚminrä   rå   rT  r   r   r   Úcheck_atomic_min„  s    

z TestCudaAtomics.check_atomic_minc                 C   s   | j tjddd� d S rY  )rp  r³   r  r¶   r   r   r   Útest_atomic_min_int32�  s    z%TestCudaAtomics.test_atomic_min_int32c                 C   s   | j tjddd� d S r^  )rp  r³   r   r¶   r   r   r   Útest_atomic_min_uint32�  s    z&TestCudaAtomics.test_atomic_min_uint32c                 C   s   | j tjddd� d S rY  )rp  r³   rã   r¶   r   r   r   Útest_atomic_min_int64“  s    z%TestCudaAtomics.test_atomic_min_int64c                 C   s   | j tjddd� d S r^  )rp  r³   r   r¶   r   r   r   Útest_atomic_min_uint64–  s    z&TestCudaAtomics.test_atomic_min_uint64c                 C   s   | j tjddd� d S rY  )rp  r³   r   r¶   r   r   r   Útest_atomic_min_float™  s    z%TestCudaAtomics.test_atomic_min_floatc                 C   s   | j tjddd� d S rY  )rp  r³   r   r¶   r   r   r   Útest_atomic_min_doubleœ  s    z&TestCudaAtomics.test_atomic_min_doublec                 C   sd   t jjdddd� t j¡}t  dt j¡d }t d¡tƒ}|d ||ƒ t  	|¡}t j
 ||¡ d S rd  )r³   r´   r¿   rÀ   r   Úonesr   rÂ   Ú!atomic_min_double_normalizedindexro  rä   rå   rg  r   r   r   Ú&test_atomic_min_double_normalizedindexŸ  s    ÿ
z6TestCudaAtomics.test_atomic_min_double_normalizedindexc                 C   sd   t jjdddd� t j¡}t  dt j¡d }t d¡tƒ}|d ||ƒ t  	|¡}t j
 ||¡ d S ri  )r³   r´   r¿   rÀ   r   rw  r   rÂ   Úatomic_min_double_oneindexro  rä   rå   rg  r   r   r   Útest_atomic_min_double_oneindex©  s    ÿ
z/TestCudaAtomics.test_atomic_min_double_oneindexc                 C   s`   t  d¡|ƒ}tjjdddd� tj¡}t dtj¡tj }|d ||ƒ tj	 
|tjg¡ d S )Nre  r   rj  ©r*   r*   r»   r*   )r   rÂ   r³   r´   r¿   rÀ   r   rÃ   Únanrä   rå   )r·   rž   rì   rW  r¨   r   r   r   Ú _test_atomic_minmax_nan_location¾  s
    z0TestCudaAtomics._test_atomic_minmax_nan_locationc                 C   sd   t  d¡|ƒ}tjjdddd� tj¡}| ¡ }t dtj¡tj	 }|d ||ƒ tj
 ||¡ d S )Nre  r   rj  r*   r»   r|  )r   rÂ   r³   r´   r¿   rÀ   r   rÁ   rÃ   r}  rä   rå   )r·   rž   rì   r¨   rÉ   rW  r   r   r   Ú_test_atomic_minmax_nan_valÇ  s    z+TestCudaAtomics._test_atomic_minmax_nan_valc                 C   s   |   t¡ d S r   )r~  rn  r¶   r   r   r   Útest_atomic_min_nan_locationÑ  s    z,TestCudaAtomics.test_atomic_min_nan_locationc                 C   s   |   t¡ d S r   )r~  rR  r¶   r   r   r   Útest_atomic_max_nan_locationÔ  s    z,TestCudaAtomics.test_atomic_max_nan_locationc                 C   s   |   t¡ d S r   )r  rn  r¶   r   r   r   Útest_atomic_min_nan_val×  s    z'TestCudaAtomics.test_atomic_min_nan_valc                 C   s   |   t¡ d S r   )r  rR  r¶   r   r   r   Útest_atomic_max_nan_valÚ  s    z'TestCudaAtomics.test_atomic_max_nan_valc                 C   sd   t jjdddd� t j¡}t  dt j¡}d}t |¡tƒ}|d ||ƒ t  	|¡}t j
 ||¡ d S ©Nr   r4   r»   r*   rk  r¼   )r³   r´   r¿   rÀ   r   rÃ   r   rÂ   Úatomic_max_double_sharedrS  rä   rå   ©r·   rW  r¨   rï   rì   rÉ   r   r   r   Útest_atomic_max_double_sharedÝ  s    
z-TestCudaAtomics.test_atomic_max_double_sharedc                 C   sh   t jjdddd� t j¡}t  dt j¡d }d}t |¡tƒ}|d ||ƒ t  	|¡}t j
 ||¡ d S r„  )r³   r´   r¿   rÀ   r   rw  r   rÂ   Úatomic_min_double_sharedro  rä   rå   r†  r   r   r   Útest_atomic_min_double_sharedç  s    
z-TestCudaAtomics.test_atomic_min_double_sharedr*   c                 C   sø   |g|d  |g|d   }t j |¡ t j||d�}|dkrDd|_t  |¡}t jjdd|jd� |j¡}	||k}
||k}t  |¡}|	|
 ||
< |||< | 	¡ }t
 |¡}|dkrÆ|d |||	|ƒ n|d |||	|ƒ t j ||¡ t j ||¡ d S )	Nr®   r½   )é
   éÿÿÿÿr*   rŠ  r»   ©rŠ  rŠ  )rŒ  rŒ  )r³   r´   ÚshuffleZasarrayr0   Z
zeros_liker¿   rÀ   r¾   rÁ   r   rÂ   rä   Zassert_array_equal)r·   ÚnÚfillÚunfillr¾   Úcas_funcÚndimr¨   Úoutr   Z	fill_maskZunfill_maskZ
expect_resZ
expect_outrì   r   r   r   Ú	check_casñ  s&    


zTestCudaAtomics.check_casc                 C   s   | j dddtjtd� d S ©NrK  é�ÿÿÿr‹  ©rŽ  r�  r�  r¾   r‘  )r”  r³   r  r«   r¶   r   r   r   Útest_atomic_compare_and_swap  s    ÿz,TestCudaAtomics.test_atomic_compare_and_swapc                 C   s   | j dddtjtd� d S ©NrK  éÓÿÿÿr‹  r—  )r”  r³   rã   r«   r¶   r   r   r   Útest_atomic_compare_and_swap2  s    ÿz-TestCudaAtomics.test_atomic_compare_and_swap2c                 C   sB   t jjddt jd�}t jjddt jd�}| jd||t jtd� d S ©NrJ  r  r½   r*   é   rK  r—  )r³   r´   r¿   r   r”  r«   ©r·   ZrfillZrunfillr   r   r   Útest_atomic_compare_and_swap3  s
    ÿz-TestCudaAtomics.test_atomic_compare_and_swap3c                 C   sB   t jjddt jd�}t jjddt jd�}| jd||t jtd� d S rœ  )r³   r´   r¿   r   r”  r«   rž  r   r   r   Útest_atomic_compare_and_swap4  s
    ÿz-TestCudaAtomics.test_atomic_compare_and_swap4c                 C   s   | j dddtjtd� d S r•  )r”  r³   r  r­   r¶   r   r   r   Útest_atomic_cas_1dim   s    ÿz$TestCudaAtomics.test_atomic_cas_1dimc                 C   s   | j dddtjtdd� d S )NrK  r–  r‹  r®   ©rŽ  r�  r�  r¾   r‘  r’  )r”  r³   r  r¯   r¶   r   r   r   Útest_atomic_cas_2dim$  s     ÿz$TestCudaAtomics.test_atomic_cas_2dimc                 C   s   | j dddtjtd� d S r™  )r”  r³   rã   r­   r¶   r   r   r   Útest_atomic_cas2_1dim(  s    ÿz%TestCudaAtomics.test_atomic_cas2_1dimc                 C   s   | j dddtjtdd� d S )NrK  rš  r‹  r®   r¢  )r”  r³   rã   r¯   r¶   r   r   r   Útest_atomic_cas2_2dim,  s     ÿz%TestCudaAtomics.test_atomic_cas2_2dimc                 C   sB   t jjddt jd�}t jjddt jd�}| jd||t jtd� d S rœ  )r³   r´   r¿   r   r”  r­   rž  r   r   r   Útest_atomic_cas3_1dim0  s
    ÿz%TestCudaAtomics.test_atomic_cas3_1dimc                 C   sD   t jjddt jd�}t jjddt jd�}| jd||t jtdd� d S ©	NrJ  r  r½   r*   r�  rK  r®   r¢  )r³   r´   r¿   r   r”  r¯   rž  r   r   r   Útest_atomic_cas3_2dim6  s     ÿz%TestCudaAtomics.test_atomic_cas3_2dimc                 C   sB   t jjddt jd�}t jjddt jd�}| jd||t jtd� d S rœ  )r³   r´   r¿   r   r”  r­   rž  r   r   r   Útest_atomic_cas4_1dim<  s
    ÿz%TestCudaAtomics.test_atomic_cas4_1dimc                 C   sD   t jjddt jd�}t jjddt jd�}| jd||t jtdd� d S r§  )r³   r´   r¿   r   r”  r¯   rž  r   r   r   Útest_atomic_cas4_2dimB  s     ÿz%TestCudaAtomics.test_atomic_cas4_2dimc                 C   sX   t jdt jd�}||d< |d |ƒ t  |¡rD|  t  |d ¡¡ n|  |d |¡ d S )Nr®   r½   r   r|  r*   )r³   rÃ   r   ÚisnanrÅ   ÚassertEqual©r·   rÝ   Úinitialr   r   r   r   Ú_test_atomic_returns_oldM  s    
z(TestCudaAtomics._test_atomic_returns_oldc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S r)   )r   r6   r7   ©r   r   r   r   rÝ   W  s    z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernelrŠ  ©r   rÂ   r¯  ©r·   rÝ   r   r   r   Útest_atomic_add_returns_oldV  s    
z+TestCudaAtomics.test_atomic_add_returns_oldc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S r)   ©r   r6   rS  r°  r   r   r   rÝ   ^  s    zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelrŠ  r±  r²  r   r   r   Ú"test_atomic_max_returns_no_replace]  s    
z2TestCudaAtomics.test_atomic_max_returns_no_replacec                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S ©Nr   rŠ  r*   r´  r°  r   r   r   rÝ   e  s    zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelr*   r±  r²  r   r   r   Ú#test_atomic_max_returns_old_replaced  s    
z3TestCudaAtomics.test_atomic_max_returns_old_replacec                 C   s    t jdd„ ƒ}|  |tj¡ d S )Nc                 S   s   t j | dd¡| d< d S r)   r´  r°  r   r   r   rÝ   l  s    zHTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernel©r   rÂ   r¯  r³   r}  r²  r   r   r   Ú(test_atomic_max_returns_old_nan_in_arrayk  s    
z8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dtj¡| d< d S r)   )r   r6   rS  r³   r}  r°  r   r   r   rÝ   s  s    zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelrŠ  r±  r²  r   r   r   Ú#test_atomic_max_returns_old_nan_valr  s    
z3TestCudaAtomics.test_atomic_max_returns_old_nan_valc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S ©Nr   é   r*   ©r   r6   ro  r°  r   r   r   rÝ   z  s    zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernelrŠ  r±  r²  r   r   r   Ú&test_atomic_min_returns_old_no_replacey  s    
z6TestCudaAtomics.test_atomic_min_returns_old_no_replacec                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S r¶  r½  r°  r   r   r   rÝ   �  s    zCTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernelr¼  r±  r²  r   r   r   Ú#test_atomic_min_returns_old_replace€  s    
z3TestCudaAtomics.test_atomic_min_returns_old_replacec                 C   s    t jdd„ ƒ}|  |tj¡ d S )Nc                 S   s   t j | dd¡| d< d S r»  r½  r°  r   r   r   rÝ   ˆ  s    zHTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernelr¸  r²  r   r   r   Ú(test_atomic_min_returns_old_nan_in_array‡  s    
z8TestCudaAtomics.test_atomic_min_returns_old_nan_in_arrayc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dtj¡| d< d S r)   )r   r6   ro  r³   r}  r°  r   r   r   rÝ   �  s    zCTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernelr¼  r±  r²  r   r   r   Ú#test_atomic_min_returns_old_nan_valŽ  s    
z3TestCudaAtomics.test_atomic_min_returns_old_nan_valc           	      C   sj   t jj||dd� |¡}||dd d…< t jd|jd�}t t¡}|d ||ƒ t  	|¡}t j
 ||¡ d S )NrQ  r»   r*   r®   r½   )r³   r´   r¿   rÀ   rÃ   r¾   r   rÂ   Úatomic_nanmaxÚnanmaxrä   rå   ©	r·   r¾   rU  rV  Úinit_valrW  r¨   rì   rÉ   r   r   r   Úcheck_atomic_nanmax˜  s    

z#TestCudaAtomics.check_atomic_nanmaxc                 C   s   | j tjdddd� d S ©NrZ  r[  r   ©r¾   rU  rV  rÅ  )rÆ  r³   r  r¶   r   r   r   Útest_atomic_nanmax_int32¡  s    ÿz(TestCudaAtomics.test_atomic_nanmax_int32c                 C   s   | j tjdddd� d S ©Nr   r[  rÈ  )rÆ  r³   r   r¶   r   r   r   Útest_atomic_nanmax_uint32¥  s    ÿz)TestCudaAtomics.test_atomic_nanmax_uint32c                 C   s   | j tjdddd� d S rÇ  )rÆ  r³   rã   r¶   r   r   r   Útest_atomic_nanmax_int64©  s    ÿz(TestCudaAtomics.test_atomic_nanmax_int64c                 C   s   | j tjdddd� d S rÊ  )rÆ  r³   r   r¶   r   r   r   Útest_atomic_nanmax_uint64­  s    ÿz)TestCudaAtomics.test_atomic_nanmax_uint64c                 C   s   | j tjddtjd� d S ©NrZ  r[  rÈ  )rÆ  r³   r   r}  r¶   r   r   r   Útest_atomic_nanmax_float32±  s    ÿz*TestCudaAtomics.test_atomic_nanmax_float32c                 C   s   | j tjddtjd� d S rÎ  )rÆ  r³   r   r}  r¶   r   r   r   Útest_atomic_nanmax_doubleµ  s    ÿz)TestCudaAtomics.test_atomic_nanmax_doublec                 C   sx   t jjdddd� t j¡}t j|dd d…< t jdg|jd�}d}t 	|¡t
ƒ}|d ||ƒ t  |¡}t j ||¡ d S ©	Nr   r4   r»   r*   r®   r½   rk  r¼   )r³   r´   r¿   rÀ   r   r}  r   r¾   r   rÂ   Úatomic_nanmax_double_sharedrÃ  rä   rå   r†  r   r   r   Ú test_atomic_nanmax_double_shared¹  s    
z0TestCudaAtomics.test_atomic_nanmax_double_sharedc                 C   sp   t jjdddd� t j¡}t j|dd d…< t  dt j¡}t d¡t	ƒ}|d ||ƒ t  
|¡}t j ||¡ d S ©	Nr   rj  r4   r»   r*   r®   rk  r¼   )r³   r´   r¿   rÀ   r   r}  rÃ   r   rÂ   rl  rÃ  rä   rå   rg  r   r   r   Ú"test_atomic_nanmax_double_oneindexÄ  s    ÿ
z2TestCudaAtomics.test_atomic_nanmax_double_oneindexc           	      C   sl   t jj||dd� |¡}||dd d…< t jdg|jd�}t t¡}|d ||ƒ t  	|¡}t j
 ||¡ d S )NrQ  r»   r*   r®   r[  r½   )r³   r´   r¿   rÀ   r   r¾   r   rÂ   Úatomic_nanminÚnanminrä   rå   rÄ  r   r   r   Úcheck_atomic_nanminÐ  s    

z#TestCudaAtomics.check_atomic_nanminc                 C   s   | j tjdddd� d S rÇ  )rØ  r³   r  r¶   r   r   r   Útest_atomic_nanmin_int32Ú  s    ÿz(TestCudaAtomics.test_atomic_nanmin_int32c                 C   s   | j tjdddd� d S rÊ  )rØ  r³   r   r¶   r   r   r   Útest_atomic_nanmin_uint32Þ  s    ÿz)TestCudaAtomics.test_atomic_nanmin_uint32c                 C   s   | j tjdddd� d S rÇ  )rØ  r³   rã   r¶   r   r   r   Útest_atomic_nanmin_int64â  s    ÿz(TestCudaAtomics.test_atomic_nanmin_int64c                 C   s   | j tjdddd� d S rÊ  )rØ  r³   r   r¶   r   r   r   Útest_atomic_nanmin_uint64æ  s    ÿz)TestCudaAtomics.test_atomic_nanmin_uint64c                 C   s   | j tjddtjd� d S rÎ  )rØ  r³   r   r}  r¶   r   r   r   Útest_atomic_nanmin_floatê  s    ÿz(TestCudaAtomics.test_atomic_nanmin_floatc                 C   s   | j tjddtjd� d S rÎ  )rØ  r³   r   r}  r¶   r   r   r   Útest_atomic_nanmin_doubleî  s    ÿz)TestCudaAtomics.test_atomic_nanmin_doublec                 C   sx   t jjdddd� t j¡}t j|dd d…< t jdg|jd�}d}t 	|¡t
ƒ}|d ||ƒ t  |¡}t j ||¡ d S rÑ  )r³   r´   r¿   rÀ   r   r}  r   r¾   r   rÂ   Úatomic_nanmin_double_sharedr×  rä   rå   r†  r   r   r   Ú test_atomic_nanmin_double_sharedò  s    
z0TestCudaAtomics.test_atomic_nanmin_double_sharedc                 C   sr   t jjdddd� t j¡}t j|dd d…< t  dgt j¡}t d¡t	ƒ}|d ||ƒ t  
|¡}t j ||¡ d S rÔ  )r³   r´   r¿   rÀ   r   r}  r   r   rÂ   rz  r×  rä   rå   rg  r   r   r   Ú"test_atomic_nanmin_double_oneindexý  s    ÿ
z2TestCudaAtomics.test_atomic_nanmin_double_oneindexc                 C   sv   t jdt jd�}||d< t j|d< |d |ƒ t  |¡rb|  t  |d ¡¡ |  t  |d ¡¡ n|  |d |¡ d S )Nr®   r½   r   r*   r|  )r³   rÃ   r   r}  r«  ZassertFalserÅ   r¬  r­  r   r   r   Ú_test_atomic_nan_returns_old
  s    

z,TestCudaAtomics._test_atomic_nan_returns_oldc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S r)   ©r   r6   rÃ  r°  r   r   r   rÝ     s    zITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernelrŠ  ©r   rÂ   râ  r²  r   r   r   Ú)test_atomic_nanmax_returns_old_no_replace  s    
z9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replacec                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S r¶  rã  r°  r   r   r   rÝ     s    zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernelr*   rä  r²  r   r   r   Ú&test_atomic_nanmax_returns_old_replace  s    
z6TestCudaAtomics.test_atomic_nanmax_returns_old_replacec                 C   s    t jdd„ ƒ}|  |tj¡ d S )Nc                 S   s   t j | dd¡| d< d S r)   rã  r°  r   r   r   rÝ   $  s    zKTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernel©r   rÂ   râ  r³   r}  r²  r   r   r   Ú+test_atomic_nanmax_returns_old_nan_in_array#  s    
z;TestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_arrayc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dtj¡| d< d S r)   )r   r6   rÃ  r³   r}  r°  r   r   r   rÝ   +  s    zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernelrŠ  rä  r²  r   r   r   Ú&test_atomic_nanmax_returns_old_nan_val*  s    
z6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_valc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S r»  ©r   r6   r×  r°  r   r   r   rÝ   2  s    zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernelrŠ  rä  r²  r   r   r   Ú)test_atomic_nanmin_returns_old_no_replace1  s    
z9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replacec                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dd¡| d< d S r¶  rê  r°  r   r   r   rÝ   9  s    zFTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernelr¼  rä  r²  r   r   r   Ú&test_atomic_nanmin_returns_old_replace8  s    
z6TestCudaAtomics.test_atomic_nanmin_returns_old_replacec                 C   s    t jdd„ ƒ}|  |tj¡ d S )Nc                 S   s   t j | dd¡| d< d S r»  rê  r°  r   r   r   rÝ   @  s    zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelrç  r²  r   r   r   Ú+test_atomic_nanmin_returns_old_nan_in_array?  s    
z;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayc                 C   s   t jdd„ ƒ}|  |d¡ d S )Nc                 S   s   t j | dtj¡| d< d S r)   )r   r6   r×  r³   r}  r°  r   r   r   rÝ   G  s    zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelr¼  rä  r²  r   r   r   Ú&test_atomic_nanmin_returns_old_nan_valF  s    
z6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_val)T)r*   )˜Ú__name__Ú
__module__Ú__qualname__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*  r,  r0  r2  r5  r8  r9  r:  r;  r<  r=  r>  r?  r@  rA  rB  rC  rD  rE  rF  rG  rH  rI  rL  rN  rO  rP  rX  r]  r_  r`  ra  rb  rc  rh  rm  rp  rq  rr  rs  rt  ru  rv  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Ü  rÝ  rÞ  rà  rá  râ  rå  ræ  rè  ré  rë  rì  rí  rî  Ú__classcell__r   r   r¸   r   r°   Ö  s*  
	
				



	
	


	
	

r°   Ú__main__)rÚnumpyr³   Útextwrapr   Znumbar   r   r   r   r   Znumba.cuda.testingr   r	   r
   r   Z
numba.corer   rÂ   r   r   r   r'   r(   r/   r1   r2   r9   r:   r@   rA   rB   rG   rI   rL   rN   rO   rS   rT   rW   rX   rY   r[   r\   r^   r_   r`   rb   rc   rd   re   rf   rg   rh   ri   rk   rl   rm   rn   rq   rs   rt   rw   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£   rR  rf  rl  r…  rn  rx  rz  rˆ  rÂ  Z$atomic_nanmax_double_normalizedindexZatomic_nanmax_double_oneindexrÒ  rÖ  Z$atomic_nanmin_double_normalizedindexZatomic_nanmin_double_oneindexrß  r«   r­   r¯   r°   rï  Úmainr   r   r   r   Ú<module>   sæ   













	

	%ÿÿþ þ          
