diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 3f734ad71..73857965b 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py @@ -1,11 +1,19 @@ +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 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 +import inspect @skip_on_cudasim("Simulator does not produce debug dumps") @@ -26,7 +34,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 @@ -404,7 +412,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 @@ -443,7 +451,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 @@ -501,6 +509,288 @@ 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_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 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): + 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__) + + # 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.""" + + @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 + 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() + self.assertEqual(result[0], 5) + + 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 = 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) + + 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 + + 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() 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