Emit correct kernel arguments debug metadata types when NUMBA_DUMP_LLVM is on (#136)

After shifting out the return value type from the operands of 'types' field of
the 'DISubroutineType' metadata node, compiler needs to clear the cached string
representation of the corresponding MDValue object if NUMBA_DUMP_LLVM is on.

Because at that moment, the cached string has already been filled with the
tuple before kernel_fixup() which has the return value type and continue being
used when printing out the LLVM IR.

This change fixes #135

* Add a test for issue#135

Enriches an existing test by adding a variation, i.e. turn NUMBA_DUMP_LLVM on.

* Use the output capturing to prevent dumping to stdout.

---------

Co-authored-by: Graham Markall <[email protected]>

Author: jiel-nv

numba_cuda/numba/cuda/compiler.py

@@ -427,6 +427,8 @@ def kernel_fixup(kernel, debug):
                         if tm_name == 'types':
                             types = tm_value
                             types.operands = types.operands[1:]
+                            if config.DUMP_LLVM:
+                                types._clear_string_cache()
 
     # Mark as a kernel for NVVM
 

numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py

@@ -1,4 +1,4 @@
-from numba.tests.support import override_config
+from numba.tests.support import (override_config, captured_stdout)
 from numba.cuda.testing import skip_on_cudasim
 from numba import cuda
 from numba.core import types
@@ -268,7 +268,7 @@ def kernel(x, y):
         three_device_fns(kernel_debug=False, leaf_debug=True)
         three_device_fns(kernel_debug=False, leaf_debug=False)
 
-    def test_kernel_args_types(self):
+    def _test_kernel_args_types(self):
         sig = (types.int32, types.int32)
 
         @cuda.jit("void(int32, int32)", debug=True, opt=False)
@@ -298,6 +298,15 @@ def f(x, y):
         match = re.compile(pat).search(llvm_ir)
         self.assertIsNotNone(match, msg=llvm_ir)
 
+    def test_kernel_args_types(self):
+        self._test_kernel_args_types()
+
+    def test_kernel_args_types_dump(self):
+        # see issue#135
+        with override_config('DUMP_LLVM', 1):
+            with captured_stdout():
+                self._test_kernel_args_types()
+
 
 if __name__ == '__main__':
     unittest.main()