From eb8a67910cde8ea0b5e78c74682e7db053263bc1 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Tue, 29 Jul 2025 13:28:25 -0700 Subject: [PATCH 01/10] Add more debuginfo tests from upstream --- numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 3f734ad71..28e6da2f7 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -3,9 +3,15 @@ from numba import cuda from numba.core import types from numba.cuda.testing import CUDATestCase +from textwrap import dedent +import math import itertools import re import unittest +import warnings +from numba.core.errors import NumbaDebugInfoWarning +from numba.tests.support import ignore_internal_warnings +import numpy as np @skip_on_cudasim("Simulator does not produce debug dumps") From 68e2a5522f0bfdf0024f4996231bb18dd7290d50 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Wed, 30 Jul 2025 10:10:53 -0700 Subject: [PATCH 02/10] Address review comments Make the filecheck error message a little more readable and break the inlining test into two checks. --- numba_cuda/numba/cuda/testing.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/numba_cuda/numba/cuda/testing.py b/numba_cuda/numba/cuda/testing.py index 91d1cf243..60712c462 100644 --- a/numba_cuda/numba/cuda/testing.py +++ b/numba_cuda/numba/cuda/testing.py @@ -177,6 +177,8 @@ def assertFileCheckMatches( self.fail( f"FileCheck failed:\n{matcher.stderr.getvalue()}\n\n" + + f"Check prefixes:\n{check_prefixes}\n\n" + + f"Check patterns:\n{check_patterns}\n" + dump_instructions ) From bc208655c98d8c160b5753ec341e4d55a58ace79 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Sat, 2 Aug 2025 08:02:39 -0700 Subject: [PATCH 03/10] Add opt=False to debug tests --- numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py | 12 +++--------- .../numba/cuda/tests/cudapy/test_debuginfo_types.py | 2 +- 2 files changed, 4 insertions(+), 10 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 28e6da2f7..30b39926f 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -3,15 +3,9 @@ from numba import cuda from numba.core import types from numba.cuda.testing import CUDATestCase -from textwrap import dedent -import math import itertools import re import unittest -import warnings -from numba.core.errors import NumbaDebugInfoWarning -from numba.tests.support import ignore_internal_warnings -import numpy as np @skip_on_cudasim("Simulator does not produce debug dumps") @@ -32,7 +26,7 @@ def _check(self, fn, sig, expect): assertfn(match, msg=asm) def test_no_debuginfo_in_asm(self): - @cuda.jit(debug=False) + @cuda.jit(debug=False, opt=False) def foo(x): x[0] = 1 @@ -410,7 +404,7 @@ def f(x, y): self.assertIsNotNone(match, msg=llvm_ir) def test_DW_LANG(self): - @cuda.jit(debug=True) + @cuda.jit(debug=True, opt=False) def foo(): """ CHECK: distinct !DICompileUnit @@ -449,7 +443,7 @@ def test_DILocation(self): """ sig = (types.float64,) - @cuda.jit(sig, debug=True) + @cuda.jit(sig, debug=True, opt=False) def foo(a): """ CHECK-LABEL: define void @{{.+}}foo diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py index 35544b24e..0db72d91b 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py @@ -465,7 +465,7 @@ def sanitize_name(name: str) -> str: f"Test DITypes for {sanitize_name(numba_type.name)}" ): - @cuda.jit((numba_type,), debug=True) + @cuda.jit((numba_type,), debug=True, opt=False) def foo(a): pass From 3b84de70351b704a0290f4d612eaf2eb02631801 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Mon, 4 Aug 2025 08:07:43 -0700 Subject: [PATCH 04/10] Vendor in additional debug info tests This is a hard reset after a rebase went wrong and I lost my place in the PR. --- .../numba/cuda/tests/cudapy/test_debuginfo.py | 176 ++++++++++++++++++ 1 file changed, 176 insertions(+) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 30b39926f..3d75268bd 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -3,9 +3,15 @@ from numba import cuda from numba.core import types from numba.cuda.testing import CUDATestCase +from textwrap import dedent +import math import itertools import re import unittest +import warnings +from numba.core.errors import NumbaDebugInfoWarning +from numba.tests.support import ignore_internal_warnings +import numpy as np @skip_on_cudasim("Simulator does not produce debug dumps") @@ -501,6 +507,176 @@ def foo(a): ir = foo.inspect_llvm()[sig] self.assertFileCheckMatches(ir, foo.__doc__) + def test_missing_source(self): + strsrc = """ + def foo(): + pass + """ + l = dict() + exec(dedent(strsrc), {}, l) + foo = cuda.jit(debug=True, opt=False)(l["foo"]) + + with warnings.catch_warnings(record=True) as w: + warnings.simplefilter("always", NumbaDebugInfoWarning) + ignore_internal_warnings() + foo[1, 1]() + + self.assertEqual(len(w), 1) + found = w[0] + self.assertEqual(found.category, NumbaDebugInfoWarning) + msg = str(found.message) + # make sure the warning contains the right message + self.assertIn("Could not find source for function", msg) + # and refers to the offending function + self.assertIn(str(foo.py_func), msg) + + def test_irregularly_indented_source(self): + @cuda.jit(tuple(), debug=True, opt=False) + def foo(): + # NOTE: THIS COMMENT MUST START AT COLUMN 0 FOR THIS SAMPLE CODE TO BE VALID # noqa: E115, E501 + pass + + with warnings.catch_warnings(record=True) as w: + warnings.simplefilter("always", NumbaDebugInfoWarning) + ignore_internal_warnings() + foo[1, 1]() + + # No warnings + self.assertEqual(len(w), 0) + + self.assertFileCheckMatches( + foo.inspect_llvm()[tuple()], + """ + CHECK: define void @{{.+}}foo + CHECK: !DILocation( + CHECK-SAME: column: 1 + CHECK-NOT: DILocation + """, + ) + + def test_no_if_op_bools_declared(self): + @cuda.jit( + "int64(boolean, boolean)", + debug=True, + opt=False, + _dbg_optnone=True, + device=True, + ) + def choice(cond1, cond2): + """ + CHECK: define void @{{.+}}choices + """ + if cond1 and cond2: + return 1 + else: + return 2 + + ir_content = choice.inspect_llvm()[choice.signatures[0]] + # We should not declare variables used as the condition in if ops. + # See Numba PR #9888: https://github.com/numba/numba/pull/9888 + + for line in ir_content.splitlines(): + if "llvm.dbg.declare" in line: + self.assertNotIn("bool", line) + + def test_llvm_inliner_flag_conflict(self): + # bar will be marked as 'alwaysinline', but when DEBUGINFO_DEFAULT is + # set functions are marked as 'noinline' this results in a conflict. + # baz will be marked as 'noinline' as a result of DEBUGINFO_DEFAULT + + @cuda.jit(forceinline=True) + def bar(x): + return math.sin(x) + + @cuda.jit(forceinline=False) + def baz(x): + return math.cos(x) + + @cuda.jit(opt=True) + def foo(x, y): + """ + CHECK-LABEL: define void @{{.+}}foo + CHECK: call i32 @"[[BAR:.+]]"( + CHECK: call i32 @"[[BAZ:.+]]"( + + CHECK-DAG: declare i32 @"[[BAR]]"({{.+}}alwaysinline + CHECK-DAG: declare i32 @"[[BAZ]]"( + + CHECK-DAG: define linkonce_odr i32 @"[[BAR]]"({{.+}}alwaysinline + CHECK-DAG: define linkonce_odr i32 @"[[BAZ]]"( + """ + a = bar(y) + b = baz(y) + x[0] = a + b + + # check it compiles + with override_config("DEBUGINFO_DEFAULT", 1): + result = cuda.device_array(1, dtype=np.float32) + foo[1, 1](result, np.pi) + result.copy_to_host() + + result_host = math.sin(np.pi) + math.cos(np.pi) + self.assertPreciseEqual(result[0], result_host) + + ir_content = foo.inspect_llvm()[foo.signatures[0]] + self.assertFileCheckMatches(ir_content, foo.__doc__) + + def test_DILocation_versioned_variables(self): + """Tests that DILocation information for versions of variables matches + up to their definition site.""" + + @cuda.jit(debug=True, opt=False) + def foo(dest, n): + """ + CHECK: define void @{{.+}}foo + CHECK: store i64 5, i64* %"c{{.+}} !dbg ![[STORE5:.+]] + CHECK: store i64 1, i64* %"c{{.+}} !dbg ![[STORE1:.+]] + CHECK: [[STORE5]] = !DILocation( + CHECK: [[STORE1]] = !DILocation( + """ + if n: + c = 5 + else: + c = 1 + # prevents inline of return on py310 + py310_defeat1 = 1 # noqa + py310_defeat2 = 2 # noqa + py310_defeat3 = 3 # noqa + py310_defeat4 = 4 # noqa + dest[0] = c + + result = cuda.device_array(1, dtype=np.int32) + foo[1, 1](result, 1) + result.copy_to_host() + self.assertEqual(result[0], 5) + + ir_content = foo.inspect_llvm()[foo.signatures[0]] + self.assertFileCheckMatches(ir_content, foo.__doc__) + + def test_debuginfo_asm(self): + def foo(): + pass + + foo_debug = cuda.jit(debug=True, opt=False)(foo) + foo_debug[1, 1]() + asm = foo_debug.inspect_asm()[foo_debug.signatures[0]] + self.assertFileCheckMatches( + asm, + """ + CHECK: .section{{.+}}.debug + """, + ) + + foo_nodebug = cuda.jit(debug=False)(foo) + foo_nodebug[1, 1]() + asm = foo_nodebug.inspect_asm()[foo_nodebug.signatures[0]] + self.assertFileCheckMatches( + asm, + """ + CHECK-NOT: .section{{.+}}.debug + """, + ) + if __name__ == "__main__": unittest.main() From 7ce3c56e94c5425db5cfce554f4d9f7104490b01 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Mon, 4 Aug 2025 08:13:01 -0700 Subject: [PATCH 05/10] Remove incorrect change from rebase --- numba_cuda/numba/cuda/testing.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/numba_cuda/numba/cuda/testing.py b/numba_cuda/numba/cuda/testing.py index 60712c462..91d1cf243 100644 --- a/numba_cuda/numba/cuda/testing.py +++ b/numba_cuda/numba/cuda/testing.py @@ -177,8 +177,6 @@ def assertFileCheckMatches( self.fail( f"FileCheck failed:\n{matcher.stderr.getvalue()}\n\n" - + f"Check prefixes:\n{check_prefixes}\n\n" - + f"Check patterns:\n{check_patterns}\n" + dump_instructions ) From b41097c15090dc8091044fceec29c9461096783e Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Mon, 4 Aug 2025 08:58:15 -0700 Subject: [PATCH 06/10] Force irregular indenting Test test_irregularly_indented_source checks for debuginfo on irregularly indented lines, but ruff wants to indent this line. I bypassed the precommit checks to force this formatting, hopefully we can prevent ruff from re-indenting this line with `noqa` comments. --- numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 3d75268bd..969e1595f 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -533,7 +533,7 @@ def foo(): def test_irregularly_indented_source(self): @cuda.jit(tuple(), debug=True, opt=False) def foo(): - # NOTE: THIS COMMENT MUST START AT COLUMN 0 FOR THIS SAMPLE CODE TO BE VALID # noqa: E115, E501 +# NOTE: THIS COMMENT MUST START AT COLUMN 0 FOR THIS SAMPLE CODE TO BE VALID # noqa: E115, E501 pass with warnings.catch_warnings(record=True) as w: From 4888fd0f4be8ff2e354d35cbd173f32fa9ac3a9a Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Mon, 4 Aug 2025 10:03:11 -0700 Subject: [PATCH 07/10] Remove irregular indented test A test case in upstream numba contains an irregularly indented comment which I cannot get ruff to ignore, which causes CI to fail. I removed this test case for now, until we can figure out how to allow this irregularly indented comment. `# fmt: on/off/skip` and `# noqa:` comments have not worked thus far. Also clarify description comments for test for conflicting inliner settings. --- .../numba/cuda/tests/cudapy/test_debuginfo.py | 29 ++----------------- 1 file changed, 3 insertions(+), 26 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 969e1595f..c0397c158 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -530,30 +530,6 @@ def foo(): # and refers to the offending function self.assertIn(str(foo.py_func), msg) - def test_irregularly_indented_source(self): - @cuda.jit(tuple(), debug=True, opt=False) - def foo(): -# NOTE: THIS COMMENT MUST START AT COLUMN 0 FOR THIS SAMPLE CODE TO BE VALID # noqa: E115, E501 - pass - - with warnings.catch_warnings(record=True) as w: - warnings.simplefilter("always", NumbaDebugInfoWarning) - ignore_internal_warnings() - foo[1, 1]() - - # No warnings - self.assertEqual(len(w), 0) - - self.assertFileCheckMatches( - foo.inspect_llvm()[tuple()], - """ - CHECK: define void @{{.+}}foo - CHECK: !DILocation( - CHECK-SAME: column: 1 - CHECK-NOT: DILocation - """, - ) - def test_no_if_op_bools_declared(self): @cuda.jit( "int64(boolean, boolean)", @@ -581,8 +557,9 @@ def choice(cond1, cond2): def test_llvm_inliner_flag_conflict(self): # bar will be marked as 'alwaysinline', but when DEBUGINFO_DEFAULT is - # set functions are marked as 'noinline' this results in a conflict. - # baz will be marked as 'noinline' as a result of DEBUGINFO_DEFAULT + # set functions are not marked as 'alwaysinline' and this results in a + # conflict. baz will not be marked as 'alwaysinline' as a result of + # DEBUGINFO_DEFAULT @cuda.jit(forceinline=True) def bar(x): From 1cfebf31beae6fb6866da8a56e5a415f7a2fd176 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Mon, 4 Aug 2025 14:38:06 -0700 Subject: [PATCH 08/10] Add more rigorous checks for math functions Test test_llvm_inliner_flag_conflict calls math functions in device functions. This change incorporates more rigorous checks for the kernels and math functions. --- .../numba/cuda/tests/cudapy/test_debuginfo.py | 35 ++++++++++++++++++- 1 file changed, 34 insertions(+), 1 deletion(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index c0397c158..776cbc52a 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -578,7 +578,6 @@ def foo(x, y): CHECK-DAG: declare i32 @"[[BAR]]"({{.+}}alwaysinline CHECK-DAG: declare i32 @"[[BAZ]]"( - CHECK-DAG: define linkonce_odr i32 @"[[BAR]]"({{.+}}alwaysinline CHECK-DAG: define linkonce_odr i32 @"[[BAZ]]"( """ @@ -598,6 +597,40 @@ def foo(x, y): ir_content = foo.inspect_llvm()[foo.signatures[0]] self.assertFileCheckMatches(ir_content, foo.__doc__) + # Check that the device functions call the appropriate device + # math functions and have the correct attributes. + self.assertFileCheckMatches( + ir_content, + """ + CHECK: define linkonce_odr i32 @{{.+}}bar + CHECK-SAME: alwaysinline + CHECK-NEXT: { + CHECK-NEXT: {{.*}}: + CHECK-NEXT: br label {{.*}} + CHECK-NEXT: {{.*}}: + CHECK-NEXT: call double @"__nv_sin" + CHECK-NEXT: store double {{.*}}, double* {{.*}} + CHECK-NEXT: ret i32 0 + CHECK-NEXT: } + """, + ) + + self.assertFileCheckMatches( + ir_content, + """ + CHECK: define linkonce_odr i32 @{{.+}}baz + CHECK-NOT: alwaysinline + CHECK-NEXT: { + CHECK-NEXT: {{.*}}: + CHECK-NEXT: br label {{.*}} + CHECK-NEXT: {{.*}}: + CHECK-NEXT: call double @"__nv_cos" + CHECK-NEXT: store double {{.*}}, double* {{.*}} + CHECK-NEXT: ret i32 0 + CHECK-NEXT: } + """, + ) + def test_DILocation_versioned_variables(self): """Tests that DILocation information for versions of variables matches up to their definition site.""" From 502ff7619761b5cc75fc59b4bff5da84d4cec0e3 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Wed, 6 Aug 2025 08:16:00 -0700 Subject: [PATCH 09/10] Add checks for line numbers matching Python source Inspects the Python source and ensures that the line number information in the LLVM IR debuginfo metadata matches the original line numbers. --- .../numba/cuda/tests/cudapy/test_debuginfo.py | 118 +++++++++++++++++- 1 file changed, 113 insertions(+), 5 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 776cbc52a..f78315cde 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -1,3 +1,4 @@ +from collections import namedtuple from numba.cuda.tests.support import override_config, captured_stdout from numba.cuda.testing import skip_on_cudasim from numba import cuda @@ -12,6 +13,7 @@ from numba.core.errors import NumbaDebugInfoWarning from numba.tests.support import ignore_internal_warnings import numpy as np +import inspect @skip_on_cudasim("Simulator does not produce debug dumps") @@ -648,13 +650,12 @@ def foo(dest, n): c = 5 else: c = 1 - # prevents inline of return on py310 - py310_defeat1 = 1 # noqa - py310_defeat2 = 2 # noqa - py310_defeat3 = 3 # noqa - py310_defeat4 = 4 # noqa dest[0] = c + foo_source_lines, foo_source_lineno = inspect.getsourcelines( + foo.py_func + ) + result = cuda.device_array(1, dtype=np.int32) foo[1, 1](result, 1) result.copy_to_host() @@ -663,6 +664,113 @@ def foo(dest, n): ir_content = foo.inspect_llvm()[foo.signatures[0]] self.assertFileCheckMatches(ir_content, foo.__doc__) + # Collect lines pertaining to the function `foo` and debuginfo + # metadata + lines = ir_content.splitlines() + debuginfo_equals = re.compile(r"!(\d+) = ") + debug_info_lines = list( + filter(lambda x: debuginfo_equals.search(x), lines) + ) + + function_start_regex = re.compile(r"define void @.+foo") + function_start_lines = list( + filter( + lambda x: function_start_regex.search(x[1]), enumerate(lines) + ) + ) + function_end_lines = list( + filter(lambda x: x[1] == "}", enumerate(lines)) + ) + foo_ir_lines = lines[ + function_start_lines[0][0] : function_end_lines[0][0] + ] + + # Check the if condition's debuginfo + cond_branch = list(filter(lambda x: "br i1" in x, foo_ir_lines)) + self.assertEqual(len(cond_branch), 1) + self.assertIn("!dbg", cond_branch[0]) + cond_branch_dbginfo_node = cond_branch[0].split("!dbg")[1].strip() + cond_branch_dbginfos = list( + filter( + lambda x: cond_branch_dbginfo_node + " = " in x, + debug_info_lines, + ) + ) + self.assertEqual(len(cond_branch_dbginfos), 1) + cond_branch_dbginfo = cond_branch_dbginfos[0] + + # Check debuginfo for the store instructions + store_1_lines, store_5_lines = [], [] + for line in foo_ir_lines: + if "store i64 1" in line: + store_1_lines.append(line) + elif "store i64 5" in line: + store_5_lines.append(line) + + self.assertEqual(len(store_1_lines), 2) + self.assertEqual(len(store_5_lines), 2) + + store_1_dbginfo_set = set( + map(lambda x: x.split("!dbg")[1].strip(), store_1_lines) + ) + store_5_dbginfo_set = set( + map(lambda x: x.split("!dbg")[1].strip(), store_5_lines) + ) + self.assertEqual(len(store_1_dbginfo_set), 1) + self.assertEqual(len(store_5_dbginfo_set), 1) + store_1_dbginfo_node = store_1_dbginfo_set.pop() + store_5_dbginfo_node = store_5_dbginfo_set.pop() + store_1_dbginfos = list( + filter( + lambda x: store_1_dbginfo_node + " = " in x, debug_info_lines + ) + ) + store_5_dbginfos = list( + filter( + lambda x: store_5_dbginfo_node + " = " in x, debug_info_lines + ) + ) + self.assertEqual(len(store_1_dbginfos), 1) + self.assertEqual(len(store_5_dbginfos), 1) + store_1_dbginfo = store_1_dbginfos[0] + store_5_dbginfo = store_5_dbginfos[0] + + # Ensure the line numbers match what we expect based on the Python source + line_number_regex = re.compile(r"line: (\d+)") + LineNumbers = namedtuple( + "LineNumbers", ["cond_branch", "store_5", "store_1"] + ) + line_number_matches = LineNumbers( + *map( + lambda x: line_number_regex.search(x), + [cond_branch_dbginfo, store_5_dbginfo, store_1_dbginfo], + ) + ) + self.assertTrue( + all( + map( + lambda x: x is not None, + line_number_matches, + ) + ) + ) + line_numbers = LineNumbers( + *map( + lambda x: int(x.group(1)), + line_number_matches, + ) + ) + source_line_numbers = LineNumbers( + *map( + lambda x: x[0] + foo_source_lineno, + filter( + lambda x: "c = " in x[1] or "if n:" in x[1], + enumerate(foo_source_lines), + ), + ) + ) + self.assertEqual(line_numbers, source_line_numbers) + def test_debuginfo_asm(self): def foo(): pass From 5f9162a3675cd172ccd4256771f8976fd63762e7 Mon Sep 17 00:00:00 2001 From: Asher Mancinelli Date: Wed, 6 Aug 2025 08:38:13 -0700 Subject: [PATCH 10/10] Use more consistent filtering --- numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index f78315cde..73857965b 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -700,12 +700,8 @@ def foo(dest, n): cond_branch_dbginfo = cond_branch_dbginfos[0] # Check debuginfo for the store instructions - store_1_lines, store_5_lines = [], [] - for line in foo_ir_lines: - if "store i64 1" in line: - store_1_lines.append(line) - elif "store i64 5" in line: - store_5_lines.append(line) + store_1_lines = list(filter(lambda x: "store i64 1" in x, foo_ir_lines)) + store_5_lines = list(filter(lambda x: "store i64 5" in x, foo_ir_lines)) self.assertEqual(len(store_1_lines), 2) self.assertEqual(len(store_5_lines), 2)