We are currently solving the following issue within our repository. Here is the issue text:
--- BEGIN ISSUE ---
`compile_ptx()` allows compilation of kernel with non-`void` return type
For example:

```python
from numba import cuda, int32

def f(x, y):
    return x[0] + y[0]

ptx, resty = cuda.compile_ptx(f, (int32[::1], int32[::1]))

print(resty)
# int64
```

compiles, and generates PTX where the function body returns nothing (heavily edited for clarity, but the idea is represented:

```assembly
.visible .entry f (
    // args omitted for brevity
)
{
	ret;
}
```

Usually we check that the kernel has a void return type in the `@cuda.jit` decorator, which is why this slips by in `compile_ptx`. The check should probably be pushed a bit deeper to cover both uses.

(cc @brandonwillard )

--- END ISSUE ---

Below are some code segments, each from a relevant file. One or more of these files may contain bugs.
--- BEGIN FILES ---
Path: `numba/cuda/compiler.py`
Content:
```
1 from numba.core.typing.templates import ConcreteTemplate
2 from numba.core import types, typing, funcdesc, config, compiler
3 from numba.core.compiler import (sanitize_compile_result_entries, CompilerBase,
4                                  DefaultPassBuilder, Flags, Option,
5                                  CompileResult)
6 from numba.core.compiler_lock import global_compiler_lock
7 from numba.core.compiler_machinery import (LoweringPass, AnalysisPass,
8                                            PassManager, register_pass)
9 from numba.core.errors import NumbaInvalidConfigWarning, TypingError
10 from numba.core.typed_passes import (IRLegalization, NativeLowering,
11                                      AnnotateTypes)
12 from warnings import warn
13 from numba.cuda.api import get_current_device
14 
15 
16 def _nvvm_options_type(x):
17     if x is None:
18         return None
19 
20     else:
21         assert isinstance(x, dict)
22         return x
23 
24 
25 class CUDAFlags(Flags):
26     nvvm_options = Option(
27         type=_nvvm_options_type,
28         default=None,
29         doc="NVVM options",
30     )
31 
32 
33 # The CUDACompileResult (CCR) has a specially-defined entry point equal to its
34 # id.  This is because the entry point is used as a key into a dict of
35 # overloads by the base dispatcher. The id of the CCR is the only small and
36 # unique property of a CompileResult in the CUDA target (cf. the CPU target,
37 # which uses its entry_point, which is a pointer value).
38 #
39 # This does feel a little hackish, and there are two ways in which this could
40 # be improved:
41 #
42 # 1. We could change the core of Numba so that each CompileResult has its own
43 #    unique ID that can be used as a key - e.g. a count, similar to the way in
44 #    which types have unique counts.
45 # 2. At some future time when kernel launch uses a compiled function, the entry
46 #    point will no longer need to be a synthetic value, but will instead be a
47 #    pointer to the compiled function as in the CPU target.
48 
49 class CUDACompileResult(CompileResult):
50     @property
51     def entry_point(self):
52         return id(self)
53 
54 
55 def cuda_compile_result(**entries):
56     entries = sanitize_compile_result_entries(entries)
57     return CUDACompileResult(**entries)
58 
59 
60 @register_pass(mutates_CFG=True, analysis_only=False)
61 class CUDABackend(LoweringPass):
62 
63     _name = "cuda_backend"
64 
65     def __init__(self):
66         LoweringPass.__init__(self)
67 
68     def run_pass(self, state):
69         """
70         Back-end: Packages lowering output in a compile result
71         """
72         lowered = state['cr']
73         signature = typing.signature(state.return_type, *state.args)
74 
75         state.cr = cuda_compile_result(
76             typing_context=state.typingctx,
77             target_context=state.targetctx,
78             typing_error=state.status.fail_reason,
79             type_annotation=state.type_annotation,
80             library=state.library,
81             call_helper=lowered.call_helper,
82             signature=signature,
83             fndesc=lowered.fndesc,
84         )
85         return True
86 
87 
88 @register_pass(mutates_CFG=False, analysis_only=False)
89 class CreateLibrary(LoweringPass):
90     """
91     Create a CUDACodeLibrary for the NativeLowering pass to populate. The
92     NativeLowering pass will create a code library if none exists, but we need
93     to set it up with nvvm_options from the flags if they are present.
94     """
95 
96     _name = "create_library"
97 
98     def __init__(self):
99         LoweringPass.__init__(self)
100 
101     def run_pass(self, state):
102         codegen = state.targetctx.codegen()
103         name = state.func_id.func_qualname
104         nvvm_options = state.flags.nvvm_options
105         state.library = codegen.create_library(name, nvvm_options=nvvm_options)
106         # Enable object caching upfront so that the library can be serialized.
107         state.library.enable_object_caching()
108 
109         return True
110 
111 
112 @register_pass(mutates_CFG=False, analysis_only=True)
113 class CUDALegalization(AnalysisPass):
114 
115     _name = "cuda_legalization"
116 
117     def __init__(self):
118         AnalysisPass.__init__(self)
119 
120     def run_pass(self, state):
121         # Early return if NVVM 7
122         from numba.cuda.cudadrv.nvvm import NVVM
123         if NVVM().is_nvvm70:
124             return False
125         # NVVM < 7, need to check for charseq
126         typmap = state.typemap
127 
128         def check_dtype(dtype):
129             if isinstance(dtype, (types.UnicodeCharSeq, types.CharSeq)):
130                 msg = (f"{k} is a char sequence type. This type is not "
131                        "supported with CUDA toolkit versions < 11.2. To "
132                        "use this type, you need to update your CUDA "
133                        "toolkit - try 'conda install cudatoolkit=11' if "
134                        "you are using conda to manage your environment.")
135                 raise TypingError(msg)
136             elif isinstance(dtype, types.Record):
137                 for subdtype in dtype.fields.items():
138                     # subdtype is a (name, _RecordField) pair
139                     check_dtype(subdtype[1].type)
140 
141         for k, v in typmap.items():
142             if isinstance(v, types.Array):
143                 check_dtype(v.dtype)
144         return False
145 
146 
147 class CUDACompiler(CompilerBase):
148     def define_pipelines(self):
149         dpb = DefaultPassBuilder
150         pm = PassManager('cuda')
151 
152         untyped_passes = dpb.define_untyped_pipeline(self.state)
153         pm.passes.extend(untyped_passes.passes)
154 
155         typed_passes = dpb.define_typed_pipeline(self.state)
156         pm.passes.extend(typed_passes.passes)
157         pm.add_pass(CUDALegalization, "CUDA legalization")
158 
159         lowering_passes = self.define_cuda_lowering_pipeline(self.state)
160         pm.passes.extend(lowering_passes.passes)
161 
162         pm.finalize()
163         return [pm]
164 
165     def define_cuda_lowering_pipeline(self, state):
166         pm = PassManager('cuda_lowering')
167         # legalise
168         pm.add_pass(IRLegalization,
169                     "ensure IR is legal prior to lowering")
170         pm.add_pass(AnnotateTypes, "annotate types")
171 
172         # lower
173         pm.add_pass(CreateLibrary, "create library")
174         pm.add_pass(NativeLowering, "native lowering")
175         pm.add_pass(CUDABackend, "cuda backend")
176 
177         pm.finalize()
178         return pm
179 
180 
181 @global_compiler_lock
182 def compile_cuda(pyfunc, return_type, args, debug=False, lineinfo=False,
183                  inline=False, fastmath=False, nvvm_options=None):
184     from .descriptor import cuda_target
185     typingctx = cuda_target.typing_context
186     targetctx = cuda_target.target_context
187 
188     flags = CUDAFlags()
189     # Do not compile (generate native code), just lower (to LLVM)
190     flags.no_compile = True
191     flags.no_cpython_wrapper = True
192     flags.no_cfunc_wrapper = True
193     if debug or lineinfo:
194         # Note both debug and lineinfo turn on debug information in the
195         # compiled code, but we keep them separate arguments in case we
196         # later want to overload some other behavior on the debug flag.
197         # In particular, -opt=3 is not supported with -g.
198         flags.debuginfo = True
199         flags.error_model = 'python'
200     else:
201         flags.error_model = 'numpy'
202     if inline:
203         flags.forceinline = True
204     if fastmath:
205         flags.fastmath = True
206     if nvvm_options:
207         flags.nvvm_options = nvvm_options
208 
209     # Run compilation pipeline
210     from numba.core.target_extension import target_override
211     with target_override('cuda'):
212         cres = compiler.compile_extra(typingctx=typingctx,
213                                       targetctx=targetctx,
214                                       func=pyfunc,
215                                       args=args,
216                                       return_type=return_type,
217                                       flags=flags,
218                                       locals={},
219                                       pipeline_class=CUDACompiler)
220 
221     library = cres.library
222     library.finalize()
223 
224     return cres
225 
226 
227 @global_compiler_lock
228 def compile_ptx(pyfunc, args, debug=False, lineinfo=False, device=False,
229                 fastmath=False, cc=None, opt=True):
230     """Compile a Python function to PTX for a given set of argument types.
231 
232     :param pyfunc: The Python function to compile.
233     :param args: A tuple of argument types to compile for.
234     :param debug: Whether to include debug info in the generated PTX.
235     :type debug: bool
236     :param lineinfo: Whether to include a line mapping from the generated PTX
237                      to the source code. Usually this is used with optimized
238                      code (since debug mode would automatically include this),
239                      so we want debug info in the LLVM but only the line
240                      mapping in the final PTX.
241     :type lineinfo: bool
242     :param device: Whether to compile a device function. Defaults to ``False``,
243                    to compile global kernel functions.
244     :type device: bool
245     :param fastmath: Whether to enable fast math flags (ftz=1, prec_sqrt=0,
246                      prec_div=, and fma=1)
247     :type fastmath: bool
248     :param cc: Compute capability to compile for, as a tuple ``(MAJOR, MINOR)``.
249                Defaults to ``(5, 3)``.
250     :type cc: tuple
251     :param opt: Enable optimizations. Defaults to ``True``.
252     :type opt: bool
253     :return: (ptx, resty): The PTX code and inferred return type
254     :rtype: tuple
255     """
256     if debug and opt:
257         msg = ("debug=True with opt=True (the default) "
258                "is not supported by CUDA. This may result in a crash"
259                " - set debug=False or opt=False.")
260         warn(NumbaInvalidConfigWarning(msg))
261 
262     nvvm_options = {
263         'debug': debug,
264         'lineinfo': lineinfo,
265         'fastmath': fastmath,
266         'opt': 3 if opt else 0
267     }
268 
269     cres = compile_cuda(pyfunc, None, args, debug=debug, lineinfo=lineinfo,
270                         fastmath=fastmath,
271                         nvvm_options=nvvm_options)
272     resty = cres.signature.return_type
273     if device:
274         lib = cres.library
275     else:
276         tgt = cres.target_context
277         code = pyfunc.__code__
278         filename = code.co_filename
279         linenum = code.co_firstlineno
280 
281         lib, kernel = tgt.prepare_cuda_kernel(cres.library, cres.fndesc, debug,
282                                               nvvm_options, filename, linenum)
283 
284     cc = cc or config.CUDA_DEFAULT_PTX_CC
285     ptx = lib.get_asm_str(cc=cc)
286     return ptx, resty
287 
288 
289 def compile_ptx_for_current_device(pyfunc, args, debug=False, lineinfo=False,
290                                    device=False, fastmath=False, opt=True):
291     """Compile a Python function to PTX for a given set of argument types for
292     the current device's compute capabilility. This calls :func:`compile_ptx`
293     with an appropriate ``cc`` value for the current device."""
294     cc = get_current_device().compute_capability
295     return compile_ptx(pyfunc, args, debug=debug, lineinfo=lineinfo,
296                        device=device, fastmath=fastmath, cc=cc, opt=True)
297 
298 
299 def declare_device_function(name, restype, argtypes):
300     return declare_device_function_template(name, restype, argtypes).key
301 
302 
303 def declare_device_function_template(name, restype, argtypes):
304     from .descriptor import cuda_target
305     typingctx = cuda_target.typing_context
306     targetctx = cuda_target.target_context
307     sig = typing.signature(restype, *argtypes)
308     extfn = ExternFunction(name, sig)
309 
310     class device_function_template(ConcreteTemplate):
311         key = extfn
312         cases = [sig]
313 
314     fndesc = funcdesc.ExternalFunctionDescriptor(
315         name=name, restype=restype, argtypes=argtypes)
316     typingctx.insert_user_function(extfn, device_function_template)
317     targetctx.insert_user_function(extfn, fndesc)
318 
319     return device_function_template
320 
321 
322 class ExternFunction(object):
323     def __init__(self, name, sig):
324         self.name = name
325         self.sig = sig
326
```
--- END FILES ---

Please first localize the bug based on the issue statement, and then generate a patch according to the `git diff` format fenced by three backticks.

Here is an example:

```diff
diff --git a/examples/server_async.py b/examples/server_async.py
--- a/examples/server_async.py
+++ b/examples/server_async.py
@@ -313,4 +313,4 @@
 
 
 if __name__ == "__main__":
-    asyncio.run(run_async_server("."), debug=True)
+    asyncio.run(run_async_server(), debug=True)
diff --git a/examples/server_sync.py b/examples/server_sync.py
--- a/examples/server_sync.py
+++ b/examples/server_sync.py
@@ -313,5 +313,5 @@
 
 
 if __name__ == "__main__":
-    server = run_sync_server(".")
+    server = run_sync_server()
     server.shutdown()

```

diff --git a/numba/cuda/compiler.py b/numba/cuda/compiler.py
--- a/numba/cuda/compiler.py
+++ b/numba/cuda/compiler.py
@@ -270,6 +270,10 @@
                         fastmath=fastmath,
                         nvvm_options=nvvm_options)
     resty = cres.signature.return_type
+
+    if resty and not device and resty != types.void:
+        raise TypeError("CUDA kernel must have void return type.")
+
     if device:
         lib = cres.library
     else:

{"golden_diff": "diff --git a/numba/cuda/compiler.py b/numba/cuda/compiler.py\n--- a/numba/cuda/compiler.py\n+++ b/numba/cuda/compiler.py\n@@ -270,6 +270,10 @@\n                         fastmath=fastmath,\n                         nvvm_options=nvvm_options)\n     resty = cres.signature.return_type\n+\n+    if resty and not device and resty != types.void:\n+        raise TypeError(\"CUDA kernel must have void return type.\")\n+\n     if device:\n         lib = cres.library\n     else:\n", "issue": "`compile_ptx()` allows compilation of kernel with non-`void` return type\nFor example:\r\n\r\n```python\r\nfrom numba import cuda, int32\r\n\r\ndef f(x, y):\r\n    return x[0] + y[0]\r\n\r\nptx, resty = cuda.compile_ptx(f, (int32[::1], int32[::1]))\r\n\r\nprint(resty)\r\n# int64\r\n```\r\n\r\ncompiles, and generates PTX where the function body returns nothing (heavily edited for clarity, but the idea is represented:\r\n\r\n```assembly\r\n.visible .entry f (\r\n    // args omitted for brevity\r\n)\r\n{\r\n\tret;\r\n}\r\n```\r\n\r\nUsually we check that the kernel has a void return type in the `@cuda.jit` decorator, which is why this slips by in `compile_ptx`. The check should probably be pushed a bit deeper to cover both uses.\r\n\r\n(cc @brandonwillard )\n", "before_files": [{"content": "from numba.core.typing.templates import ConcreteTemplate\nfrom numba.core import types, typing, funcdesc, config, compiler\nfrom numba.core.compiler import (sanitize_compile_result_entries, CompilerBase,\n                                 DefaultPassBuilder, Flags, Option,\n                                 CompileResult)\nfrom numba.core.compiler_lock import global_compiler_lock\nfrom numba.core.compiler_machinery import (LoweringPass, AnalysisPass,\n                                           PassManager, register_pass)\nfrom numba.core.errors import NumbaInvalidConfigWarning, TypingError\nfrom numba.core.typed_passes import (IRLegalization, NativeLowering,\n                                     AnnotateTypes)\nfrom warnings import warn\nfrom numba.cuda.api import get_current_device\n\n\ndef _nvvm_options_type(x):\n    if x is None:\n        return None\n\n    else:\n        assert isinstance(x, dict)\n        return x\n\n\nclass CUDAFlags(Flags):\n    nvvm_options = Option(\n        type=_nvvm_options_type,\n        default=None,\n        doc=\"NVVM options\",\n    )\n\n\n# The CUDACompileResult (CCR) has a specially-defined entry point equal to its\n# id.  This is because the entry point is used as a key into a dict of\n# overloads by the base dispatcher. The id of the CCR is the only small and\n# unique property of a CompileResult in the CUDA target (cf. the CPU target,\n# which uses its entry_point, which is a pointer value).\n#\n# This does feel a little hackish, and there are two ways in which this could\n# be improved:\n#\n# 1. We could change the core of Numba so that each CompileResult has its own\n#    unique ID that can be used as a key - e.g. a count, similar to the way in\n#    which types have unique counts.\n# 2. At some future time when kernel launch uses a compiled function, the entry\n#    point will no longer need to be a synthetic value, but will instead be a\n#    pointer to the compiled function as in the CPU target.\n\nclass CUDACompileResult(CompileResult):\n    @property\n    def entry_point(self):\n        return id(self)\n\n\ndef cuda_compile_result(**entries):\n    entries = sanitize_compile_result_entries(entries)\n    return CUDACompileResult(**entries)\n\n\n@register_pass(mutates_CFG=True, analysis_only=False)\nclass CUDABackend(LoweringPass):\n\n    _name = \"cuda_backend\"\n\n    def __init__(self):\n        LoweringPass.__init__(self)\n\n    def run_pass(self, state):\n        \"\"\"\n        Back-end: Packages lowering output in a compile result\n        \"\"\"\n        lowered = state['cr']\n        signature = typing.signature(state.return_type, *state.args)\n\n        state.cr = cuda_compile_result(\n            typing_context=state.typingctx,\n            target_context=state.targetctx,\n            typing_error=state.status.fail_reason,\n            type_annotation=state.type_annotation,\n            library=state.library,\n            call_helper=lowered.call_helper,\n            signature=signature,\n            fndesc=lowered.fndesc,\n        )\n        return True\n\n\n@register_pass(mutates_CFG=False, analysis_only=False)\nclass CreateLibrary(LoweringPass):\n    \"\"\"\n    Create a CUDACodeLibrary for the NativeLowering pass to populate. The\n    NativeLowering pass will create a code library if none exists, but we need\n    to set it up with nvvm_options from the flags if they are present.\n    \"\"\"\n\n    _name = \"create_library\"\n\n    def __init__(self):\n        LoweringPass.__init__(self)\n\n    def run_pass(self, state):\n        codegen = state.targetctx.codegen()\n        name = state.func_id.func_qualname\n        nvvm_options = state.flags.nvvm_options\n        state.library = codegen.create_library(name, nvvm_options=nvvm_options)\n        # Enable object caching upfront so that the library can be serialized.\n        state.library.enable_object_caching()\n\n        return True\n\n\n@register_pass(mutates_CFG=False, analysis_only=True)\nclass CUDALegalization(AnalysisPass):\n\n    _name = \"cuda_legalization\"\n\n    def __init__(self):\n        AnalysisPass.__init__(self)\n\n    def run_pass(self, state):\n        # Early return if NVVM 7\n        from numba.cuda.cudadrv.nvvm import NVVM\n        if NVVM().is_nvvm70:\n            return False\n        # NVVM < 7, need to check for charseq\n        typmap = state.typemap\n\n        def check_dtype(dtype):\n            if isinstance(dtype, (types.UnicodeCharSeq, types.CharSeq)):\n                msg = (f\"{k} is a char sequence type. This type is not \"\n                       \"supported with CUDA toolkit versions < 11.2. To \"\n                       \"use this type, you need to update your CUDA \"\n                       \"toolkit - try 'conda install cudatoolkit=11' if \"\n                       \"you are using conda to manage your environment.\")\n                raise TypingError(msg)\n            elif isinstance(dtype, types.Record):\n                for subdtype in dtype.fields.items():\n                    # subdtype is a (name, _RecordField) pair\n                    check_dtype(subdtype[1].type)\n\n        for k, v in typmap.items():\n            if isinstance(v, types.Array):\n                check_dtype(v.dtype)\n        return False\n\n\nclass CUDACompiler(CompilerBase):\n    def define_pipelines(self):\n        dpb = DefaultPassBuilder\n        pm = PassManager('cuda')\n\n        untyped_passes = dpb.define_untyped_pipeline(self.state)\n        pm.passes.extend(untyped_passes.passes)\n\n        typed_passes = dpb.define_typed_pipeline(self.state)\n        pm.passes.extend(typed_passes.passes)\n        pm.add_pass(CUDALegalization, \"CUDA legalization\")\n\n        lowering_passes = self.define_cuda_lowering_pipeline(self.state)\n        pm.passes.extend(lowering_passes.passes)\n\n        pm.finalize()\n        return [pm]\n\n    def define_cuda_lowering_pipeline(self, state):\n        pm = PassManager('cuda_lowering')\n        # legalise\n        pm.add_pass(IRLegalization,\n                    \"ensure IR is legal prior to lowering\")\n        pm.add_pass(AnnotateTypes, \"annotate types\")\n\n        # lower\n        pm.add_pass(CreateLibrary, \"create library\")\n        pm.add_pass(NativeLowering, \"native lowering\")\n        pm.add_pass(CUDABackend, \"cuda backend\")\n\n        pm.finalize()\n        return pm\n\n\n@global_compiler_lock\ndef compile_cuda(pyfunc, return_type, args, debug=False, lineinfo=False,\n                 inline=False, fastmath=False, nvvm_options=None):\n    from .descriptor import cuda_target\n    typingctx = cuda_target.typing_context\n    targetctx = cuda_target.target_context\n\n    flags = CUDAFlags()\n    # Do not compile (generate native code), just lower (to LLVM)\n    flags.no_compile = True\n    flags.no_cpython_wrapper = True\n    flags.no_cfunc_wrapper = True\n    if debug or lineinfo:\n        # Note both debug and lineinfo turn on debug information in the\n        # compiled code, but we keep them separate arguments in case we\n        # later want to overload some other behavior on the debug flag.\n        # In particular, -opt=3 is not supported with -g.\n        flags.debuginfo = True\n        flags.error_model = 'python'\n    else:\n        flags.error_model = 'numpy'\n    if inline:\n        flags.forceinline = True\n    if fastmath:\n        flags.fastmath = True\n    if nvvm_options:\n        flags.nvvm_options = nvvm_options\n\n    # Run compilation pipeline\n    from numba.core.target_extension import target_override\n    with target_override('cuda'):\n        cres = compiler.compile_extra(typingctx=typingctx,\n                                      targetctx=targetctx,\n                                      func=pyfunc,\n                                      args=args,\n                                      return_type=return_type,\n                                      flags=flags,\n                                      locals={},\n                                      pipeline_class=CUDACompiler)\n\n    library = cres.library\n    library.finalize()\n\n    return cres\n\n\n@global_compiler_lock\ndef compile_ptx(pyfunc, args, debug=False, lineinfo=False, device=False,\n                fastmath=False, cc=None, opt=True):\n    \"\"\"Compile a Python function to PTX for a given set of argument types.\n\n    :param pyfunc: The Python function to compile.\n    :param args: A tuple of argument types to compile for.\n    :param debug: Whether to include debug info in the generated PTX.\n    :type debug: bool\n    :param lineinfo: Whether to include a line mapping from the generated PTX\n                     to the source code. Usually this is used with optimized\n                     code (since debug mode would automatically include this),\n                     so we want debug info in the LLVM but only the line\n                     mapping in the final PTX.\n    :type lineinfo: bool\n    :param device: Whether to compile a device function. Defaults to ``False``,\n                   to compile global kernel functions.\n    :type device: bool\n    :param fastmath: Whether to enable fast math flags (ftz=1, prec_sqrt=0,\n                     prec_div=, and fma=1)\n    :type fastmath: bool\n    :param cc: Compute capability to compile for, as a tuple ``(MAJOR, MINOR)``.\n               Defaults to ``(5, 3)``.\n    :type cc: tuple\n    :param opt: Enable optimizations. Defaults to ``True``.\n    :type opt: bool\n    :return: (ptx, resty): The PTX code and inferred return type\n    :rtype: tuple\n    \"\"\"\n    if debug and opt:\n        msg = (\"debug=True with opt=True (the default) \"\n               \"is not supported by CUDA. This may result in a crash\"\n               \" - set debug=False or opt=False.\")\n        warn(NumbaInvalidConfigWarning(msg))\n\n    nvvm_options = {\n        'debug': debug,\n        'lineinfo': lineinfo,\n        'fastmath': fastmath,\n        'opt': 3 if opt else 0\n    }\n\n    cres = compile_cuda(pyfunc, None, args, debug=debug, lineinfo=lineinfo,\n                        fastmath=fastmath,\n                        nvvm_options=nvvm_options)\n    resty = cres.signature.return_type\n    if device:\n        lib = cres.library\n    else:\n        tgt = cres.target_context\n        code = pyfunc.__code__\n        filename = code.co_filename\n        linenum = code.co_firstlineno\n\n        lib, kernel = tgt.prepare_cuda_kernel(cres.library, cres.fndesc, debug,\n                                              nvvm_options, filename, linenum)\n\n    cc = cc or config.CUDA_DEFAULT_PTX_CC\n    ptx = lib.get_asm_str(cc=cc)\n    return ptx, resty\n\n\ndef compile_ptx_for_current_device(pyfunc, args, debug=False, lineinfo=False,\n                                   device=False, fastmath=False, opt=True):\n    \"\"\"Compile a Python function to PTX for a given set of argument types for\n    the current device's compute capabilility. This calls :func:`compile_ptx`\n    with an appropriate ``cc`` value for the current device.\"\"\"\n    cc = get_current_device().compute_capability\n    return compile_ptx(pyfunc, args, debug=debug, lineinfo=lineinfo,\n                       device=device, fastmath=fastmath, cc=cc, opt=True)\n\n\ndef declare_device_function(name, restype, argtypes):\n    return declare_device_function_template(name, restype, argtypes).key\n\n\ndef declare_device_function_template(name, restype, argtypes):\n    from .descriptor import cuda_target\n    typingctx = cuda_target.typing_context\n    targetctx = cuda_target.target_context\n    sig = typing.signature(restype, *argtypes)\n    extfn = ExternFunction(name, sig)\n\n    class device_function_template(ConcreteTemplate):\n        key = extfn\n        cases = [sig]\n\n    fndesc = funcdesc.ExternalFunctionDescriptor(\n        name=name, restype=restype, argtypes=argtypes)\n    typingctx.insert_user_function(extfn, device_function_template)\n    targetctx.insert_user_function(extfn, fndesc)\n\n    return device_function_template\n\n\nclass ExternFunction(object):\n    def __init__(self, name, sig):\n        self.name = name\n        self.sig = sig\n", "path": "numba/cuda/compiler.py"}], "after_files": [{"content": "from numba.core.typing.templates import ConcreteTemplate\nfrom numba.core import types, typing, funcdesc, config, compiler\nfrom numba.core.compiler import (sanitize_compile_result_entries, CompilerBase,\n                                 DefaultPassBuilder, Flags, Option,\n                                 CompileResult)\nfrom numba.core.compiler_lock import global_compiler_lock\nfrom numba.core.compiler_machinery import (LoweringPass, AnalysisPass,\n                                           PassManager, register_pass)\nfrom numba.core.errors import NumbaInvalidConfigWarning, TypingError\nfrom numba.core.typed_passes import (IRLegalization, NativeLowering,\n                                     AnnotateTypes)\nfrom warnings import warn\nfrom numba.cuda.api import get_current_device\n\n\ndef _nvvm_options_type(x):\n    if x is None:\n        return None\n\n    else:\n        assert isinstance(x, dict)\n        return x\n\n\nclass CUDAFlags(Flags):\n    nvvm_options = Option(\n        type=_nvvm_options_type,\n        default=None,\n        doc=\"NVVM options\",\n    )\n\n\n# The CUDACompileResult (CCR) has a specially-defined entry point equal to its\n# id.  This is because the entry point is used as a key into a dict of\n# overloads by the base dispatcher. The id of the CCR is the only small and\n# unique property of a CompileResult in the CUDA target (cf. the CPU target,\n# which uses its entry_point, which is a pointer value).\n#\n# This does feel a little hackish, and there are two ways in which this could\n# be improved:\n#\n# 1. We could change the core of Numba so that each CompileResult has its own\n#    unique ID that can be used as a key - e.g. a count, similar to the way in\n#    which types have unique counts.\n# 2. At some future time when kernel launch uses a compiled function, the entry\n#    point will no longer need to be a synthetic value, but will instead be a\n#    pointer to the compiled function as in the CPU target.\n\nclass CUDACompileResult(CompileResult):\n    @property\n    def entry_point(self):\n        return id(self)\n\n\ndef cuda_compile_result(**entries):\n    entries = sanitize_compile_result_entries(entries)\n    return CUDACompileResult(**entries)\n\n\n@register_pass(mutates_CFG=True, analysis_only=False)\nclass CUDABackend(LoweringPass):\n\n    _name = \"cuda_backend\"\n\n    def __init__(self):\n        LoweringPass.__init__(self)\n\n    def run_pass(self, state):\n        \"\"\"\n        Back-end: Packages lowering output in a compile result\n        \"\"\"\n        lowered = state['cr']\n        signature = typing.signature(state.return_type, *state.args)\n\n        state.cr = cuda_compile_result(\n            typing_context=state.typingctx,\n            target_context=state.targetctx,\n            typing_error=state.status.fail_reason,\n            type_annotation=state.type_annotation,\n            library=state.library,\n            call_helper=lowered.call_helper,\n            signature=signature,\n            fndesc=lowered.fndesc,\n        )\n        return True\n\n\n@register_pass(mutates_CFG=False, analysis_only=False)\nclass CreateLibrary(LoweringPass):\n    \"\"\"\n    Create a CUDACodeLibrary for the NativeLowering pass to populate. The\n    NativeLowering pass will create a code library if none exists, but we need\n    to set it up with nvvm_options from the flags if they are present.\n    \"\"\"\n\n    _name = \"create_library\"\n\n    def __init__(self):\n        LoweringPass.__init__(self)\n\n    def run_pass(self, state):\n        codegen = state.targetctx.codegen()\n        name = state.func_id.func_qualname\n        nvvm_options = state.flags.nvvm_options\n        state.library = codegen.create_library(name, nvvm_options=nvvm_options)\n        # Enable object caching upfront so that the library can be serialized.\n        state.library.enable_object_caching()\n\n        return True\n\n\n@register_pass(mutates_CFG=False, analysis_only=True)\nclass CUDALegalization(AnalysisPass):\n\n    _name = \"cuda_legalization\"\n\n    def __init__(self):\n        AnalysisPass.__init__(self)\n\n    def run_pass(self, state):\n        # Early return if NVVM 7\n        from numba.cuda.cudadrv.nvvm import NVVM\n        if NVVM().is_nvvm70:\n            return False\n        # NVVM < 7, need to check for charseq\n        typmap = state.typemap\n\n        def check_dtype(dtype):\n            if isinstance(dtype, (types.UnicodeCharSeq, types.CharSeq)):\n                msg = (f\"{k} is a char sequence type. This type is not \"\n                       \"supported with CUDA toolkit versions < 11.2. To \"\n                       \"use this type, you need to update your CUDA \"\n                       \"toolkit - try 'conda install cudatoolkit=11' if \"\n                       \"you are using conda to manage your environment.\")\n                raise TypingError(msg)\n            elif isinstance(dtype, types.Record):\n                for subdtype in dtype.fields.items():\n                    # subdtype is a (name, _RecordField) pair\n                    check_dtype(subdtype[1].type)\n\n        for k, v in typmap.items():\n            if isinstance(v, types.Array):\n                check_dtype(v.dtype)\n        return False\n\n\nclass CUDACompiler(CompilerBase):\n    def define_pipelines(self):\n        dpb = DefaultPassBuilder\n        pm = PassManager('cuda')\n\n        untyped_passes = dpb.define_untyped_pipeline(self.state)\n        pm.passes.extend(untyped_passes.passes)\n\n        typed_passes = dpb.define_typed_pipeline(self.state)\n        pm.passes.extend(typed_passes.passes)\n        pm.add_pass(CUDALegalization, \"CUDA legalization\")\n\n        lowering_passes = self.define_cuda_lowering_pipeline(self.state)\n        pm.passes.extend(lowering_passes.passes)\n\n        pm.finalize()\n        return [pm]\n\n    def define_cuda_lowering_pipeline(self, state):\n        pm = PassManager('cuda_lowering')\n        # legalise\n        pm.add_pass(IRLegalization,\n                    \"ensure IR is legal prior to lowering\")\n        pm.add_pass(AnnotateTypes, \"annotate types\")\n\n        # lower\n        pm.add_pass(CreateLibrary, \"create library\")\n        pm.add_pass(NativeLowering, \"native lowering\")\n        pm.add_pass(CUDABackend, \"cuda backend\")\n\n        pm.finalize()\n        return pm\n\n\n@global_compiler_lock\ndef compile_cuda(pyfunc, return_type, args, debug=False, lineinfo=False,\n                 inline=False, fastmath=False, nvvm_options=None):\n    from .descriptor import cuda_target\n    typingctx = cuda_target.typing_context\n    targetctx = cuda_target.target_context\n\n    flags = CUDAFlags()\n    # Do not compile (generate native code), just lower (to LLVM)\n    flags.no_compile = True\n    flags.no_cpython_wrapper = True\n    flags.no_cfunc_wrapper = True\n    if debug or lineinfo:\n        # Note both debug and lineinfo turn on debug information in the\n        # compiled code, but we keep them separate arguments in case we\n        # later want to overload some other behavior on the debug flag.\n        # In particular, -opt=3 is not supported with -g.\n        flags.debuginfo = True\n        flags.error_model = 'python'\n    else:\n        flags.error_model = 'numpy'\n    if inline:\n        flags.forceinline = True\n    if fastmath:\n        flags.fastmath = True\n    if nvvm_options:\n        flags.nvvm_options = nvvm_options\n\n    # Run compilation pipeline\n    from numba.core.target_extension import target_override\n    with target_override('cuda'):\n        cres = compiler.compile_extra(typingctx=typingctx,\n                                      targetctx=targetctx,\n                                      func=pyfunc,\n                                      args=args,\n                                      return_type=return_type,\n                                      flags=flags,\n                                      locals={},\n                                      pipeline_class=CUDACompiler)\n\n    library = cres.library\n    library.finalize()\n\n    return cres\n\n\n@global_compiler_lock\ndef compile_ptx(pyfunc, args, debug=False, lineinfo=False, device=False,\n                fastmath=False, cc=None, opt=True):\n    \"\"\"Compile a Python function to PTX for a given set of argument types.\n\n    :param pyfunc: The Python function to compile.\n    :param args: A tuple of argument types to compile for.\n    :param debug: Whether to include debug info in the generated PTX.\n    :type debug: bool\n    :param lineinfo: Whether to include a line mapping from the generated PTX\n                     to the source code. Usually this is used with optimized\n                     code (since debug mode would automatically include this),\n                     so we want debug info in the LLVM but only the line\n                     mapping in the final PTX.\n    :type lineinfo: bool\n    :param device: Whether to compile a device function. Defaults to ``False``,\n                   to compile global kernel functions.\n    :type device: bool\n    :param fastmath: Whether to enable fast math flags (ftz=1, prec_sqrt=0,\n                     prec_div=, and fma=1)\n    :type fastmath: bool\n    :param cc: Compute capability to compile for, as a tuple ``(MAJOR, MINOR)``.\n               Defaults to ``(5, 3)``.\n    :type cc: tuple\n    :param opt: Enable optimizations. Defaults to ``True``.\n    :type opt: bool\n    :return: (ptx, resty): The PTX code and inferred return type\n    :rtype: tuple\n    \"\"\"\n    if debug and opt:\n        msg = (\"debug=True with opt=True (the default) \"\n               \"is not supported by CUDA. This may result in a crash\"\n               \" - set debug=False or opt=False.\")\n        warn(NumbaInvalidConfigWarning(msg))\n\n    nvvm_options = {\n        'debug': debug,\n        'lineinfo': lineinfo,\n        'fastmath': fastmath,\n        'opt': 3 if opt else 0\n    }\n\n    cres = compile_cuda(pyfunc, None, args, debug=debug, lineinfo=lineinfo,\n                        fastmath=fastmath,\n                        nvvm_options=nvvm_options)\n    resty = cres.signature.return_type\n\n    if resty and not device and resty != types.void:\n        raise TypeError(\"CUDA kernel must have void return type.\")\n\n    if device:\n        lib = cres.library\n    else:\n        tgt = cres.target_context\n        code = pyfunc.__code__\n        filename = code.co_filename\n        linenum = code.co_firstlineno\n\n        lib, kernel = tgt.prepare_cuda_kernel(cres.library, cres.fndesc, debug,\n                                              nvvm_options, filename, linenum)\n\n    cc = cc or config.CUDA_DEFAULT_PTX_CC\n    ptx = lib.get_asm_str(cc=cc)\n    return ptx, resty\n\n\ndef compile_ptx_for_current_device(pyfunc, args, debug=False, lineinfo=False,\n                                   device=False, fastmath=False, opt=True):\n    \"\"\"Compile a Python function to PTX for a given set of argument types for\n    the current device's compute capabilility. This calls :func:`compile_ptx`\n    with an appropriate ``cc`` value for the current device.\"\"\"\n    cc = get_current_device().compute_capability\n    return compile_ptx(pyfunc, args, debug=debug, lineinfo=lineinfo,\n                       device=device, fastmath=fastmath, cc=cc, opt=True)\n\n\ndef declare_device_function(name, restype, argtypes):\n    return declare_device_function_template(name, restype, argtypes).key\n\n\ndef declare_device_function_template(name, restype, argtypes):\n    from .descriptor import cuda_target\n    typingctx = cuda_target.typing_context\n    targetctx = cuda_target.target_context\n    sig = typing.signature(restype, *argtypes)\n    extfn = ExternFunction(name, sig)\n\n    class device_function_template(ConcreteTemplate):\n        key = extfn\n        cases = [sig]\n\n    fndesc = funcdesc.ExternalFunctionDescriptor(\n        name=name, restype=restype, argtypes=argtypes)\n    typingctx.insert_user_function(extfn, device_function_template)\n    targetctx.insert_user_function(extfn, fndesc)\n\n    return device_function_template\n\n\nclass ExternFunction(object):\n    def __init__(self, name, sig):\n        self.name = name\n        self.sig = sig\n", "path": "numba/cuda/compiler.py"}]}