Ë
    üÇ:jm0  ã                   óÂ  — 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$y)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y)ÚDim3z;
    Used to implement thread/block indices/dimensions
    c                 ó.   — || _         || _        || _        y ©N©ÚxÚyÚz)Úselfr   r   r   s       ús/home/mcse/projects/srt_converter/srt-converter-venv/lib/python3.12/site-packages/numba/cuda/simulator/kernelapi.pyÚ__init__zDim3.__init__   s   € ØˆŒØˆŒØˆ�ó    c                 óV   — d| j                   ›d| j                  ›d| j                  ›d�S )Nú(ú, ú)r   ©r   s    r   Ú__str__zDim3.__str__   s   � Ø!%§£¨¯«°·³Ð8Ð8r   c                 óV   — d| j                   ›d| j                  ›d| j                  ›d�S )NzDim3(r   r   r   r   s    r   Ú__repr__zDim3.__repr__   s   � Ø%)§V£V¨T¯V«V°T·V³VÐ<Ð<r   c              #   ó`   K  — | j                   –— | j                  –— | j                  –— y ­wr   r   r   s    r   Ú__iter__zDim3.__iter__!   s!   è ø€ Ø�f‰fŠØ�f‰fŠØ�f‰f‹ùs   ‚,.N)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   r   r   r   © r   r   r	   r	      s   „ ñòò
9ò=ór   r	   c                   ó   — e Zd ZdZd„ Zy)Ú	GridGroupz+
    Used to implement the grid group.
    c                 óH   — t        j                  «       j                  «        y r   ©Ú	threadingÚcurrent_threadÚsyncthreadsr   s    r   ÚsynczGridGroup.sync,   s   € ô 	× Ñ Ó"×.Ñ.Õ0r   N)r   r   r    r!   r*   r"   r   r   r$   r$   '   s   „ ñó1r   r$   c                   ó   — e Zd ZdZd„ Zy)Ú
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y)Ú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?   s1   € Ü�eœUŸZ™ZÔ(Ü!×*Ñ*¨5Ó1ˆEÜ�x‰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y)Ú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y)Ú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                 ót   — i | _         || _        t        j                  |t        j                  ¬«      | _        y ©N)r8   )Ú_allocationsÚ_dynshared_sizer5   ÚzerosÚbyteÚ
_dynshared)r   Údynshared_sizes     r   r   zFakeCUDAShared.__init___   s(   € ØˆÔØ-ˆÔÜŸ(™( >¼¿¹ÔAˆ�r   c                 óÚ  — t        |t        j                  «      rt        j                  |«      }|dk(  rE| j
                  |j                  z  }t        j                  | j                  j                  ||¬«      S t        j                  t        j                  «       «      }|d   dd }| j                  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™ZÔ(Ü!×*Ñ*¨5Ó1ˆEð �AŠ:ð ×(Ñ(¨E¯N©NÑ:ˆEÜ—=‘= §¡×!5Ñ!5¸UÈ%ÔPÐPô
 ×'Ñ'¬¯©«Ó8ˆØ�r‘˜1˜Q�ˆØ×Ñ×#Ñ# FÓ+ˆØˆ;Ü—(‘(˜5 %Ó(ˆCØ(+ˆD×Ñ˜fÑ%Øˆ
r   N)r   r   r    r!   r   r9   r"   r   r   r@   r@   M   s   „ ñò"Bó
r   r@   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y)ÚFakeCUDAAtomicc                 ób   — t         5  ||   }||xx   |z  cc<   d d d «       |S # 1 sw Y   S xY wr   )Úaddlock©r   r9   ÚindexÚvalÚolds        r   ÚaddzFakeCUDAAtomic.add‹   ó:   € Üñ 	 Ø˜‘,ˆCØ�%‹L˜CÑ‹L÷	 ð ˆ
÷	 ð ˆ
úó   ‡$¤.c                 ób   — t         5  ||   }||xx   |z  cc<   d d d «       |S # 1 sw Y   S xY wr   )Úsublockr\   s        r   ÚsubzFakeCUDAAtomic.sub‘   ra   rb   c                 ób   — t         5  ||   }||xx   |z  cc<   d d d «       |S # 1 sw Y   S xY wr   )Úandlockr\   s        r   Úand_zFakeCUDAAtomic.and_—   ra   rb   c                 ób   — t         5  ||   }||xx   |z  cc<   d d d «       |S # 1 sw Y   S xY wr   )Úorlockr\   s        r   Úor_zFakeCUDAAtomic.or_�   s:   € Üñ 	 Ø˜‘,ˆCØ�%‹L˜CÑ‹L÷	 ð ˆ
÷	 ð ˆ
úrb   c                 ób   — t         5  ||   }||xx   |z  cc<   d d d «       |S # 1 sw Y   S xY wr   )Úxorlockr\   s        r   ÚxorzFakeCUDAAtomic.xor£   ra   rb   c                 óx   — t         5  ||   }||k\  rd||<   n||xx   dz  cc<   d d d «       |S # 1 sw Y   S xY w©Nr   r   )Úinclockr\   s        r   ÚinczFakeCUDAAtomic.inc©   sJ   € Üñ 	"Ø˜‘,ˆCØ�cŠzØ ��e’à�e“ Ñ!“÷	"ð ˆ
÷	"ð ˆ
ús   ‡/¯9c                 ó‚   — t         5  ||   }|dk(  s||kD  r|||<   n||xx   dz  cc<   d d d «       |S # 1 sw Y   S xY wrp   )Údeclockr\   s        r   ÚdeczFakeCUDAAtomic.dec²   sP   € Üñ 	"Ø˜‘,ˆCØ�q’˜c CšiØ"��e’à�e“ Ñ!“÷	"ð ˆ
÷	"ð ˆ
úó   ‡#4´>c                 óR   — t         5  ||   }|||<   d d d «       |S # 1 sw Y   S xY wr   )Úexchlockr\   s        r   ÚexchzFakeCUDAAtomic.exch»   s5   € Üñ 	Ø˜‘,ˆCØˆE�%‰L÷	ð ˆ
÷	ð ˆ
ús   ‡œ&c                 óf   — t         5  ||   }t        ||«      ||<   d d d «       |S # 1 sw Y   S xY wr   )ÚmaxlockÚmaxr\   s        r   r|   zFakeCUDAAtomic.maxÁ   ó;   € Üñ 	)Ø˜‘,ˆCÜ˜s C›=ˆE�%‰L÷	)ð ˆ
÷	)ð ˆ
úó   ‡&¦0c                 óf   — t         5  ||   }t        ||«      ||<   d d d «       |S # 1 sw Y   S xY wr   )ÚminlockÚminr\   s        r   r�   zFakeCUDAAtomic.minÇ   r}   r~   c                 ó‚   — t         5  ||   }t        j                  ||   |g«      ||<   d d d «       |S # 1 sw Y   S xY wr   )r{   r5   Únanmaxr\   s        r   rƒ   zFakeCUDAAtomic.nanmaxÍ   óG   € Üñ 	:Ø˜‘,ˆCÜŸ9™9 e¨E¡l°CÐ%8Ó9ˆE�%‰L÷	:ð ˆ
÷	:ð ˆ
úrv   c                 ó‚   — t         5  ||   }t        j                  ||   |g«      ||<   d d d «       |S # 1 sw Y   S xY wr   )r€   r5   Únanminr\   s        r   r†   zFakeCUDAAtomic.nanminÓ   r„   rv   c                 óz   — t         5  d|j                  z  }||   }||k(  r|||<   |cd d d «       S # 1 sw Y   y xY w)N)r   )Úcompare_and_swaplockÚndim)r   r9   r_   r^   r]   Úloadeds         r   Úcompare_and_swapzFakeCUDAAtomic.compare_and_swapÙ   sD   € Ü!ñ 	Ø˜5Ÿ:™:Ñ%ˆEØ˜5‘\ˆFØ˜Š}Ø"��e‘Ø÷	÷ 	ò 	ús   ‡ 1±:c                 ó\   — t         5  ||   }||k(  r|||<   |cd d d «       S # 1 sw Y   y xY wr   )Úcaslock)r   r9   r]   r_   r^   rŠ   s         r   ÚcaszFakeCUDAAtomic.casá   s6   € Üñ 	Ø˜5‘\ˆFØ˜Š}Ø"��e‘Ø÷		÷ 	ò 	ús   ‡"¢+N)r   r   r    r`   re   rh   rk   rn   rr   ru   ry   r|   r�   rƒ   r†   r‹   rŽ   r"   r   r   rY   rY   Š   sH   „ òòòòòòòòòòòòòó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 y) Ú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ˆr   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   ÚsinÚfloat16©r   r   s     r   ÚhsinzFakeCUDAFp16.hsinÿ   ó   € Ü�v‰v�aœrŸz™zÔ*Ð*r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Úcosrª   r«   s     r   ÚhcoszFakeCUDAFp16.hcos  r­   r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Úlogrª   r«   s     r   ÚhlogzFakeCUDAFp16.hlog  r­   r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Úlog2rª   r«   s     r   Úhlog2zFakeCUDAFp16.hlog2  ó   € Ü�w‰w�q¤§
¡
Ô+Ð+r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Úlog10rª   r«   s     r   Úhlog10zFakeCUDAFp16.hlog10  ó   € Ü�x‰x˜¤§¡Ô,Ð,r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Úexprª   r«   s     r   ÚhexpzFakeCUDAFp16.hexp  r­   r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Úexp2rª   r«   s     r   Úhexp2zFakeCUDAFp16.hexp2  r·   r   c                 ó2   — t        j                  d|z  «      S )Né
   ©r5   rª   r«   s     r   Úhexp10zFakeCUDAFp16.hexp10  s   € Ü�z‰z˜" ™'Ó"Ð"r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Úsqrtrª   r«   s     r   ÚhsqrtzFakeCUDAFp16.hsqrt  r·   r   c                 ó2   — t        j                  |dz  «      S )Ng      à¿rÄ   r«   s     r   ÚhrsqrtzFakeCUDAFp16.hrsqrt  s   € Ü�z‰z˜!˜t™)Ó$Ð$r   c                 óL   — t        j                  |t         j                  ¬«      S rB   ©r5   Úceilrª   r«   s     r   ÚhceilzFakeCUDAFp16.hceil  r·   r   c                 óL   — t        j                  |t         j                  ¬«      S rB   rÌ   r«   s     r   ÚhfloorzFakeCUDAFp16.hfloor   r·   r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Ú
reciprocalrª   r«   s     r   ÚhrcpzFakeCUDAFp16.hrcp#  s   € Ü�}‰}˜Q¤b§j¡jÔ1Ð1r   c                 óL   — t        j                  |t         j                  ¬«      S rB   )r5   Útruncrª   r«   s     r   ÚhtrunczFakeCUDAFp16.htrunc&  r»   r   c                 óL   — 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                 ó   — ||k7  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                 ó   — ||kD  S r   r"   r’   s      r   ÚhgtzFakeCUDAFp16.hgt5  r–   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Ð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ò-ò,òòòòòòòór   r�   c                   ó  — 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y)Ú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 ]0  \  }}t#        | ||«       |j$                  D ]  }t#        | ||«       Œ Œ2 y 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ˆŒÜ% nÓ5ˆŒÜ#“oˆŒÜ%Ó'ˆŒÜ!“^ˆŒ
ô '×,Ñ,Ó.ò 	+‰JˆD�$Ü�D˜$ Ô%ØŸ™ò +�Ü˜˜e TÕ*ñ+ñ	+r   c                 ó   — | j                   S r   )rñ   r   s    r   ÚcgzFakeCUDAModule.cgc  s   € à�x‰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‰zÐr   c                 ó>   — t        j                  «       j                  S r   )r'   r(   Ú	threadIdxr   s    r   r  zFakeCUDAModule.threadIdx{  s   € ä×'Ñ'Ó)×3Ñ3Ð3r   c                 ó>   — t        j                  «       j                  S r   )r'   r(   ÚblockIdxr   s    r   r  zFakeCUDAModule.blockIdx  s   € ä×'Ñ'Ó)×2Ñ2Ð2r   c                  ó   — y©Né    r"   r   s    r   ÚwarpsizezFakeCUDAModule.warpsizeƒ  s   € àr   c                 óD   — t        j                  «       j                  dz  S r  )r'   r(   Ú	thread_idr   s    r   ÚlaneidzFakeCUDAModule.laneid‡  s   € ä×'Ñ'Ó)×3Ñ3°bÑ8Ð8r   c                 óH   — t        j                  «       j                  «        y r   r&   r   s    r   r)   zFakeCUDAModule.syncthreads‹  s   € Ü× Ñ Ó"×.Ñ.Õ0r   c                  ó   — y r   r"   r   s    r   ÚthreadfencezFakeCUDAModule.threadfenceŽ  ó   € àr   c                  ó   — y r   r"   r   s    r   Úthreadfence_blockz FakeCUDAModule.threadfence_block’  r  r   c                  ó   — y r   r"   r   s    r   Úthreadfence_systemz!FakeCUDAModule.threadfence_system–  r  r   c                 óH   — t        j                  «       j                  |«      S r   )r'   r(   Úsyncthreads_count©r   r^   s     r   r!  z FakeCUDAModule.syncthreads_countš  s   € Ü×'Ñ'Ó)×;Ñ;¸CÓ@Ð@r   c                 óH   — t        j                  «       j                  |«      S r   )r'   r(   Úsyncthreads_andr"  s     r   r$  zFakeCUDAModule.syncthreads_and�  s   € Ü×'Ñ'Ó)×9Ñ9¸#Ó>Ð>r   c                 óH   — t        j                  «       j                  |«      S r   )r'   r(   Úsyncthreads_orr"  s     r   r&  zFakeCUDAModule.syncthreads_or   s   € Ü×'Ñ'Ó)×8Ñ8¸Ó=Ð=r   c                 ó6   — t        |«      j                  d«      S )NÚ1)ÚbinrJ   r"  s     r   ÚpopczFakeCUDAModule.popc£  s   € Ü�3‹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                 óD   — t        dj                  |«      d d d…   d«      S )Nú{:032b}éÿÿÿÿrL   )ÚintÚformatr"  s     r   ÚbrevzFakeCUDAModule.brev¬  s#   € Ü�9×#Ñ# CÓ(©¨2¨Ñ.°Ó2Ð2r   c                 óp   — dj                  |«      }t        |«      t        |j                  d«      «      z
  S )Nr0  Ú0)r3  ÚlenÚlstrip)r   r^   Úss      r   ÚclzzFakeCUDAModule.clz¯  s.   € Ø×Ñ˜SÓ!ˆÜ�1‹vœ˜AŸH™H S›MÓ*Ñ*Ð*r   c                 ó€   — dj                  |«      }t        |«      t        |j                  d«      «      z
  dz   dz  }|S )Nr0  r6  r   é!   )r3  r7  Úrstrip)r   r^   r9  Úrs       r   ÚffszFakeCUDAModule.ffs³  s>   € ð ×Ñ˜SÓ!ˆÜ�‹V”c˜!Ÿ(™( 3›-Ó(Ñ(¨1Ñ,°Ñ2ˆØˆr   c                 ó   — |r|S |S r   r"   rž   s       r   ÚselpzFakeCUDAModule.selp½  s   € Ùˆ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‰mˆØ�n‰nˆØ�E‰E�D—F‘F‰N˜SŸU™UÑ"ˆØ�Š6ØˆHØ�E‰E�D—F‘F‰N˜SŸU™UÑ"ˆØ�Š6Ø�q�6ˆMØ�E‰E�D—F‘F‰N˜SŸU™UÑ"ˆØ�Š6Ø�q˜!�9ÐäÐGÈ!ÑKÓ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   rC  z,Global grid has 1-3 dimensions. %d requested)rð   rï   r   r   r   rD  )r   rE  rF  Úgdimr   r   r   s          r   ÚgridsizezFakeCUDAModule.gridsizeÐ  s‡   € Ø�}‰}ˆØ�|‰|ˆØ�F‰F�T—V‘V‰OˆØ�Š6ØˆHØ�F‰F�T—V‘V‰OˆØ�Š6Ø�q�6ˆMØ�F‰F�T—V‘V‰OˆØ�Š6Ø�q˜!�9ÐäÐIÈAÑMÓ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.  r4  r:  r?  rA  rI  rL  r"   r   r   rí   rí   E  s'  „ ñò+ð& ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñó ðð ñ4ó ð4ð ñ3ó ð3ð ñó ðð ñ9ó ð9ò1òòòòAò?ò>ò#òòò3ò+òòòMó Or   rí   c              #   ó6  ‡‡K  — ddl mŠ | j                  }t        ˆfd„|j	                  «       D «       «      }t        ˆfd„|j	                  «       D «       «      }|j                  |«       	 d –— |j                  |«       y # |j                  |«       w xY w­w)Nr   )Úcudac              3   ó6   •K  — | ]  \  }}|‰u sŒ||f–— Œ y ­wr   r"   )Ú.0ÚkÚvrO  s      €r   ú	<genexpr>z&swapped_cuda_module.<locals>.<genexpr>æ  s   øè ø€ ÒA™4˜1˜a°q¸D²y��A”ÑAùs   ƒ�	c              3   ó,   •K  — | ]  \  }}|‰f–— Œ y ­wr   r"   )rQ  rR  rS  Úfake_cuda_modules      €r   rT  z&swapped_cuda_module.<locals>.<genexpr>è  s   øè ø€ Ò?©$¨!¨Q�Ð$Ô%Ñ?ùs   ƒ)ÚnumbarO  Ú__globals__Údictr÷   Úupdate)ÚfnrV  Úfn_globsÚorigÚreplrO  s    `   @r   Úswapped_cuda_moduler_  à  sn   ùè ø€ åà�~‰~€HäÓA 8§>¡>Ó#3ÔAÓA€DäÓ?°$·*±*³,Ô?Ó?€Dà‡O�O�DÔðÛð 	�‰˜Õøˆ�‰˜Õüs   „A(BÁ-B Á1BÂBÂB)%r!   Ú
contextlibr   rR   r'   rP   Ú
numba.corer   Únumpyr5   Únumba.npr   r   Úobjectr	   r$   r,   r0   r;   r@   ÚLockr[   rd   rg   rj   rm   r{   r€   rˆ   r�   rq   rt   rx   rY   r�   rí   r_  r"   r   r   ú<module>rf     sW  ðñõ
 &Û 
Û Û Ý Û å "å &ôˆ6ô ÷*	1ñ 	1÷ñ ô&�Fô &ô�Fô ô,�Vô ,ð^ ˆ)�.‰.Ó
€Ø
ˆ)�.‰.Ó
€Ø
ˆ)�.‰.Ó
€Ø	ˆ�‰Ó	€Ø
ˆ)�.‰.Ó
€Ø
ˆ)�.‰.Ó
€Ø
ˆ)�.‰.Ó
€Ø%�y—~‘~Ó'Ð Ø
ˆ)�.‰.Ó
€Ø
ˆ)�.‰.Ó
€Ø
ˆ)�.‰.Ó
€Øˆ9�>‰>Ó€ô\�Vô \ô~Y�6ô YôxXO�Vô XOðv ñó ñr   