o 8ήc]@s|ddlmZddlmZmZmZddlmZmZddl m Z ddl m Z ddl mZdZdd d gdd d fd d Zd dZdS))warn)typesconfigsigutils)DeprecationErrorNumbaInvalidConfigWarning)declare_device_function)CUDADispatcherFakeCUDAKernelz`Deprecated keyword argument `{0}`. Signatures should be passed as the first positional argument.NFTc sr tjr tddrtdddur"td}t|ddur2td}t|ddurBtd}t|durItjndd d grcrcd }tt |rnd rnt d t |rtjrfdd} | St |\} | rs| tjkrtdf dd} | S|durtjrfdd} | Sfdd} | Stjrt|dS} | d<| d<| d <| d<| d<| d <t|| d}r||S)a JIT compile a python function conforming to the CUDA Python specification. If a signature is supplied, then a function is returned that takes a function to compile. :param func_or_sig: A function to JIT compile, or a signature of a function to compile. If a function is supplied, then a :class:`numba.cuda.compiler.AutoJitCUDAKernel` is returned. If a signature is supplied, then a function is returned. The returned function accepts another function, which it will compile and then return a :class:`numba.cuda.compiler.AutoJitCUDAKernel`. .. 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 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 `. :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 CUDAargtypesNrestypebindfastmathF 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.linkz(link keyword invalid for device functionct|dSNdevicerr funcr.jitwrapperz'CUDA kernel must have void return type.cs}|d<|d<|d<|d<|d<|d<t||d}r(|rKdd lm}|| |Wdn1sEwYn|d |_| |S) Ndebugroptrrr targetoptionsr) typeinferT) copyr enable_caching numba.corer register_dispatchercompile_devicecompile _specializeddisable_compile)rrdispr r cacherrrrkwsrrrr_jit`s(     zjit.._jitcrrr rrrrautojitwrapper~szjit..autojitwrappercst|fdS)N)rrrrr+)jitr)r+rrr,rrrrr.s rrrrr)rENABLE_CUDASIMNotImplementedErrorget_msg_deprecated_signature_argformatrCUDA_DEBUGINFO_DEFAULTrr ValueErrorr is_signaturenormalize_signaturervoid TypeErrorr r!r r") func_or_sigrinlinerrrr+r,msgrrr-r.rr)rr*rr/sf *         r/cCs.t|\}}|durd}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)rr8r:r)namesigr rr=rrrdeclare_devices   r@)warningsrr#rrrnumba.core.errorsrrnumba.cuda.compilerrnumba.cuda.dispatcherr numba.cuda.simulator.kernelr r3r/r@rrrrs