§
    1@Äj\2  ã                   óÐ  — 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	e¦  «        Z G d
„ d¦  «        Z G d„ d¦  «        Z G d„ de¦  «        Z G d„ de¦  «        Z G d„ de¦  «        Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z ej        ¦   «         Z  G d„ de¦  «        Z! G d„ de¦  «        Z" G d„ de¦  «        Z#e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                   ó*   — e Zd ZdZd„ Zd„ Zd„ Zd„ ZdS )ÚDim3z;
    Used to implement thread/block indices/dimensions
    c                 ó0   — || _         || _        || _        d S ©N©ÚxÚyÚz)Úselfr   r   r   s       úOC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/simulator/kernelapi.pyÚ__init__zDim3.__init__   s   € ØˆŒØˆŒØˆŒˆˆó    c                 ó8   — d| j         ›d| j        ›d| j        ›d�S )Nú(ú, ú)r   ©r   s    r   Ú__str__zDim3.__str__   s#   € € Ø!%¤  ¨¬¨¨°´°°Ð8Ð8r   c                 ó8   — d| j         ›d| j        ›d| j        ›d�S )NzDim3(r   r   r   r   s    r   Ú__repr__zDim3.__repr__   s#   € € Ø%)¤V V V¨T¬V¨V¨V°T´V°V°VÐ<Ð<r   c              #   ó@   K  — | j         V — | j        V — | j        V — d S r   r   r   s    r   Ú__iter__zDim3.__iter__!   s+   è è € ØŒfˆˆˆØŒfˆˆˆØŒfˆˆˆˆˆr   N)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   r   r   r   © r   r   r	   r	      sZ   € € € € € ðð ðð ð ð
9ð 9ð 9ð=ð =ð =ðð ð ð ð r   r	   c                   ó   — e Zd ZdZd„ ZdS )Ú	GridGroupz+
    Used to implement the grid group.
    c                 óP   — t          j        ¦   «                              ¦   «          d S r   ©Ú	threadingÚcurrent_threadÚsyncthreadsr   s    r   ÚsynczGridGroup.sync,   s%   € õ 	Ô Ñ"Ô"×.Ò.Ñ0Ô0Ð0Ð0Ð0r   N)r   r   r    r!   r*   r"   r   r   r$   r$   '   s-   € € € € € ðð ð1ð 1ð 1ð 1ð 1r   r$   c                   ó   — e Zd ZdZd„ ZdS )Ú
FakeCUDACgz!
    CUDA Cooperative Groups
    c                 ó   — t          ¦   «         S r   )r$   r   s    r   Ú	this_gridzFakeCUDACg.this_grid7   s   € Ý‰{Œ{Ðr   N)r   r   r    r!   r.   r"   r   r   r,   r,   3   s-   € € € € € ðð ðð ð ð ð r   r,   c                   ó   — e Zd ZdZd„ ZdS )ÚFakeCUDALocalz
    CUDA Local arrays
    c                 óˆ   — t          |t          j        ¦  «        rt          j        |¦  «        }t          j        ||¦  «        S r   )Ú
isinstancer   ÚTyper   Úas_dtypeÚnpÚempty)r   ÚshapeÚdtypes      r   ÚarrayzFakeCUDALocal.array?   s8   € Ý�e�UœZÑ(Ô(ð 	2Ý!Ô*¨5Ñ1Ô1ˆEÝŒx˜˜uÑ%Ô%Ð%r   N)r   r   r    r!   r9   r"   r   r   r0   r0   ;   s-   € € € € € ðð ð&ð &ð &ð &ð &r   r0   c                   ó   — e Zd ZdZd„ ZdS )ÚFakeCUDAConstz
    CUDA Const arrays
    c                 ó   — |S r   r"   )r   Úarys     r   Ú
array_likezFakeCUDAConst.array_likeI   s   € Øˆ
r   N)r   r   r    r!   r>   r"   r   r   r;   r;   E   s-   € € € € € ðð ðð ð ð ð r   r;   c                   ó   — e Zd ZdZd„ Z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                 ól   — i | _         || _        t          j        |t          j        ¬¦  «        | _        d S ©N)r8   )Ú_allocationsÚ_dynshared_sizer5   ÚzerosÚbyteÚ
_dynshared)r   Údynshared_sizes     r   r   zFakeCUDAShared.__init___   s.   € ØˆÔØ-ˆÔÝœ( >½¼ÐAÑAÔAˆŒˆˆr   c                 ó®  — t          |t          j        ¦  «        rt          j        |¦  «        }|dk    r0| j        |j        z  }t          j        | j	        j
        ||¬¦  «        S t          j        t          j        ¦   «         ¦  «        }|d         dd…         }| j                             |¦  «        }|€t          j        ||¦  «        }|| j        |<   |S )Nr   )r8   Úcountéþÿÿÿé   )r2   r   r3   r   r4   rD   Úitemsizer5   Ú
frombufferrG   ÚdataÚ	tracebackÚextract_stackÚsysÚ	_getframerC   Úgetr6   )r   r7   r8   rJ   ÚstackÚcallerÚress          r   r9   zFakeCUDAShared.arrayd   sÅ   € Ý�e�UœZÑ(Ô(ð 	2Ý!Ô*¨5Ñ1Ô1ˆEð �AŠ:ˆ:ð Ô(¨E¬NÑ:ˆEÝ”= ¤Ô!5¸UÈ%ÐPÑPÔPÐPõ
 Ô'­¬©¬Ñ8Ô8ˆØ�r”˜1˜Q˜3”ˆØÔ×#Ò# FÑ+Ô+ˆØˆ;Ý”(˜5 %Ñ(Ô(ˆCØ(+ˆDÔ˜fÑ%Øˆ
r   N)r   r   r    r!   r   r9   r"   r   r   r@   r@   M   s?   € € € € € ðð ð"Bð Bð Bð
ð ð ð ð r   r@   c                   ób   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z	d„ Z
d	„ Zd
„ Zd„ Zd„ Zd„ Zd„ ZdS )ÚFakeCUDAAtomicc                 ót   — t           5  ||         }||xx         |z  cc<   d d d ¦  «         n# 1 swxY w Y   |S r   )Úaddlock©r   r9   ÚindexÚvalÚolds        r   ÚaddzFakeCUDAAtomic.add‹   óƒ   € Ýð 	 ð 	 Ø˜”,ˆCØ�%ˆLˆLŒL˜CÑˆLˆL‰Lð	 ð 	 ð 	 ñ 	 ô 	 ð 	 ð 	 ð 	 ð 	 ð 	 ð 	 øøøð 	 ð 	 ð 	 ð 	 ð ˆ
ó   ˆ-­1´1c                 ót   — t           5  ||         }||xx         |z  cc<   d d d ¦  «         n# 1 swxY w Y   |S r   )Úsublockr\   s        r   ÚsubzFakeCUDAAtomic.sub‘   ra   rb   c                 ót   — t           5  ||         }||xx         |z  cc<   d d d ¦  «         n# 1 swxY w Y   |S r   )Úandlockr\   s        r   Úand_zFakeCUDAAtomic.and_—   ra   rb   c                 ót   — t           5  ||         }||xx         |z  cc<   d d d ¦  «         n# 1 swxY w Y   |S r   )Úorlockr\   s        r   Úor_zFakeCUDAAtomic.or_�   sƒ   € Ýð 	 ð 	 Ø˜”,ˆCØ�%ˆLˆLŒL˜CÑˆLˆL‰Lð	 ð 	 ð 	 ñ 	 ô 	 ð 	 ð 	 ð 	 ð 	 ð 	 ð 	 øøøð 	 ð 	 ð 	 ð 	 ð ˆ
rb   c                 ót   — t           5  ||         }||xx         |z  cc<   d d d ¦  «         n# 1 swxY w Y   |S r   )Úxorlockr\   s        r   ÚxorzFakeCUDAAtomic.xor£   ra   rb   c                 óŒ   — t           5  ||         }||k    rd||<   n||xx         dz  cc<   d d d ¦  «         n# 1 swxY w Y   |S ©Nr   r   )Úinclockr\   s        r   ÚinczFakeCUDAAtomic.inc©   s—   € Ýð 	"ð 	"Ø˜”,ˆCØ�cŠzˆzØ ��e‘�à�e��” Ñ!��‘ð	"ð 	"ð 	"ñ 	"ô 	"ð 	"ð 	"ð 	"ð 	"ð 	"ð 	"øøøð 	"ð 	"ð 	"ð 	"ð ˆ
s   ˆ%9¹=Á =c                 ó˜   — t           5  ||         }|dk    s||k    r|||<   n||xx         dz  cc<   d d d ¦  «         n# 1 swxY w Y   |S rp   )Údeclockr\   s        r   ÚdeczFakeCUDAAtomic.dec²   sŸ   € Ýð 	"ð 	"Ø˜”,ˆCØ�q’�˜c Cši˜iØ"��e‘�à�e��” Ñ!��‘ð	"ð 	"ð 	"ñ 	"ô 	"ð 	"ð 	"ð 	"ð 	"ð 	"ð 	"øøøð 	"ð 	"ð 	"ð 	"ð ˆ
s   ˆ+?¿AÁAc                 ó^   — t           5  ||         }|||<   d d d ¦  «         n# 1 swxY w Y   |S r   )Úexchlockr\   s        r   ÚexchzFakeCUDAAtomic.exch»   sv   € Ýð 	ð 	Ø˜”,ˆCØˆE�%‰Lð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð ˆ
s   ˆ"¢&©&c                 óz   — t           5  ||         }t          ||¦  «        ||<   d d d ¦  «         n# 1 swxY w Y   |S r   )ÚmaxlockÚmaxr\   s        r   r{   zFakeCUDAAtomic.maxÁ   ó~   € Ýð 	)ð 	)Ø˜”,ˆCÝ˜s C™=œ=ˆE�%‰Lð	)ð 	)ð 	)ñ 	)ô 	)ð 	)ð 	)ð 	)ð 	)ð 	)ð 	)øøøð 	)ð 	)ð 	)ð 	)ð ˆ
ó   ˆ0°4·4c                 óz   — t           5  ||         }t          ||¦  «        ||<   d d d ¦  «         n# 1 swxY w Y   |S r   )ÚminlockÚminr\   s        r   r€   zFakeCUDAAtomic.minÇ   r|   r}   c                 ó’   — t           5  ||         }t          j        ||         |g¦  «        ||<   d d d ¦  «         n# 1 swxY w Y   |S r   )rz   r5   Únanmaxr\   s        r   r‚   zFakeCUDAAtomic.nanmaxÍ   ó‰   € Ýð 	:ð 	:Ø˜”,ˆCÝœ9 e¨E¤l°CÐ%8Ñ9Ô9ˆE�%‰Lð	:ð 	:ð 	:ñ 	:ô 	:ð 	:ð 	:ð 	:ð 	:ð 	:ð 	:øøøð 	:ð 	:ð 	:ð 	:ð ˆ
ó   ˆ(<¼A ÁA c                 ó’   — t           5  ||         }t          j        ||         |g¦  «        ||<   d d d ¦  «         n# 1 swxY w Y   |S r   )r   r5   Únanminr\   s        r   r†   zFakeCUDAAtomic.nanminÓ   rƒ   r„   c                 ó‚   — t           5  d|j        z  }||         }||k    r|||<   |cd d d ¦  «         S # 1 swxY w Y   d S )N)r   )Úcompare_and_swaplockÚndim)r   r9   r_   r^   r]   Úloadeds         r   Úcompare_and_swapzFakeCUDAAtomic.compare_and_swapÙ   s–   € Ý!ð 	ð 	Ø˜5œ:Ñ%ˆEØ˜5”\ˆFØ˜Š}ˆ}Ø"��e‘Øð	ð 	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	ð 	s   ˆ4´8»8c                 ón   — t           5  ||         }||k    r|||<   |cd d d ¦  «         S # 1 swxY w Y   d S r   )Úcaslock)r   r9   r]   r_   r^   rŠ   s         r   ÚcaszFakeCUDAAtomic.casá   sŠ   € Ýð 	ð 	Ø˜5”\ˆFØ˜Š}ˆ}Ø"��e‘Øð		ð 	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	ð 	s   ˆ*ª.±.N)r   r   r    r`   re   rh   rk   rn   rr   ru   rx   r{   r€   r‚   r†   r‹   rŽ   r"   r   r   rY   rY   Š   sæ   € € € € € ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð ð r   rY   c                   óÂ   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z	d„ Z
d	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z dS ) ÚFakeCUDAFp16c                 ó   — ||z   S r   r"   ©r   ÚaÚbs      r   ÚhaddzFakeCUDAFp16.haddê   ó   € Ø�1‰uˆr   c                 ó   — ||z
  S r   r"   r’   s      r   ÚhsubzFakeCUDAFp16.hsubí   r–   r   c                 ó   — ||z  S r   r"   r’   s      r   ÚhmulzFakeCUDAFp16.hmulð   r–   r   c                 ó   — ||z  S r   r"   r’   s      r   ÚhdivzFakeCUDAFp16.hdivó   r–   r   c                 ó   — ||z  |z   S r   r"   ©r   r“   r”   Úcs       r   ÚhfmazFakeCUDAFp16.hfmaö   ó   € Ø�1‰u�q‰yÐr   c                 ó   — | S r   r"   ©r   r“   s     r   ÚhnegzFakeCUDAFp16.hnegù   s	   € Øˆrˆ	r   c                 ó    — t          |¦  «        S r   )Úabsr£   s     r   ÚhabszFakeCUDAFp16.habsü   s   € Ý�1‰vŒvˆr   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   ÚsinÚfloat16©r   r   s     r   ÚhsinzFakeCUDAFp16.hsinÿ   ó   € ÝŒv�a�rœzÐ*Ñ*Ô*Ð*r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úcosrª   r«   s     r   ÚhcoszFakeCUDAFp16.hcos  r­   r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úlogrª   r«   s     r   ÚhlogzFakeCUDAFp16.hlog  r­   r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úlog2rª   r«   s     r   Úhlog2zFakeCUDAFp16.hlog2  ó   € ÝŒw�q¥¤
Ð+Ñ+Ô+Ð+r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úlog10rª   r«   s     r   Úhlog10zFakeCUDAFp16.hlog10  ó   € ÝŒx˜¥¤Ð,Ñ,Ô,Ð,r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úexprª   r«   s     r   ÚhexpzFakeCUDAFp16.hexp  r­   r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úexp2rª   r«   s     r   Úhexp2zFakeCUDAFp16.hexp2  r·   r   c                 ó0   — t          j        d|z  ¦  «        S )Né
   ©r5   rª   r«   s     r   Úhexp10zFakeCUDAFp16.hexp10  s   € ÝŒz˜" ™'Ñ"Ô"Ð"r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úsqrtrª   r«   s     r   ÚhsqrtzFakeCUDAFp16.hsqrt  r·   r   c                 ó0   — t          j        |dz  ¦  «        S )Ng      à¿rÄ   r«   s     r   ÚhrsqrtzFakeCUDAFp16.hrsqrt  s   € ÝŒz˜!˜t™)Ñ$Ô$Ð$r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   ©r5   Úceilrª   r«   s     r   ÚhceilzFakeCUDAFp16.hceil  r·   r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   rÌ   r«   s     r   ÚhfloorzFakeCUDAFp16.hfloor   r·   r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Ú
reciprocalrª   r«   s     r   ÚhrcpzFakeCUDAFp16.hrcp#  s   € ÝŒ}˜Q¥b¤jÐ1Ñ1Ô1Ð1r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Útruncrª   r«   s     r   ÚhtrunczFakeCUDAFp16.htrunc&  r»   r   c                 óB   — t          j        |t           j        ¬¦  «        S rB   )r5   Úrintrª   r«   s     r   ÚhrintzFakeCUDAFp16.hrint)  r·   r   c                 ó   — ||k    S r   r"   r’   s      r   ÚheqzFakeCUDAFp16.heq,  ó   € Ø�AŠvˆr   c                 ó   — ||k    S r   r"   r’   s      r   ÚhnezFakeCUDAFp16.hne/  rÜ   r   c                 ó   — ||k    S r   r"   r’   s      r   ÚhgezFakeCUDAFp16.hge2  rÜ   r   c                 ó   — ||k    S r   r"   r’   s      r   ÚhgtzFakeCUDAFp16.hgt5  ó   € Ø�1Šuˆr   c                 ó   — ||k    S r   r"   r’   s      r   ÚhlezFakeCUDAFp16.hle8  rÜ   r   c                 ó   — ||k     S r   r"   r’   s      r   ÚhltzFakeCUDAFp16.hlt;  rã   r   c                 ó"   — t          ||¦  «        S r   )r{   r’   s      r   ÚhmaxzFakeCUDAFp16.hmax>  ó   € Ý�1�a‰yŒyÐr   c                 ó"   — t          ||¦  «        S r   )r€   r’   s      r   ÚhminzFakeCUDAFp16.hminA  rê   r   N)!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�   é   sÖ  € € € € € ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð+ð +ð +ð+ð +ð +ð+ð +ð +ð,ð ,ð ,ð-ð -ð -ð+ð +ð +ð,ð ,ð ,ð#ð #ð #ð,ð ,ð ,ð%ð %ð %ð,ð ,ð ,ð,ð ,ð ,ð2ð 2ð 2ð-ð -ð -ð,ð ,ð ,ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð ð r   r�   c                   óT  — e Zd ZdZd„ Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Z	ed„ ¦   «         Z
ed„ ¦   «         Zed	„ ¦   «         Zed
„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ ZdS )Ú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                 ó®  — t          |Ž | _        t          |Ž | _        t          ¦   «         | _        t          ¦   «         | _        t          |¦  «        | _        t          ¦   «         | _
        t          ¦   «         | _        t          ¦   «         | _        t          j        ¦   «         D ]1\  }}t#          | ||¦  «         |j        D ]}t#          | ||¦  «         ŒŒ2d S r   )r	   ÚgridDimÚblockDimr,   Ú_cgr0   Ú_localr@   Ú_sharedr;   Ú_constrY   Ú_atomicr�   Ú_fp16r   ÚitemsÚsetattrÚaliases)r   Úgrid_dimÚ	block_dimrH   ÚnameÚsvtyÚaliass          r   r   zFakeCUDAModule.__init__P  sÆ   € Ý˜X�ˆŒÝ˜iÐ(ˆŒÝ‘<”<ˆŒÝ#‘o”oˆŒÝ% nÑ5Ô5ˆŒÝ#‘o”oˆŒÝ%Ñ'Ô'ˆŒÝ!‘^”^ˆŒ
õ 'Ô,Ñ.Ô.ð 	+ð 	+‰JˆD�$Ý�D˜$ Ñ%Ô%Ð%Øœð +ð +�Ý˜˜e TÑ*Ô*Ð*Ð*ð+ð	+ð 	+r   c                 ó   — | j         S r   )rò   r   s    r   ÚcgzFakeCUDAModule.cgc  s	   € àŒxˆr   c                 ó   — | j         S r   )ró   r   s    r   ÚlocalzFakeCUDAModule.localg  ó
   € àŒ{Ðr   c                 ó   — | j         S r   )rô   r   s    r   ÚsharedzFakeCUDAModule.sharedk  ó
   € àŒ|Ðr   c                 ó   — | j         S r   )rõ   r   s    r   ÚconstzFakeCUDAModule.consto  r  r   c                 ó   — | j         S r   )rö   r   s    r   ÚatomiczFakeCUDAModule.atomics  r  r   c                 ó   — | j         S r   )r÷   r   s    r   Úfp16zFakeCUDAModule.fp16w  s
   € àŒzÐr   c                 ó2   — t          j        ¦   «         j        S r   )r'   r(   Ú	threadIdxr   s    r   r  zFakeCUDAModule.threadIdx{  s   € åÔ'Ñ)Ô)Ô3Ð3r   c                 ó2   — t          j        ¦   «         j        S r   )r'   r(   ÚblockIdxr   s    r   r  zFakeCUDAModule.blockIdx  s   € åÔ'Ñ)Ô)Ô2Ð2r   c                 ó   — dS ©Né    r"   r   s    r   ÚwarpsizezFakeCUDAModule.warpsizeƒ  s   € àˆrr   c                 ó8   — t          j        ¦   «         j        dz  S r  )r'   r(   Ú	thread_idr   s    r   ÚlaneidzFakeCUDAModule.laneid‡  s   € åÔ'Ñ)Ô)Ô3°bÑ8Ð8r   c                 óP   — t          j        ¦   «                              ¦   «          d S r   r&   r   s    r   r)   zFakeCUDAModule.syncthreads‹  s#   € ÝÔ Ñ"Ô"×.Ò.Ñ0Ô0Ð0Ð0Ð0r   c                 ó   — d S r   r"   r   s    r   ÚthreadfencezFakeCUDAModule.threadfenceŽ  ó   € àˆr   c                 ó   — d S r   r"   r   s    r   Úthreadfence_blockz FakeCUDAModule.threadfence_block’  r  r   c                 ó   — d S r   r"   r   s    r   Úthreadfence_systemz!FakeCUDAModule.threadfence_system–  r  r   c                 óN   — t          j        ¦   «                              |¦  «        S r   )r'   r(   Úsyncthreads_count©r   r^   s     r   r"  z FakeCUDAModule.syncthreads_countš  s   € ÝÔ'Ñ)Ô)×;Ò;¸CÑ@Ô@Ð@r   c                 óN   — t          j        ¦   «                              |¦  «        S r   )r'   r(   Úsyncthreads_andr#  s     r   r%  zFakeCUDAModule.syncthreads_and�  s   € ÝÔ'Ñ)Ô)×9Ò9¸#Ñ>Ô>Ð>r   c                 óN   — t          j        ¦   «                              |¦  «        S r   )r'   r(   Úsyncthreads_orr#  s     r   r'  zFakeCUDAModule.syncthreads_or   s   € ÝÔ'Ñ)Ô)×8Ò8¸Ñ=Ô=Ð=r   c                 óF   — t          |¦  «                             d¦  «        S )NÚ1)ÚbinrJ   r#  s     r   ÚpopczFakeCUDAModule.popc£  s   € Ý�3‰xŒx�~Š~˜cÑ"Ô"Ð"r   c                 ó   — ||z  |z   S r   r"   rž   s       r   ÚfmazFakeCUDAModule.fma¦  r¡   r   c                 ó   — |dz  S )NgUUUUUUÕ?r"   r£   s     r   ÚcbrtzFakeCUDAModule.cbrt©  s   € Ø�U‰|Ðr   c                 óZ   — t          d                     |¦  «        d d d…         d¦  «        S )Nú{:032b}éÿÿÿÿrL   )ÚintÚformatr#  s     r   ÚbrevzFakeCUDAModule.brev¬  s+   € Ý�9×#Ò# CÑ(Ô(¨¨¨2¨Ô.°Ñ2Ô2Ð2r   c                 ó�   — d                      |¦  «        }t          |¦  «        t          |                     d¦  «        ¦  «        z
  S )Nr1  Ú0)r4  ÚlenÚlstrip)r   r^   Úss      r   ÚclzzFakeCUDAModule.clz¯  s8   € Ø×Ò˜SÑ!Ô!ˆÝ�1‰vŒv�˜AŸHšH S™MœMÑ*Ô*Ñ*Ð*r   c                 ó    — d                      |¦  «        }t          |¦  «        t          |                     d¦  «        ¦  «        z
  dz   dz  }|S )Nr1  r7  r   é!   )r4  r8  Úrstrip)r   r^   r:  Úrs       r   ÚffszFakeCUDAModule.ffs³  sH   € ð ×Ò˜SÑ!Ô!ˆÝ�‰VŒV•c˜!Ÿ(š( 3™-œ-Ñ(Ô(Ñ(¨1Ñ,°Ñ2ˆØˆr   c                 ó   — |r|n|S r   r"   rž   s       r   ÚselpzFakeCUDAModule.selp½  s   € Øˆ}ˆqˆq˜1Ðr   c                 ó  — | j         }| j        }| j        }|j        |j        z  |j        z   }|dk    r|S |j        |j        z  |j        z   }|dk    r||fS |j        |j        z  |j        z   }|dk    r|||fS t          d|z  ¦  «        ‚)Nr   rL   é   z*Global ID has 1-3 dimensions. %d requested)rñ   r  r  r   r   r   ÚRuntimeError)r   ÚnÚbdimÚbidÚtidr   r   r   s           r   ÚgridzFakeCUDAModule.gridÀ  sœ   € ØŒ}ˆØŒmˆØŒnˆØŒE�D”F‰N˜SœUÑ"ˆØ�Š6ˆ6ØˆHØŒE�D”F‰N˜SœUÑ"ˆØ�Š6ˆ6Ø�q�6ˆMØŒE�D”F‰N˜SœUÑ"ˆØ�Š6ˆ6Ø�q˜!�9ÐåÐGÈ!ÑKÑLÔLÐLr   c                 óÖ   — | j         }| j        }|j        |j        z  }|dk    r|S |j        |j        z  }|dk    r||fS |j        |j        z  }|dk    r|||fS t          d|z  ¦  «        ‚)Nr   rL   rD  z,Global grid has 1-3 dimensions. %d requested)rñ   rð   r   r   r   rE  )r   rF  rG  Úgdimr   r   r   s          r   ÚgridsizezFakeCUDAModule.gridsizeÐ  s€   € ØŒ}ˆØŒ|ˆØŒF�T”V‰OˆØ�Š6ˆ6ØˆHØŒF�T”V‰OˆØ�Š6ˆ6Ø�q�6ˆMØŒF�T”V‰OˆØ�Š6ˆ6Ø�q˜!�9ÐåÐIÈAÑMÑNÔNÐNr   N) 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/  r5  r;  r@  rB  rJ  rM  r"   r   r   rî   rî   E  s6  € € € € € ðð ð+ð +ð +ð& ðð ñ „Xðð ðð ñ „Xðð ðð ñ „Xðð ðð ñ „Xðð ðð ñ „Xðð ðð ñ „Xðð ð4ð 4ñ „Xð4ð ð3ð 3ñ „Xð3ð ðð ñ „Xðð ð9ð 9ñ „Xð9ð1ð 1ð 1ðð ð ðð ð ðð ð ðAð Að Að?ð ?ð ?ð>ð >ð >ð#ð #ð #ðð ð ðð ð ð3ð 3ð 3ð+ð +ð +ðð ð ðð ð ðMð Mð Mð Oð Oð Oð Oð Or   rî   c              #   ón  ‡‡K  — ddl mŠ | j        }t          ˆfd„|                     ¦   «         D ¦   «         ¦  «        }t          ˆfd„|                     ¦   «         D ¦   «         ¦  «        }|                     |¦  «         	 d V — |                     |¦  «         d S # |                     |¦  «         w xY w)Nr   )Úcudac              3   ó.   •K  — | ]\  }}|‰u ¯	||fV — Œd S r   r"   )Ú.0ÚkÚvrP  s      €r   ú	<genexpr>z&swapped_cuda_module.<locals>.<genexpr>æ  s/   øè è € ÐAÐA™4˜1˜a°q¸D°y°y��A�°y°y°y°yÐAÐAr   c              3   ó&   •K  — | ]\  }}|‰fV — Œd S r   r"   )rR  rS  rT  Úfake_cuda_modules      €r   rU  z&swapped_cuda_module.<locals>.<genexpr>è  s-   øè è € Ð?Ð?©$¨!¨Q�Ð$Ð%Ð?Ð?Ð?Ð?Ð?Ð?r   )ÚnumbarP  Ú__globals__Údictrø   Úupdate)ÚfnrW  Úfn_globsÚorigÚreplrP  s    `   @r   Úswapped_cuda_moduler`  à  sÉ   øøè è € àÐÐÐÐÐàŒ~€HåÐAÐAÐAÐA 8§>¢>Ñ#3Ô#3ÐAÑAÔAÑAÔA€DåÐ?Ð?Ð?Ð?°$·*²*±,´,Ð?Ñ?Ô?Ñ?Ô?€Dà‡O‚O�DÑÔÐðØˆˆˆð 	�Š˜ÑÔÐÐÐøˆ�Š˜ÑÔÐÐøøøs   ÂB ÂB4)%r!   Ú
contextlibr   rR   r'   rP   Ú
numba.corer   Únumpyr5   Únumba.npr   r   Úobjectr	   r$   r,   r0   r;   r@   ÚLockr[   rd   rg   rj   rm   rz   r   rˆ   r�   rq   rt   rw   rY   r�   rî   r`  r"   r   r   ú<module>rg     sü  ððð ð
 &Ð %Ð %Ð %Ð %Ð %Ø 
€
€
€
Ø Ð Ð Ð Ø Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð à "Ð "Ð "Ð "Ð "Ð "à &Ð &Ð &Ð &Ð &Ð &ðð ð ð ð ˆ6ñ ô ð ð*	1ð 	1ð 	1ð 	1ð 	1ñ 	1ô 	1ð 	1ðð ð ð ð ñ ô ð ð&ð &ð &ð &ð &�Fñ &ô &ð &ðð ð ð ð �Fñ ô ð ð,ð ,ð ,ð ,ð ,�Vñ ,ô ,ð ,ð^ ˆ)Œ.Ñ
Ô
€Ø
ˆ)Œ.Ñ
Ô
€Ø
ˆ)Œ.Ñ
Ô
€Ø	ˆŒÑ	Ô	€Ø
ˆ)Œ.Ñ
Ô
€Ø
ˆ)Œ.Ñ
Ô
€Ø
ˆ)Œ.Ñ
Ô
€Ø%�y”~Ñ'Ô'Ð Ø
ˆ)Œ.Ñ
Ô
€Ø
ˆ)Œ.Ñ
Ô
€Ø
ˆ)Œ.Ñ
Ô
€Øˆ9Œ>ÑÔ€ð\ð \ð \ð \ð \�Vñ \ô \ð \ð~Yð Yð Yð Yð Y�6ñ Yô Yð YðxXOð XOð XOð XOð XO�Vñ XOô XOð XOðv ðð ñ „ðð ð r   