
    im0                        S r SSKJr  SSKrSSKrSSKrSSKJr  SSKr	SSK
Jr  SSKJr   " S S	\5      r " S
 S5      r " S S5      r " S S\5      r " S S\5      r " S S\5      r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r\R(                  " 5       r  " S S\5      r! " S S\5      r" " S S\5      r#\S 5       r$g)zf
Implements the cuda module as called from within an executing kernel
(@cuda.jit-decorated function).
    )contextmanagerN)types)numpy_support   )vector_typesc                   0    \ rS rSrSrS rS rS rS rSr	g)	Dim3   z3
Used to implement thread/block indices/dimensions
c                 (    Xl         X l        X0l        g Nxyz)selfr   r   r   s       Y/var/www/html/trading/venv/lib/python3.13/site-packages/numba/cuda/simulator/kernelapi.py__init__Dim3.__init__   s        c                 \    SU R                   < SU R                  < SU R                  < S3$ )N(, )r   r   s    r   __str__Dim3.__str__   s    !%88r   c                 \    SU R                   < SU R                  < SU R                  < S3$ )NzDim3(r   r   r   r   s    r   __repr__Dim3.__repr__   s    %)VVTVVTVV<<r   c              #   `   #    U R                   v   U R                  v   U R                  v   g 7fr   r   r   s    r   __iter__Dim3.__iter__!   s      ffffffs   ,.r   N)
__name__
__module____qualname____firstlineno____doc__r   r   r   r!   __static_attributes__ r   r   r	   r	      s    
9=r   r	   c                       \ rS rSrSrS rSrg)	GridGroup'   z#
Used to implement the grid group.
c                 J    [         R                  " 5       R                  5         g r   	threadingcurrent_threadsyncthreadsr   s    r   syncGridGroup.sync,   s     	  "..0r   r)   N)r#   r$   r%   r&   r'   r2   r(   r)   r   r   r+   r+   '   s    1r   r+   c                       \ rS rSrSrS rSrg)
FakeCUDACg3   z
CUDA Cooperative Groups
c                     [        5       $ r   )r+   r   s    r   	this_gridFakeCUDACg.this_grid7   s
    {r   r)   N)r#   r$   r%   r&   r'   r8   r(   r)   r   r   r5   r5   3   s    r   r5   c                       \ rS rSrSrS rSrg)FakeCUDALocal;   z
CUDA Local arrays
c                     [        U[        R                  5      (       a  [        R                  " U5      n[
        R                  " X5      $ r   )
isinstancer   Typer   as_dtypenpempty)r   shapedtypes      r   arrayFakeCUDALocal.array?   s2    eUZZ((!**51Exx%%r   r)   N)r#   r$   r%   r&   r'   rE   r(   r)   r   r   r;   r;   ;   s    &r   r;   c                       \ rS rSrSrS rSrg)FakeCUDAConstE   z
CUDA Const arrays
c                     U$ r   r)   )r   arys     r   
array_likeFakeCUDAConst.array_likeI   s    
r   r)   N)r#   r$   r%   r&   r'   rL   r(   r)   r   r   rH   rH   E   s    r   rH   c                   $    \ rS rSrSrS rS rSrg)FakeCUDASharedM   a  
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                 n    0 U l         Xl        [        R                  " U[        R                  S9U l        g N)rD   )_allocations_dynshared_sizerA   zerosbyte
_dynshared)r   dynshared_sizes     r   r   FakeCUDAShared.__init___   s&    -((>Ar   c                    [        U[        R                  5      (       a  [        R                  " U5      nUS:X  aB  U R
                  UR                  -  n[        R                  " U R                  R                  X#S9$ [        R                  " [        R                  " 5       5      nUS   SS nU R                  R!                  U5      nUc$  [        R"                  " X5      nX`R                  U'   U$ )Nr   )rD   count   )r>   r   r?   r   r@   rT   itemsizerA   
frombufferrW   data	tracebackextract_stacksys	_getframerS   getrB   )r   rC   rD   r[   stackcallerress          r   rE   FakeCUDAShared.arrayd   s    eUZZ((!**51E A: ((ENN:E==!5!5UPP
 ''8r1Q##F+;((5(C(+f%
r   )rS   rW   rT   N)r#   r$   r%   r&   r'   r   rE   r(   r)   r   r   rO   rO   M   s    "B
r   rO   c                   h    \ rS rSrS rS rS rS rS rS r	S r
S	 rS
 rS rS rS rS rS rSrg)FakeCUDAAtomic   c                 h    [            X   nX==   U-  ss'   S S S 5        U$ ! , (       d  f       W$ = fr   )addlockr   rE   indexvalolds        r   addFakeCUDAAtomic.add   3    ,CLCL  
 W 
   "
1c                 h    [            X   nX==   U-  ss'   S S S 5        U$ ! , (       d  f       W$ = fr   )sublockro   s        r   subFakeCUDAAtomic.sub   ru   rv   c                 h    [            X   nX==   U-  ss'   S S S 5        U$ ! , (       d  f       W$ = fr   )andlockro   s        r   and_FakeCUDAAtomic.and_   ru   rv   c                 h    [            X   nX==   U-  ss'   S S S 5        U$ ! , (       d  f       W$ = fr   )orlockro   s        r   or_FakeCUDAAtomic.or_   s3    ,CLCL  
 V 
rv   c                 h    [            X   nX==   U-  ss'   S S S 5        U$ ! , (       d  f       W$ = fr   )xorlockro   s        r   xorFakeCUDAAtomic.xor   ru   rv   c                 |    [            X   nXC:  a  SX'   OX==   S-  ss'   S S S 5        U$ ! , (       d  f       W$ = fNr   r   )inclockro   s        r   incFakeCUDAAtomic.inc   s?    ,Cz !  
 W 
s   ,
;c                     [            X   nUS:X  d  XC:  a  X1U'   OX==   S-  ss'   S S S 5        U$ ! , (       d  f       W$ = fr   )declockro   s        r   decFakeCUDAAtomic.dec   sE    ,Cqci"e!  
 W 
	   !2
Ac                 X    [            X   nX1U'   S S S 5        U$ ! , (       d  f       W$ = fr   )exchlockro   s        r   exchFakeCUDAAtomic.exch   s.    ,C%L  
 X 
s   	
)c                 j    [            X   n[        XC5      X'   S S S 5        U$ ! , (       d  f       W$ = fr   )maxlockmaxro   s        r   r   FakeCUDAAtomic.max   2    ,Cs=EL  
 W 
   #
2c                 j    [            X   n[        XC5      X'   S S S 5        U$ ! , (       d  f       W$ = fr   )minlockminro   s        r   r   FakeCUDAAtomic.min   r   r   c                     [            X   n[        R                  " X   U/5      X'   S S S 5        U$ ! , (       d  f       W$ = fr   )r   rA   nanmaxro   s        r   r   FakeCUDAAtomic.nanmax   >    ,C99elC%89EL  
 W 
r   c                     [            X   n[        R                  " X   U/5      X'   S S S 5        U$ ! , (       d  f       W$ = fr   )r   rA   nanminro   s        r   r   FakeCUDAAtomic.nanmin   r   r   c                     [            SUR                  -  nX   nXR:X  a  X1U'   UsS S S 5        $ ! , (       d  f       g = f)N)r   )compare_and_swaplockndim)r   rE   rr   rq   rp   loadeds         r   compare_and_swapFakeCUDAAtomic.compare_and_swap   s5    !5::%E\F}"e "!!s   /
=c                 b    [            X   nXS:X  a  XAU'   UsS S S 5        $ ! , (       d  f       g = fr   )caslock)r   rE   rp   rr   rq   r   s         r   casFakeCUDAAtomic.cas   s%    \F}"e	 WWs    
.r)   N)r#   r$   r%   r&   rs   ry   r}   r   r   r   r   r   r   r   r   r   r   r   r(   r)   r   r   rk   rk      sH    r   rk   c                       \ rS rSrS rS rS rS rS rS r	S r
S	 rS
 rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS r S r!S r"g!)"FakeCUDAFp16   c                 
    X-   $ r   r)   r   abs      r   haddFakeCUDAFp16.hadd   	    ur   c                 
    X-
  $ r   r)   r   s      r   hsubFakeCUDAFp16.hsub   r   r   c                 
    X-  $ r   r)   r   s      r   hmulFakeCUDAFp16.hmul   r   r   c                 
    X-  $ r   r)   r   s      r   hdivFakeCUDAFp16.hdiv   r   r   c                     X-  U-   $ r   r)   r   r   r   cs       r   hfmaFakeCUDAFp16.hfma       uqyr   c                     U* $ r   r)   r   r   s     r   hnegFakeCUDAFp16.hneg   s	    r	r   c                     [        U5      $ r   )absr   s     r   habsFakeCUDAFp16.habs   s    1vr   c                 H    [         R                  " U[         R                  S9$ rR   )rA   sinfloat16r   r   s     r   hsinFakeCUDAFp16.hsin       vvarzz**r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   cosr   r   s     r   hcosFakeCUDAFp16.hcos  r   r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   logr   r   s     r   hlogFakeCUDAFp16.hlog  r   r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   log2r   r   s     r   hlog2FakeCUDAFp16.hlog2      wwq

++r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   log10r   r   s     r   hlog10FakeCUDAFp16.hlog10      xx,,r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   expr   r   s     r   hexpFakeCUDAFp16.hexp  r   r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   exp2r   r   s     r   hexp2FakeCUDAFp16.hexp2  r   r   c                 4    [         R                  " SU-  5      $ )N
   rA   r   r   s     r   hexp10FakeCUDAFp16.hexp10  s    zz"'""r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   sqrtr   r   s     r   hsqrtFakeCUDAFp16.hsqrt  r   r   c                 4    [         R                  " US-  5      $ )Ng      r   r   s     r   hrsqrtFakeCUDAFp16.hrsqrt  s    zz!t)$$r   c                 H    [         R                  " U[         R                  S9$ rR   rA   ceilr   r   s     r   hceilFakeCUDAFp16.hceil  r   r   c                 H    [         R                  " U[         R                  S9$ rR   r   r   s     r   hfloorFakeCUDAFp16.hfloor   r   r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   
reciprocalr   r   s     r   hrcpFakeCUDAFp16.hrcp#  s    }}Qbjj11r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   truncr   r   s     r   htruncFakeCUDAFp16.htrunc&  r   r   c                 H    [         R                  " U[         R                  S9$ rR   )rA   rintr   r   s     r   hrintFakeCUDAFp16.hrint)  r   r   c                 
    X:H  $ r   r)   r   s      r   heqFakeCUDAFp16.heq,  	    vr   c                 
    X:g  $ r   r)   r   s      r   hneFakeCUDAFp16.hne/  r  r   c                 
    X:  $ r   r)   r   s      r   hgeFakeCUDAFp16.hge2  r  r   c                 
    X:  $ r   r)   r   s      r   hgtFakeCUDAFp16.hgt5  r   r   c                 
    X:*  $ r   r)   r   s      r   hleFakeCUDAFp16.hle8  r  r   c                 
    X:  $ r   r)   r   s      r   hltFakeCUDAFp16.hlt;  r   r   c                     [        X5      $ r   )r   r   s      r   hmaxFakeCUDAFp16.hmax>      1yr   c                     [        X5      $ r   )r   r   s      r   hminFakeCUDAFp16.hminA  r(  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   r   r      s    +++,-+,#,%,,2-,r   r   c                      \ rS rSrSrS r\S 5       r\S 5       r\S 5       r	\S 5       r
\S 5       r\S	 5       r\S
 5       r\S 5       r\S 5       r\S 5       rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS r Sr!g) FakeCUDAModuleiE  a  
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                    [        U6 U l        [        U6 U l        [        5       U l        [        5       U l        [        U5      U l        [        5       U l
        [        5       U l        [        5       U l        [        R                   " 5        H0  u  pE[#        XU5        UR$                   H  n[#        XU5        M     M2     g r   )r	   gridDimblockDimr5   _cgr;   _localrO   _sharedrH   _constrk   _atomicr   _fp16r   itemssetattraliases)r   grid_dim	block_dimrX   namesvtyaliass          r   r   FakeCUDAModule.__init__P  s    Xi(<#o%n5#o%'!^
 ',,.JDD%T* & /r   c                     U R                   $ r   )r1  r   s    r   cgFakeCUDAModule.cgc  s    xxr   c                     U R                   $ r   )r2  r   s    r   localFakeCUDAModule.localg      {{r   c                     U R                   $ r   )r3  r   s    r   sharedFakeCUDAModule.sharedk      ||r   c                     U R                   $ r   )r4  r   s    r   constFakeCUDAModule.consto  rF  r   c                     U R                   $ r   )r5  r   s    r   atomicFakeCUDAModule.atomics  rJ  r   c                     U R                   $ r   )r6  r   s    r   fp16FakeCUDAModule.fp16w  s    zzr   c                 @    [         R                  " 5       R                  $ r   )r/   r0   	threadIdxr   s    r   rU  FakeCUDAModule.threadIdx{  s    '')333r   c                 @    [         R                  " 5       R                  $ r   )r/   r0   blockIdxr   s    r   rX  FakeCUDAModule.blockIdx  s    '')222r   c                     gN    r)   r   s    r   warpsizeFakeCUDAModule.warpsize  s    r   c                 F    [         R                  " 5       R                  S-  $ r[  )r/   r0   	thread_idr   s    r   laneidFakeCUDAModule.laneid  s    '')33b88r   c                 J    [         R                  " 5       R                  5         g r   r.   r   s    r   r1   FakeCUDAModule.syncthreads  s      "..0r   c                     g r   r)   r   s    r   threadfenceFakeCUDAModule.threadfence      r   c                     g r   r)   r   s    r   threadfence_block FakeCUDAModule.threadfence_block  rh  r   c                     g r   r)   r   s    r   threadfence_system!FakeCUDAModule.threadfence_system  rh  r   c                 J    [         R                  " 5       R                  U5      $ r   )r/   r0   syncthreads_countr   rq   s     r   rp   FakeCUDAModule.syncthreads_count  s    '');;C@@r   c                 J    [         R                  " 5       R                  U5      $ r   )r/   r0   syncthreads_andrq  s     r   rt  FakeCUDAModule.syncthreads_and  s    '')99#>>r   c                 J    [         R                  " 5       R                  U5      $ r   )r/   r0   syncthreads_orrq  s     r   rw  FakeCUDAModule.syncthreads_or  s    '')88==r   c                 6    [        U5      R                  S5      $ )N1)binr[   rq  s     r   popcFakeCUDAModule.popc  s    3x~~c""r   c                     X-  U-   $ r   r)   r   s       r   fmaFakeCUDAModule.fma  r   r   c                     US-  $ )NgUUUUUU?r)   r   s     r   cbrtFakeCUDAModule.cbrt  s    U|r   c                 D    [        SR                  U5      S S S2   S5      $ )N{:032b}r]   )intformatrq  s     r   brevFakeCUDAModule.brev  s#    9##C(2.22r   c                 p    SR                  U5      n[        U5      [        UR                  S5      5      -
  $ )Nr  0)r  lenlstrip)r   rq   ss      r   clzFakeCUDAModule.clz  s.    S!1vAHHSM***r   c                     SR                  U5      n[        U5      [        UR                  S5      5      -
  S-   S-  nU$ )Nr  r  r   !   )r  r  rstrip)r   rq   r  rs       r   ffsFakeCUDAModule.ffs  s>     S!Vc!((3-((1,2r   c                     U(       a  U$ U$ r   r)   r   s       r   selpFakeCUDAModule.selp  s    q1r   c                    U R                   nU R                  nU R                  nUR                  UR                  -  UR                  -   nUS:X  a  U$ UR                  UR                  -  UR                  -   nUS:X  a  XV4$ UR
                  UR
                  -  UR
                  -   nUS:X  a  XVU4$ [        SU-  5      e)Nr   r]      z*Global ID has 1-3 dimensions. %d requested)r0  rX  rU  r   r   r   RuntimeError)r   nbdimbidtidr   r   r   s           r   gridFakeCUDAModule.grid  s    }}mmnnEEDFFNSUU"6HEEDFFNSUU"66MEEDFFNSUU"6!9G!KLLr   c                    U R                   nU R                  nUR                  UR                  -  nUS:X  a  U$ UR                  UR                  -  nUS:X  a  XE4$ UR                  UR                  -  nUS:X  a  XEU4$ [        SU-  5      e)Nr   r]   r  z,Global grid has 1-3 dimensions. %d requested)r0  r/  r   r   r   r  )r   r  r  gdimr   r   r   s          r   gridsizeFakeCUDAModule.gridsize  s    }}||FFTVVO6HFFTVVO66MFFTVVO6!9IAMNNr   )r5  r1  r4  r6  r2  r3  r0  r/  N)"r#   r$   r%   r&   r'   r   propertyrA  rD  rH  rL  rO  rR  rU  rX  r]  ra  r1   rf  rj  rm  rp  rt  rw  r|  r  r  r  r  r  r  r  r  r(   r)   r   r   r-  r-  E  s'   +&             4 4 3 3   9 91A?>#3+M Or   r-  c              #   :  ^^#    SSK Jm  U R                  n[        U4S jUR	                  5        5       5      n[        U4S jUR	                  5        5       5      nUR                  U5         S v   UR                  U5        g ! UR                  U5        f = f7f)Nr   )cudac              3   <   >#    U  H  u  pUTL d  M  X4v   M     g 7fr   r)   ).0kvr  s      r   	<genexpr>&swapped_cuda_module.<locals>.<genexpr>  s     A#341qDy#3s   
c              3   0   >#    U  H  u  pUT4v   M     g 7fr   r)   )r  r  r  fake_cuda_modules      r   r  r    s     ?,$!$%,s   )numbar  __globals__dictr7  update)fnr  fn_globsorigreplr  s    `   @r   swapped_cuda_moduler    sm     ~~HA8>>#3AAD?$**,??DOOD 	s   A*B/B 3BBB)%r'   
contextlibr   rc   r/   ra   
numba.corer   numpyrA   numba.npr   r   objectr	   r+   r5   r;   rH   rO   Lockrn   rx   r|   r   r   r   r   r   r   r   r   r   rk   r   r-  r  r)   r   r   <module>r     s?  
 & 
     " &6 *	1 	1 &F &F ,V ,^ ..

..

..
		
..

..

..
 ~~' 
..

..

..
>>\V \~Y6 YxXOV XOv  r   