Add a CUDA DI Builder (#104)

This will be the basis for customising CUDA DI generation.

This initial implementation corrects the type for boolean values.

* Boolean debug info: correct name when datamodel is present

* Add a test for the boolean type fix

---------

Co-authored-by: jiel-nv <[email protected]>

Author: gmarkall

numba_cuda/numba/cuda/debuginfo.py

@@ -0,0 +1,35 @@
+from llvmlite import ir
+from numba.core import types
+from numba.core.debuginfo import DIBuilder
+
+_BYTE_SIZE = 8
+
+
+class CUDADIBuilder(DIBuilder):
+
+    def _var_type(self, lltype, size, datamodel=None):
+        is_bool = False
+
+        if isinstance(lltype, ir.IntType):
+            if datamodel is None:
+                if size == 1:
+                    name = str(lltype)
+                    is_bool = True
+            elif isinstance(datamodel.fe_type, types.Boolean):
+                name = str(datamodel.fe_type)
+                is_bool = True
+
+        # Booleans should use our implementation until upstream Numba is fixed
+        if is_bool:
+            m = self.module
+            bitsize = _BYTE_SIZE * size
+            ditok = "DW_ATE_boolean"
+
+            return m.add_debug_info('DIBasicType', {
+                'name': name,
+                'size': bitsize,
+                'encoding': ir.DIToken(ditok),
+            })
+
+        # For other cases, use upstream Numba implementation
+        return super()._var_type(lltype, size, datamodel=datamodel)

numba_cuda/numba/cuda/target.py

@@ -3,8 +3,7 @@
 import llvmlite.binding as ll
 from llvmlite import ir
 
-from numba.core import (cgutils, config, debuginfo, itanium_mangler, types,
-                        typing, utils)
+from numba.core import cgutils, config, itanium_mangler, types, typing, utils
 from numba.core.dispatcher import Dispatcher
 from numba.core.base import BaseContext
 from numba.core.callconv import BaseCallConv, MinimalCallConv
@@ -13,6 +12,7 @@
 
 from .cudadrv import nvvm
 from numba.cuda import codegen, nvvmutils, ufuncs
+from numba.cuda.debuginfo import CUDADIBuilder
 from numba.cuda.models import cuda_data_manager
 
 # -----------------------------------------------------------------------------
@@ -80,7 +80,7 @@ def enable_nrt(self):
 
     @property
     def DIBuilder(self):
-        return debuginfo.DIBuilder
+        return CUDADIBuilder
 
     @property
     def enable_boundscheck(self):

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

@@ -90,6 +90,26 @@ def f(cond):
         match = re.compile(pat).search(llvm_ir)
         self.assertIsNone(match, msg=llvm_ir)
 
+    def test_bool_type(self):
+        sig = (types.int32, types.int32)
+
+        @cuda.jit("void(int32, int32)", debug=True, opt=False)
+        def f(x, y):
+            z = x == y  # noqa: F841
+
+        llvm_ir = f.inspect_llvm(sig)
+
+        # extract the metadata node id from `type` field of DILocalVariable
+        pat = r'!DILocalVariable\(.*name:\s+"z".*type:\s+!(\d+)'
+        match = re.compile(pat).search(llvm_ir)
+        self.assertIsNotNone(match, msg=llvm_ir)
+        mdnode_id = match.group(1)
+
+        # verify the DIBasicType has correct encoding attribute DW_ATE_boolean
+        pat = rf'!{mdnode_id}\s+=\s+!DIBasicType\(.*DW_ATE_boolean'
+        match = re.compile(pat).search(llvm_ir)
+        self.assertIsNotNone(match, msg=llvm_ir)
+
     @unittest.skip("Wrappers no longer exist")
     def test_wrapper_has_debuginfo(self):
         sig = (types.int32[::1],)