U
    b»{j�ã  ã                   @   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   úc/var/www/html/bwave/python/venv/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ž   Z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�   r3  r   r   r   Útest_atomic_inc2_64È  s    z#TestCudaAtomics.test_atomic_inc2_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r2  )r"  r²   r   r+  rŽ   r3  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 r0  )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 r2  )r"  r²   r   r+  r�   r3  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 r5  )r"  r²   r   r+  r�   r3  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 r0  )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 r2  )r"  r²   r   r?  r”   r3  r   r   r   Útest_atomic_dec2_32  s    z#TestCudaAtomics.test_atomic_dec2_32c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r5  )r"  r²   r   r?  r•   r3  r   r   r   Útest_atomic_dec2_64  s    z#TestCudaAtomics.test_atomic_dec2_64c                 C   s,   |   tj¡\}}d}|  |||ddt¡ d S r2  )r"  r²   r   r?  r–   r3  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 r0  )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 r2  )r"  r²   r   r?  r˜   r3  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 r5  )r"  r²   r   r?  r˜   r3  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 )NrI  rJ  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 )NrI  rJ  r¼   r   r4   rº   r=   r>   r6  rÌ   )r²   r³   r¾   r   r¿   rÍ   r   rÁ   rœ   rã   rä   rL  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 )	NrI  rJ  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½   rT  rU  )rW  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   rZ  r[  )rW  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 rX  )rW  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]  )rW  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 rX  )rW  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 rX  )rW  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   rZ  rP  rº   r*   úvoid(float64[:], float64[:,:]))r²   r³   r¾   r¿   r   rÂ   r   rÁ   Ú!atomic_max_double_normalizedindexrR  rã   rä   ©r¶   rV  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_oneindexrR  rã   rä   rf  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 )NrP  rº   rZ  r¼   )r²   r³   r¾   r¿   r   r½   r   rÁ   Ú
atomic_minÚminrã   rä   rS  r   r   r   Úcheck_atomic_min„  s    

z TestCudaAtomics.check_atomic_minc                 C   s   | j tjddd� d S rX  )ro  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]  )ro  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 rX  )ro  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]  )ro  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 rX  )ro  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 rX  )ro  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 rc  )r²   r³   r¾   r¿   r   Úonesr   rÁ   Ú!atomic_min_double_normalizedindexrn  rã   rä   rf  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 rh  )r²   r³   r¾   r¿   r   rv  r   rÁ   Úatomic_min_double_oneindexrn  rã   rä   rf  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 )Nrd  r   ri  ©r*   r*   rº   r*   )r   rÁ   r²   r³   r¾   r¿   r   rÂ   Únanrã   rä   )r¶   rž   rë   rV  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 )Nrd  r   ri  r*   rº   r{  )r   rÁ   r²   r³   r¾   r¿   r   rÀ   rÂ   r|  rã   rä   )r¶   rž   rë   r§   rÈ   rV  r   r   r   Ú_test_atomic_minmax_nan_valÇ  s    z+TestCudaAtomics._test_atomic_minmax_nan_valc                 C   s   |   t¡ d S r   )r}  rm  rµ   r   r   r   Útest_atomic_min_nan_locationÑ  s    z,TestCudaAtomics.test_atomic_min_nan_locationc                 C   s   |   t¡ d S r   )r}  rQ  rµ   r   r   r   Útest_atomic_max_nan_locationÔ  s    z,TestCudaAtomics.test_atomic_max_nan_locationc                 C   s   |   t¡ d S r   )r~  rm  rµ   r   r   r   Útest_atomic_min_nan_val×  s    z'TestCudaAtomics.test_atomic_min_nan_valc                 C   s   |   t¡ d S r   )r~  rQ  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*   rj  r»   )r²   r³   r¾   r¿   r   rÂ   r   rÁ   Úatomic_max_double_sharedrR  rã   rä   ©r¶   rV  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   rv  r   rÁ   Úatomic_min_double_sharedrn  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 ©NrJ  é�ÿÿÿ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 ©NrJ  éÓÿÿÿ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 ©NrI  r  r¼   r*   é   rJ  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 )NrJ  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 )NrJ  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 ©	NrI  r  r¼   r*   rœ  rJ  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   rR  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   rR  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   rn  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   rn  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 )NrP  rº   r*   r­   r¼   )r²   r³   r¾   r¿   rÂ   r½   r   rÁ   Úatomic_nanmaxÚnanmaxrã   rä   ©	r¶   r½   rT  rU  Úinit_valrV  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 ©NrY  rZ  r   ©r½   rT  rU  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   rZ  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 ©NrY  rZ  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¼   rj  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   ri  r4   rº   r*   r­   rj  r»   )r²   r³   r¾   r¿   r   r|  rÂ   r   rÁ   rk  rÂ  rã   rä   rf  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 )NrP  rº   r*   r­   rZ  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Á   ry  rÖ  rã   rä   rf  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+  r/  r1  r4  r7  r8  r9  r:  r;  r<  r=  r>  r?  r@  rA  rB  rC  rD  rE  rF  rG  rH  rK  rM  rN  rO  rW  r\  r^  r_  r`  ra  rb  rg  rl  ro  rp  rq  rr  rs  rt  ru  rx  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Ü  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¢   rQ  re  rk  r„  rm  rw  ry  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æ   













	

	%ÿÿþ þ          
