o ۾i@s~ddlmZddlmZmZmZddlmZmZddl m Z ddl m Z ddl mZdZdd d gdd d 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 | rorod } tt | rzd rzt dt |r|g d n t |tr| d nd durtjrfdd} | S f dd} | S|durtjrÇfdd} | Sfdd} | Stjrt|dS} | d<| d<| d<| d <| d<| d<| d <t|| d}r||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 ` 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 `. 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 `. :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.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 functionTct|dSNdevicerr funcrI/home/ubuntu/.local/lib/python3.10/site-packages/numba/cuda/decorators.py jitwrappergszjit..jitwrapperc s}|d<|d<|d<|d<|d<|d<|d<t||d}r,| D]?}t|\}}|rDsD|tjkrDtd rhd d lm }| || ||Wdn1sbwYq.| |q. |_ ||S) Ndebuglineinforoptrrr targetoptionsz'CUDA kernel must have void return type.r) typeinfer)copyr enable_cachingrnormalize_signaturervoid TypeError numba.corer!register_dispatchercompile_devicecompile _specializeddisable_compile)rr dispsigr rr! cacherrrrkwsrrr signatures specializedrr_jitks2    zjit.._jitcrrr rrrrautojitwrapperszjit..autojitwrapperc st|fdS)N)rrrrrr0)jitr)r0rrr1rrrrrr5s rrrrrr)rENABLE_CUDASIMNotImplementedErrorget_msg_deprecated_signature_argformatrCUDA_DEBUGINFO_DEFAULTrr ValueErrorr is_signature isinstancelistr r"r r#) func_or_sigrinlinerrrrr0r1msgrr4r5r r-rr/rr6sv +           ! r6cCs.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)rr$r&r)namer.r rrCrrrdeclare_devices   rE)warningsrr'rrrnumba.core.errorsrrnumba.cuda.compilerrnumba.cuda.dispatcherr numba.cuda.simulator.kernelr r:r6rErrrrs       "