From 2daad1ba82e8e13c7d5e317dae272f2d431c9dac Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Tue, 7 Jan 2025 15:58:40 +0000 Subject: [PATCH 1/4] Add a CUDA DI Builder This will be the basis for customising CUDA DI generation. This initial implementation corrects the type for boolean values. --- numba_cuda/numba/cuda/debuginfo.py | 34 ++++++++++++++++++++++++++++++ numba_cuda/numba/cuda/target.py | 6 +++--- 2 files changed, 37 insertions(+), 3 deletions(-) create mode 100644 numba_cuda/numba/cuda/debuginfo.py diff --git a/numba_cuda/numba/cuda/debuginfo.py b/numba_cuda/numba/cuda/debuginfo.py new file mode 100644 index 000000000..886311f04 --- /dev/null +++ b/numba_cuda/numba/cuda/debuginfo.py @@ -0,0 +1,34 @@ +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: + is_bool = True + elif isinstance(datamodel.fe_type, types.Boolean): + is_bool = True + + # Booleans should use our implementation until upstream Numba is fixed + if is_bool: + m = self.module + name = str(lltype) + 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) diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index d6c9225f1..9739ad5d9 100644 --- a/numba_cuda/numba/cuda/target.py +++ b/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): From 735fcd24fee0b44183cda1779e1a5091e52cbd94 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Tue, 7 Jan 2025 16:33:26 +0000 Subject: [PATCH 2/4] Boolean debug info: correct name when datamodel is present --- numba_cuda/numba/cuda/debuginfo.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/numba_cuda/numba/cuda/debuginfo.py b/numba_cuda/numba/cuda/debuginfo.py index 886311f04..ff29f9949 100644 --- a/numba_cuda/numba/cuda/debuginfo.py +++ b/numba_cuda/numba/cuda/debuginfo.py @@ -13,14 +13,15 @@ def _var_type(self, lltype, size, datamodel=None): 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 - name = str(lltype) bitsize = _BYTE_SIZE * size ditok = "DW_ATE_boolean" From 447ba68d0f78dfdcbbfb247c57579c45454bcf70 Mon Sep 17 00:00:00 2001 From: jiel-nv Date: Thu, 20 Feb 2025 00:54:01 -0800 Subject: [PATCH 3/4] Add a test for the boolean type fix --- .../numba/cuda/tests/cudapy/test_debuginfo.py | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 2327deb90..9e5cc9ff6 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -90,6 +90,25 @@ 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 + + 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 = f'!{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],) From 451c36bc10c2ba0bce4d861d01e74c5784213d58 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Thu, 20 Feb 2025 11:02:42 +0000 Subject: [PATCH 4/4] Fix flake8 --- numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 9e5cc9ff6..df6af18b5 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -92,9 +92,10 @@ def f(cond): 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 + z = x == y # noqa: F841 llvm_ir = f.inspect_llvm(sig) @@ -105,7 +106,7 @@ def f(x, y): mdnode_id = match.group(1) # verify the DIBasicType has correct encoding attribute DW_ATE_boolean - pat = f'!{mdnode_id}\s+=\s+!DIBasicType\(.*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)