Fix the debug info of GridGroup type

The customized CUDA front end type 'GridGroup' maps to the back end
type 'ir.IntType(64)' which is defined in numba.cuda.models.

The numba DIBuilder emit variable types for simple numeric types by
looking into back end type which falls through the 'DW_ATE_float'
because the 'datamodel.fe_type' is not 'types.Integer'.

So emitting this type through the CUDADIBuilder is needed. Since the
type represents the handle of a group object, giving it the encoding
of 'DW_ATE_unsigned' should be reasonable.

A test is also added in this change.

This change addresses #107

Author: jiel-nv

numba_cuda/numba/cuda/debuginfo.py

@@ -1,6 +1,7 @@
 from llvmlite import ir
 from numba.core import types
 from numba.core.debuginfo import DIBuilder
+from numba.cuda.types import GridGroup
 
 _BYTE_SIZE = 8
 
@@ -9,21 +10,29 @@ class CUDADIBuilder(DIBuilder):
 
     def _var_type(self, lltype, size, datamodel=None):
         is_bool = False
+        is_grid_group = 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):
+            else:
                 name = str(datamodel.fe_type)
-                is_bool = True
+                if isinstance(datamodel.fe_type, types.Boolean):
+                    is_bool = True
+                elif isinstance(datamodel.fe_type, GridGroup):
+                    is_grid_group = True
 
-        # Booleans should use our implementation until upstream Numba is fixed
-        if is_bool:
+        if is_bool or is_grid_group:
             m = self.module
             bitsize = _BYTE_SIZE * size
-            ditok = "DW_ATE_boolean"
+            # Boolean type workaround until upstream Numba is fixed
+            if is_bool:
+                ditok = "DW_ATE_boolean"
+            # GridGroup type should use numba.cuda implementation
+            elif is_grid_group:
+                ditok = "DW_ATE_unsigned"
 
             return m.add_debug_info('DIBasicType', {
                 'name': name,

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

@@ -110,6 +110,19 @@ def f(x, y):
         match = re.compile(pat).search(llvm_ir)
         self.assertIsNotNone(match, msg=llvm_ir)
 
+    def test_grid_group_type(self):
+        sig = (types.int32,)
+
+        @cuda.jit(sig, debug=True, opt=False)
+        def f(x):
+            grid = cuda.cg.this_grid()  # noqa: F841
+
+        llvm_ir = f.inspect_llvm(sig)
+
+        pat = r'!DIBasicType\(.*DW_ATE_unsigned, name: "GridGroup", size: 64'
+        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],)