U
    hâËdm0  ã                   @   sP  d Z ddlmZ ddlZddlZddlZddlmZ ddlZ	ddl
mZ ddlmZ G dd	„ d	eƒZG d
d„ dƒZG dd„ dƒZG dd„ deƒZG dd„ deƒZG dd„ deƒZe ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Ze ¡ Z G dd„ deƒZ!G dd„ deƒZ"G dd„ deƒZ#edd„ ƒZ$dS )zf
Implements the cuda module as called from within an executing kernel
(@cuda.jit-decorated function).
é    )ÚcontextmanagerN)Útypes)Únumpy_supporté   )Úvector_typesc                   @   s0   e Zd ZdZdd„ Zdd„ Zdd„ Zdd	„ Zd
S )ÚDim3z;
    Used to implement thread/block indices/dimensions
    c                 C   s   || _ || _|| _d S ©N©ÚxÚyÚz)Úselfr
   r   r   © r   úW/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/simulator/kernelapi.pyÚ__init__   s    zDim3.__init__c                 C   s   d| j | j| jf S )Nz(%s, %s, %s)r	   ©r   r   r   r   Ú__str__   s    zDim3.__str__c                 C   s   d| j | j| jf S )NzDim3(%s, %s, %s)r	   r   r   r   r   Ú__repr__   s    zDim3.__repr__c                 c   s   | j V  | jV  | jV  d S r   r	   r   r   r   r   Ú__iter__!   s    zDim3.__iter__N)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   r   r   r   r   r   r   r   r      s
   r   c                   @   s   e Zd ZdZdd„ ZdS )Ú	GridGroupz+
    Used to implement the grid group.
    c                 C   s   t  ¡  ¡  d S r   ©Ú	threadingÚcurrent_threadÚsyncthreadsr   r   r   r   Úsync,   s    zGridGroup.syncN)r   r   r   r   r   r   r   r   r   r   '   s   r   c                   @   s   e Zd ZdZdd„ ZdS )Ú
FakeCUDACgz!
    CUDA Cooperative Groups
    c                 C   s   t ƒ S r   )r   r   r   r   r   Ú	this_grid7   s    zFakeCUDACg.this_gridN)r   r   r   r   r    r   r   r   r   r   3   s   r   c                   @   s   e Zd ZdZdd„ ZdS )ÚFakeCUDALocalz
    CUDA Local arrays
    c                 C   s"   t |tjƒrt |¡}t ||¡S r   )Ú
isinstancer   ÚTyper   Úas_dtypeÚnpÚempty)r   ÚshapeÚdtyper   r   r   Úarray?   s    
zFakeCUDALocal.arrayN)r   r   r   r   r)   r   r   r   r   r!   ;   s   r!   c                   @   s   e Zd ZdZdd„ ZdS )ÚFakeCUDAConstz
    CUDA Const arrays
    c                 C   s   |S r   r   )r   Zaryr   r   r   Ú
array_likeI   s    zFakeCUDAConst.array_likeN)r   r   r   r   r+   r   r   r   r   r*   E   s   r*   c                   @   s    e Zd ZdZdd„ Zdd„ ZdS )ÚFakeCUDASharedaæ  
    CUDA Shared arrays.

    Limitations: assumes that only one call to cuda.shared.array is on a line,
    and that that line is only executed once per thread. i.e.::

        a = cuda.shared.array(...); b = cuda.shared.array(...)

    will erroneously alias a and b, and::

        for i in range(10):
            sharedarrs[i] = cuda.shared.array(...)

    will alias all arrays created at that point (though it is not certain that
    this would be supported by Numba anyway).
    c                 C   s"   i | _ || _tj|tjd�| _d S ©N)r(   )Ú_allocationsÚ_dynshared_sizer%   ZzerosÚbyteÚ
_dynshared)r   Údynshared_sizer   r   r   r   _   s    zFakeCUDAShared.__init__c                 C   sŠ   t |tjƒrt |¡}|dkr>| j|j }tj| j	j
||d�S t t ¡ ¡}|d dd… }| j |¡}|d kr†t ||¡}|| j|< |S )Nr   )r(   Úcountéþÿÿÿé   )r"   r   r#   r   r$   r/   Úitemsizer%   Z
frombufferr1   ÚdataÚ	tracebackÚextract_stackÚsysÚ	_getframer.   Úgetr&   )r   r'   r(   r3   ÚstackZcallerÚresr   r   r   r)   d   s    

zFakeCUDAShared.arrayN)r   r   r   r   r   r)   r   r   r   r   r,   M   s   r,   c                   @   s|   e Z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S )ÚFakeCUDAAtomicc              	   C   s,   t � || }||  |7  < W 5 Q R X |S r   )Úaddlock©r   r)   ÚindexÚvalÚoldr   r   r   Úadd‹   s    zFakeCUDAAtomic.addc              	   C   s,   t � || }||  |8  < W 5 Q R X |S r   )ÚsublockrA   r   r   r   Úsub‘   s    zFakeCUDAAtomic.subc              	   C   s,   t � || }||  |M  < W 5 Q R X |S r   )ÚandlockrA   r   r   r   Úand_—   s    zFakeCUDAAtomic.and_c              	   C   s,   t � || }||  |O  < W 5 Q R X |S r   )ÚorlockrA   r   r   r   Úor_�   s    zFakeCUDAAtomic.or_c              	   C   s,   t � || }||  |N  < W 5 Q R X |S r   )ÚxorlockrA   r   r   r   Úxor£   s    zFakeCUDAAtomic.xorc              	   C   s>   t �0 || }||kr d||< n||  d7  < W 5 Q R X |S ©Nr   r   )ÚinclockrA   r   r   r   Úinc©   s    
zFakeCUDAAtomic.incc              	   C   sF   t �8 || }|dks||kr(|||< n||  d8  < W 5 Q R X |S rN   )ÚdeclockrA   r   r   r   Údec²   s    
zFakeCUDAAtomic.decc              	   C   s$   t � || }|||< W 5 Q R X |S r   )ÚexchlockrA   r   r   r   Úexch»   s    zFakeCUDAAtomic.exchc              	   C   s*   t � || }t||ƒ||< W 5 Q R X |S r   )ÚmaxlockÚmaxrA   r   r   r   rV   Á   s    zFakeCUDAAtomic.maxc              	   C   s*   t � || }t||ƒ||< W 5 Q R X |S r   )ÚminlockÚminrA   r   r   r   rX   Ç   s    zFakeCUDAAtomic.minc              	   C   s2   t �$ || }t || |g¡||< W 5 Q R X |S r   )rU   r%   ÚnanmaxrA   r   r   r   rY   Í   s     zFakeCUDAAtomic.nanmaxc              	   C   s2   t �$ || }t || |g¡||< W 5 Q R X |S r   )rW   r%   ÚnanminrA   r   r   r   rZ   Ó   s     zFakeCUDAAtomic.nanminc              
   C   sB   t �4 d|j }|| }||kr(|||< |W  5 Q R £ S Q R X d S )N)r   )Úcompare_and_swaplockÚndim)r   r)   rD   rC   rB   Úloadedr   r   r   Úcompare_and_swapÙ   s    
zFakeCUDAAtomic.compare_and_swapc              
   C   s8   t �* || }||kr|||< |W  5 Q R £ S Q R X d S r   )Úcaslock)r   r)   rB   rD   rC   r]   r   r   r   Úcasá   s
    zFakeCUDAAtomic.casN)r   r   r   rE   rG   rI   rK   rM   rP   rR   rT   rV   rX   rY   rZ   r^   r`   r   r   r   r   r?   Š   s   		r?   c                   @   sü   e Z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/d0„ Zd1d2„ Zd3d4„ Zd5d6„ Zd7d8„ Zd9d:„ Zd;d<„ Z d=S )>ÚFakeCUDAFp16c                 C   s   || S r   r   ©r   ÚaÚbr   r   r   Úhaddê   s    zFakeCUDAFp16.haddc                 C   s   || S r   r   rb   r   r   r   Úhsubí   s    zFakeCUDAFp16.hsubc                 C   s   || S r   r   rb   r   r   r   Úhmulð   s    zFakeCUDAFp16.hmulc                 C   s   || S r   r   rb   r   r   r   Úhdivó   s    zFakeCUDAFp16.hdivc                 C   s   || | S r   r   ©r   rc   rd   Úcr   r   r   Úhfmaö   s    zFakeCUDAFp16.hfmac                 C   s   | S r   r   ©r   rc   r   r   r   Úhnegù   s    zFakeCUDAFp16.hnegc                 C   s   t |ƒS r   )Úabsrl   r   r   r   Úhabsü   s    zFakeCUDAFp16.habsc                 C   s   t j|t jd�S r-   )r%   ÚsinÚfloat16©r   r
   r   r   r   Úhsinÿ   s    zFakeCUDAFp16.hsinc                 C   s   t j|t jd�S r-   )r%   Úcosrq   rr   r   r   r   Úhcos  s    zFakeCUDAFp16.hcosc                 C   s   t j|t jd�S r-   )r%   Úlogrq   rr   r   r   r   Úhlog  s    zFakeCUDAFp16.hlogc                 C   s   t j|t jd�S r-   )r%   Úlog2rq   rr   r   r   r   Úhlog2  s    zFakeCUDAFp16.hlog2c                 C   s   t j|t jd�S r-   )r%   Úlog10rq   rr   r   r   r   Úhlog10  s    zFakeCUDAFp16.hlog10c                 C   s   t j|t jd�S r-   )r%   Úexprq   rr   r   r   r   Úhexp  s    zFakeCUDAFp16.hexpc                 C   s   t j|t jd�S r-   )r%   Zexp2rq   rr   r   r   r   Úhexp2  s    zFakeCUDAFp16.hexp2c                 C   s   t  d| ¡S )Né
   ©r%   rq   rr   r   r   r   Úhexp10  s    zFakeCUDAFp16.hexp10c                 C   s   t j|t jd�S r-   )r%   Úsqrtrq   rr   r   r   r   Úhsqrt  s    zFakeCUDAFp16.hsqrtc                 C   s   t  |d ¡S )Ng      à¿r€   rr   r   r   r   Úhrsqrt  s    zFakeCUDAFp16.hrsqrtc                 C   s   t j|t jd�S r-   ©r%   Úceilrq   rr   r   r   r   Úhceil  s    zFakeCUDAFp16.hceilc                 C   s   t j|t jd�S r-   r…   rr   r   r   r   Úhfloor   s    zFakeCUDAFp16.hfloorc                 C   s   t j|t jd�S r-   )r%   Z
reciprocalrq   rr   r   r   r   Úhrcp#  s    zFakeCUDAFp16.hrcpc                 C   s   t j|t jd�S r-   )r%   Útruncrq   rr   r   r   r   Úhtrunc&  s    zFakeCUDAFp16.htruncc                 C   s   t j|t jd�S r-   )r%   Zrintrq   rr   r   r   r   Úhrint)  s    zFakeCUDAFp16.hrintc                 C   s   ||kS r   r   rb   r   r   r   Úheq,  s    zFakeCUDAFp16.heqc                 C   s   ||kS r   r   rb   r   r   r   Úhne/  s    zFakeCUDAFp16.hnec                 C   s   ||kS r   r   rb   r   r   r   Úhge2  s    zFakeCUDAFp16.hgec                 C   s   ||kS r   r   rb   r   r   r   Úhgt5  s    zFakeCUDAFp16.hgtc                 C   s   ||kS r   r   rb   r   r   r   Úhle8  s    zFakeCUDAFp16.hlec                 C   s   ||k S r   r   rb   r   r   r   Úhlt;  s    zFakeCUDAFp16.hltc                 C   s
   t ||ƒS r   )rV   rb   r   r   r   Úhmax>  s    zFakeCUDAFp16.hmaxc                 C   s
   t ||ƒS r   )rX   rb   r   r   r   ÚhminA  s    zFakeCUDAFp16.hminN)!r   r   r   re   rf   rg   rh   rk   rm   ro   rs   ru   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   ra   é   s<   ra   c                   @   s  e Zd ZdZdd„ Zedd„ ƒZedd„ ƒZedd	„ ƒZed
d„ ƒZ	edd„ ƒZ
edd„ ƒZedd„ ƒZedd„ ƒZe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d.d/„ Zd0d1„ Zd2d3„ Zd4d5„ Zd6d7„ Zd8S )9ÚFakeCUDAModulea7  
    An instance of this class will be injected into the __globals__ for an
    executing function in order to implement calls to cuda.*. This will fail to
    work correctly if the user code does::

        from numba import cuda as something_else

    In other words, the CUDA module must be called cuda.
    c                 C   s€   t |Ž | _t |Ž | _tƒ | _tƒ | _t|ƒ| _t	ƒ | _
tƒ | _tƒ | _t ¡ D ],\}}t| ||ƒ |jD ]}t| ||ƒ qhqNd S r   )r   ÚgridDimÚblockDimr   Ú_cgr!   Ú_localr,   Ú_sharedr*   Ú_constr?   Ú_atomicra   Ú_fp16r   ÚitemsÚsetattrÚaliases)r   Zgrid_dimZ	block_dimr2   ÚnameZsvtyÚaliasr   r   r   r   P  s    



zFakeCUDAModule.__init__c                 C   s   | j S r   )r˜   r   r   r   r   Úcgc  s    zFakeCUDAModule.cgc                 C   s   | j S r   )r™   r   r   r   r   Úlocalg  s    zFakeCUDAModule.localc                 C   s   | j S r   )rš   r   r   r   r   Úsharedk  s    zFakeCUDAModule.sharedc                 C   s   | j S r   )r›   r   r   r   r   Úconsto  s    zFakeCUDAModule.constc                 C   s   | j S r   )rœ   r   r   r   r   Úatomics  s    zFakeCUDAModule.atomicc                 C   s   | j S r   )r�   r   r   r   r   Úfp16w  s    zFakeCUDAModule.fp16c                 C   s
   t  ¡ jS r   )r   r   Ú	threadIdxr   r   r   r   r©   {  s    zFakeCUDAModule.threadIdxc                 C   s
   t  ¡ jS r   )r   r   ÚblockIdxr   r   r   r   rª     s    zFakeCUDAModule.blockIdxc                 C   s   dS ©Né    r   r   r   r   r   Úwarpsizeƒ  s    zFakeCUDAModule.warpsizec                 C   s   t  ¡ jd S r«   )r   r   Ú	thread_idr   r   r   r   Úlaneid‡  s    zFakeCUDAModule.laneidc                 C   s   t  ¡  ¡  d S r   r   r   r   r   r   r   ‹  s    zFakeCUDAModule.syncthreadsc                 C   s   d S r   r   r   r   r   r   ÚthreadfenceŽ  s    zFakeCUDAModule.threadfencec                 C   s   d S r   r   r   r   r   r   Úthreadfence_block’  s    z FakeCUDAModule.threadfence_blockc                 C   s   d S r   r   r   r   r   r   Úthreadfence_system–  s    z!FakeCUDAModule.threadfence_systemc                 C   s   t  ¡  |¡S r   )r   r   Úsyncthreads_count©r   rC   r   r   r   r³   š  s    z FakeCUDAModule.syncthreads_countc                 C   s   t  ¡  |¡S r   )r   r   Úsyncthreads_andr´   r   r   r   rµ   �  s    zFakeCUDAModule.syncthreads_andc                 C   s   t  ¡  |¡S r   )r   r   Úsyncthreads_orr´   r   r   r   r¶      s    zFakeCUDAModule.syncthreads_orc                 C   s   t |ƒ d¡S )NÚ1)Úbinr3   r´   r   r   r   Úpopc£  s    zFakeCUDAModule.popcc                 C   s   || | S r   r   ri   r   r   r   Úfma¦  s    zFakeCUDAModule.fmac                 C   s   |d S )NgUUUUUUÕ?r   rl   r   r   r   Úcbrt©  s    zFakeCUDAModule.cbrtc                 C   s   t d |¡d d d… dƒS )Nú{:032b}éÿÿÿÿr5   )ÚintÚformatr´   r   r   r   Úbrev¬  s    zFakeCUDAModule.brevc                 C   s    d  |¡}t|ƒt| d¡ƒ S )Nr¼   Ú0)r¿   ÚlenÚlstrip)r   rC   Úsr   r   r   Úclz¯  s    
zFakeCUDAModule.clzc                 C   s,   d  |¡}t|ƒt| d¡ƒ d d }|S )Nr¼   rÁ   r   é!   )r¿   rÂ   Úrstrip)r   rC   rÄ   Úrr   r   r   Úffs³  s    
zFakeCUDAModule.ffsc                 C   s   |r|S |S r   r   ri   r   r   r   Úselp½  s    zFakeCUDAModule.selpc                 C   s†   | j }| j}| j}|j|j |j }|dkr0|S |j|j |j }|dkrR||fS |j|j |j }|dkrv|||fS td| ƒ‚d S )Nr   r5   é   z*Global ID has 1-3 dimensions. %d requested)r—   rª   r©   r
   r   r   ÚRuntimeError)r   ÚnÚbdimÚbidÚtidr
   r   r   r   r   r   ÚgridÀ  s    
zFakeCUDAModule.gridc                 C   sn   | j }| j}|j|j }|dkr$|S |j|j }|dkr@||fS |j|j }|dkr^|||fS td| ƒ‚d S )Nr   r5   rË   z,Global grid has 1-3 dimensions. %d requested)r—   r–   r
   r   r   rÌ   )r   rÍ   rÎ   Zgdimr
   r   r   r   r   r   ÚgridsizeÐ  s    
zFakeCUDAModule.gridsizeN) r   r   r   r   r   Úpropertyr£   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•   E  sL   











r•   c              	   #   sl   ddl m‰  | j}t‡ fdd„| ¡ D ƒƒ}t‡fdd„| ¡ D ƒƒ}| |¡ z
d V  W 5 | |¡ X d S )Nr   ©Úcudac                 3   s"   | ]\}}|ˆ kr||fV  qd S r   r   ©Ú.0ÚkÚvrÔ   r   r   Ú	<genexpr>æ  s      z&swapped_cuda_module.<locals>.<genexpr>c                 3   s   | ]\}}|ˆ fV  qd S r   r   rÖ   )Úfake_cuda_moduler   r   rÚ   è  s     )ZnumbarÕ   Ú__globals__Údictrž   Úupdate)ÚfnrÛ   Zfn_globsÚorigÚreplr   )rÕ   rÛ   r   Úswapped_cuda_moduleà  s    

râ   )%r   Ú
contextlibr   r:   r   r8   Z
numba.corer   Únumpyr%   Znumba.npr   r   Úobjectr   r   r   r!   r*   r,   ÚLockr@   rF   rH   rJ   rL   rU   rW   r[   r_   rO   rQ   rS   r?   ra   r•   râ   r   r   r   r   Ú<module>   s@   
/_\ 