Pass correct linkage_name attribute to DIBuilder

numba_dppy/dppy_debuginfo.py

@@ -0,0 +1,38 @@
+# Copyright 2021 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#      http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+from numba.core.debuginfo import DIBuilder
+from llvmlite import ir
+
+
+class DPPYDIBuilder(DIBuilder):
+    def __init__(self, module, filepath, linkage_name):
+        DIBuilder.__init__(self, module, filepath)
+        self.linkage_name = linkage_name
+
+    def mark_subprogram(self, function, name, loc):
+        di_subp = self._add_subprogram(
+            name=name, linkagename=self.linkage_name, line=loc.line
+        )
+        function.set_metadata("dbg", di_subp)
+        # disable inlining for this function for easier debugging
+        function.attributes.add("noinline")
+
+    def _di_compile_unit(self):
+        di = super()._di_compile_unit()
+        operands = dict(di.operands)
+        operands["language"] = ir.DIToken("DW_LANG_C_plus_plus")
+        operands["producer"] = "numba-dppy"
+        di.operands = tuple(operands.items())
+        return di

numba_dppy/dppy_lowerer.py

@@ -1223,8 +1223,30 @@ def __init__(self, context, library, fndesc, func_ir, metadata=None):
         cpu_context = (
             context.cpu_context if isinstance(context, DPPYTargetContext) else context
         )
-        self.gpu_lower = Lower(context, library, fndesc, func_ir, metadata)
-        self.cpu_lower = Lower(cpu_context, library, fndesc_cpu, func_ir_cpu, metadata)
+        self.gpu_lower = self._lower(context, library, fndesc, func_ir, metadata)
+        self.cpu_lower = self._lower(
+            cpu_context, library, fndesc_cpu, func_ir_cpu, metadata
+        )
+
+    def _lower(self, context, library, fndesc, func_ir, metadata):
+        """Create Lower with changed linkageName in debug info"""
+        lower = Lower(context, library, fndesc, func_ir, metadata)
+
+        # Debuginfo
+        if context.enable_debuginfo:
+            from numba.core.funcdesc import qualifying_prefix, default_mangler
+            from numba_dppy.dppy_debuginfo import DPPYDIBuilder
+
+            qualprefix = qualifying_prefix(fndesc.modname, fndesc.qualname)
+            mangled_qualname = default_mangler(qualprefix, fndesc.argtypes)
+
+            lower.debuginfo = DPPYDIBuilder(
+                module=lower.module,
+                filepath=func_ir.loc.filename,
+                linkage_name=mangled_qualname,
+            )
+
+        return lower
 
     def lower(self):
         """Numba-dppy's custom lowering function.

numba_dppy/tests/test_debuginfo.py

@@ -193,3 +193,48 @@ def data_parallel_sum(a, b, c):
 
     for tag in ir_tags:
         assert debug_option == make_check(kernel_ir, tag)
+
+
+def test_debuginfo_DISubprogram_linkageName():
+    @dppy.kernel
+    def func(a, b):
+        i = dppy.get_global_id(0)
+        b[i] = a[i]
+
+    ir_tags = [
+        r'\!DISubprogram\(.*linkageName: ".*e4func.*"',  # e4func is func(), e8func$241 is func$1()
+    ]
+
+    sycl_queue = dpctl.get_current_queue()
+    sig = (
+        types.float32[:],
+        types.float32[:],
+    )
+
+    kernel_ir = get_kernel_ir(sycl_queue, func, sig, debug=True)
+
+    for tag in ir_tags:
+        assert make_check(kernel_ir, tag)
+
+
+def test_debuginfo_DICompileUnit_language_and_producer():
+    @dppy.kernel
+    def func(a, b):
+        i = dppy.get_global_id(0)
+        b[i] = a[i]
+
+    ir_tags = [
+        r"\!DICompileUnit\(language: DW_LANG_C_plus_plus,",
+        r'\!DICompileUnit\(.*producer: "numba-dppy"',
+    ]
+
+    sycl_queue = dpctl.get_current_queue()
+    sig = (
+        types.float32[:],
+        types.float32[:],
+    )
+
+    kernel_ir = get_kernel_ir(sycl_queue, func, sig, debug=True)
+
+    for tag in ir_tags:
+        assert make_check(kernel_ir, tag)