
    sh                     |    S SK Jr  S SKJrJrJr  S SKJrJr  S SK	J
r
  S SKJr  S SKJr  SrSS	S	/ SS
S	S	4S jrS rg)    )warn)typesconfigsigutils)DeprecationErrorNumbaInvalidConfigWarning)declare_device_function)CUDADispatcherFakeCUDAKernelz`Deprecated keyword argument `{0}`. Signatures should be passed as the first positional argument.NFTc                 N  ^^^^^^^^^^^ T(       a   [         R                  (       a  [        S5      eTR                  S5      (       a  [        S5      eTR                  S5      b   [        R                  S5      n	[        U	5      eTR                  S5      b   [        R                  S5      n	[        U	5      eTR                  S5      b   [        R                  S5      n	[        U	5      eTc  [         R                  OTmTR                  SS	5      mTR                  S
/ 5      mT(       a  T(       a  Sn	[        [        U	5      5        T(       a  T(       a  Sn	[        [        U	5      5        T(       a!  TR                  S5      (       a  [        S5      e[        R                  " U 5      (       a  U /mSmO[        U [        5      (       a  U mS	mOSmTb0  [         R                  (       a	  UU4S jn
U
$ UUUUUUUUUUU4S jnU$ U c,  [         R                  (       a	  UU4S jnU$ UUUUUUU4S jnU$ [         R                  (       a  [        U TTS9$ TR!                  5       nTUS'   TUS'   TUS'   TUS'   TUS'   TUS'   TUS
'   [#        XS9nT(       a  UR%                  5         U$ )a>  
JIT compile a Python function for CUDA GPUs.

:param func_or_sig: A function to JIT compile, or *signatures* of a
   function to compile. If a function is supplied, then a
   :class:`Dispatcher <numba.cuda.dispatcher.CUDADispatcher>` is returned.
   Otherwise, ``func_or_sig`` may be a signature or a list of signatures,
   and a function is returned. The returned function accepts another
   function, which it will compile and then return a :class:`Dispatcher
   <numba.cuda.dispatcher.CUDADispatcher>`. See :ref:`jit-decorator` for
   more information about passing signatures.

   .. note:: A kernel cannot have any return value.
:param device: Indicates whether this is a device function.
:type device: bool
:param link: A list of files containing PTX or CUDA C/C++ source to link
   with the function
:type link: list
:param debug: If True, check for exceptions thrown when executing the
   kernel. Since this degrades performance, this should only be used for
   debugging purposes. If set to True, then ``opt`` should be set to False.
   Defaults to False.  (The default value can be overridden by setting
   environment variable ``NUMBA_CUDA_DEBUGINFO=1``.)
:param fastmath: When True, enables fastmath optimizations as outlined in
   the :ref:`CUDA Fast Math documentation <cuda-fast-math>`.
:param max_registers: Request that the kernel is limited to using at most
   this number of registers per thread. The limit may not be respected if
   the ABI requires a greater number of registers than that requested.
   Useful for increasing occupancy.
:param opt: Whether to compile from LLVM IR to PTX with optimization
            enabled. When ``True``, ``-opt=3`` is passed to NVVM. When
            ``False``, ``-opt=0`` is passed to NVVM. Defaults to ``True``.
:type opt: bool
:param lineinfo: If True, generate a line mapping between source code and
   assembly code. This enables inspection of the source code in NVIDIA
   profiling tools and correlation with program counter sampling.
:type lineinfo: bool
:param cache: If True, enables the file-based cache for this function.
:type cache: bool
z Cannot link PTX in the simulatorboundscheckz)bounds checking is not supported for CUDAargtypesNrestypebindfastmathF
extensionsz{debug=True with opt=True (the default) is not supported by CUDA. This may result in a crash - set debug=False or opt=False.zdebug and lineinfo are mutually exclusive. Use debug to get full debug info (this disables some optimizations), or lineinfo for line info only with code generation unaffected.linkz(link keyword invalid for device functionTc                    > [        U TTS9$ Ndevicer   r   funcr   r   s    i/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/decorators.py
jitwrapperjit.<locals>.jitwrapperg   s    %d6HMM    c                 2  > TR                  5       nTUS'   TUS'   TUS'   TUS'   TUS'   T	US'   T
US'   [        XS9nT(       a  UR                  5         T H  n[        R                  " U5      u  pEU(       a&  T	(       d  U[
        R                  :w  a  [        S	5      eT	(       a3  S
SKJ	n  UR                  U5         UR                  XE5        S S S 5        M  UR                  U5        M     TUl        UR                  5         U$ ! , (       d  f       M  = f)Ndebuglineinfor   optr   r   r   targetoptionsz'CUDA kernel must have void return type.r   )	typeinfer)copyr
   enable_cachingr   normalize_signaturer   void	TypeError
numba.corer%   register_dispatchercompile_devicecompile_specializeddisable_compile)r   r$   dispsigr   r   r%   cacher    r   r   r   kwsr!   r   r"   
signaturesspecializeds          r   _jitjit.<locals>._jitk   s   HHJM%*M'"(0M*%$(M&!#&M% (0M*%&,M(#*4M,'!$DD##%!$,$@$@$E!6g.C#$MNN4"66t<++H> =< LL* " !,D  "K =<s   =D
D	c                    > [        U TTS9$ r   r   r   s    r   autojitwrapperjit.<locals>.autojitwrapper   s    )$v3;= =r   c           
      *   > [        U 4TTTTTTS.TD6$ )N)r   r    r"   r!   r   r3   )jit)r   r3   r    r   r4   r!   r   r"   s    r   r:   r;      s0    t QF%S(0t5QLOQ Qr   r   r    r!   r"   r   r#   )r   ENABLE_CUDASIMNotImplementedErrorget_msg_deprecated_signature_argformatr   CUDA_DEBUGINFO_DEFAULTr   r   
ValueErrorr   is_signature
isinstancelistr   r&   r
   r'   )func_or_sigr   inliner   r    r"   r!   r3   r4   msgr   r7   r:   r$   r1   r   r   r5   r6   s    ` ``````      @@@@r   r=   r=      sa   V %%!"DEE
ww}!"MNN
wwz&+22:>s##
wwy%+229=s##
wwv"+226:s##-2]F))Ewwz5)Hr*J2 	&s+,N 	&s+,#''&//CDD[))!]
	K	&	& 

  N	 	B $$= "!	Q Q "! $$%k&/79 9 !$
).g&,4j)'*e$(,f%,4j)*0h'.8l+%kO'')r   c                 j    [         R                  " U5      u  p#Uc  Sn[        U5      e[        XU5      $ )z
Declare the signature of a foreign function. Returns a descriptor that can
be used to call the function from a Python kernel.

:param name: The name of the foreign function.
:type name: str
:param sig: The Numba signature of the function.
z4Return type must be provided for device declarations)r   r(   r*   r	   )namer2   r   r   rJ   s        r   declare_devicerM      s8     !44S9HDn"4(;;r   )warningsr   r+   r   r   r   numba.core.errorsr   r   numba.cuda.compilerr	   numba.cuda.dispatcherr
   numba.cuda.simulator.kernelr   rA   r=   rM    r   r   <module>rT      sB     . . I 7 0 6"8 
 u2T5^B<r   