U
    ?h                     @   s~   d dl mZ d dlmZmZmZ d dlmZmZ d dl	m
Z
 d dlmZ d dlmZ dZdd	d	g dd
d	d	fddZdd ZdS )    )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                    s$  rt jrtddr$tdddk	rDtd}	t|	ddk	rdtd}	t|	ddk	rtd}	t|	dkrt jndd	d
g rƈrd}	tt	|	 rވrd}	tt	|	 rdrt
dt| r| g	d
nt| tr"| 	d	
nd		dk	rnt jrJfdd}
|
S  	
fdd}|S | dkrt jrfdd}n fdd}|S t jrt| dS  }|d< |d< |d< |d< |d< |d< |d
< t| |d} r|  |S dS )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 simulatorZ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                    s   t |  dS Ndevicer   r
   funcr    G/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/decorators.py
jitwrapperg   s    zjit.<locals>.jitwrapperc              
      s     }|d< |d< |d< |d< |d< |d< |d< t| |d} rX|  	D ]j}t|\}}|rs|tjkrtd	rd
dlm	} |
| ||| W 5 Q R X q\|| q\
|_|  |S )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   Zregister_dispatcherZcompile_devicecompileZ_specializedZdisable_compile)r   r   dispsigr   r   r   cacher   r   r   r   kwsr   r   r   Z
signaturesZspecializedr   r   _jitk   s.    zjit.<locals>._jitc                    s   t |  dS r   r
   r   r   r   r   autojitwrapper   s    zjit.<locals>.autojitwrapperc              	      s   t | f dS )N)r   r   r   r   r   r*   )jitr   )r*   r   r   r+   r   r   r   r   r   r-      s      r   r   r   r   r   r   )r   ZENABLE_CUDASIMNotImplementedErrorget_msg_deprecated_signature_argformatr   ZCUDA_DEBUGINFO_DEFAULTr   r   
ValueErrorr   Zis_signature
isinstancelistr   r    r	   r!   )Zfunc_or_sigr   inliner   r   r   r   r*   r+   msgr   r,   r-   r   r'   r   r)   r   r.      st    +





 !
r.   c                 C   s.   t |\}}|dkr"d}t|t| ||S )a  
    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.
    Nz4Return type must be provided for device declarations)r   r"   r$   r   )namer(   r   r   r7   r   r   r   declare_device   s
    	r9   )warningsr   r%   r   r   r   Znumba.core.errorsr   r   Znumba.cuda.compilerr   Znumba.cuda.dispatcherr	   Znumba.cuda.simulator.kernelr   r1   r.   r9   r   r   r   r   <module>   s   
  
 "