From 1a07b5bc0ba5e687beb6ee58ad23a90701002a11 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Fri, 3 Dec 2021 14:40:06 +0000 Subject: [PATCH 1/3] CUDA Linker test: remove unneeded @require_context Tests that operate through the normal APIs (e.g. `@cuda.jit`) don't need to have a prior context available. --- numba/cuda/tests/cudadrv/test_linker.py | 6 ------ 1 file changed, 6 deletions(-) diff --git a/numba/cuda/tests/cudadrv/test_linker.py b/numba/cuda/tests/cudadrv/test_linker.py index 0a4dc1f8ad2..04f75c6f9d6 100644 --- a/numba/cuda/tests/cudadrv/test_linker.py +++ b/numba/cuda/tests/cudadrv/test_linker.py @@ -66,7 +66,6 @@ def test_linker_basic(self): linker = Linker.new() del linker - @require_context def test_linking(self): global bar # must be a global; other it is recognized as a freevar bar = cuda.declare_device('bar', 'int32(int32)') @@ -85,7 +84,6 @@ def foo(x, y): self.assertTrue(A[0] == 123 + 2 * 321) - @require_context def test_try_to_link_nonexistent(self): with self.assertRaises(LinkerError) as e: @cuda.jit('void(int32[::1])', link=['nonexistent.a']) @@ -93,7 +91,6 @@ def f(x): x[0] = 0 self.assertIn('nonexistent.a not found', e.exception.args) - @require_context def test_set_registers_no_max(self): """Ensure that the jitted kernel used in the test_set_registers_* tests uses more than 57 registers - this ensures that test_set_registers_* @@ -103,19 +100,16 @@ def test_set_registers_no_max(self): compiled = compiled.specialize(np.empty(32), *range(6)) self.assertGreater(compiled.get_regs_per_thread(), 57) - @require_context def test_set_registers_57(self): compiled = cuda.jit(max_registers=57)(func_with_lots_of_registers) compiled = compiled.specialize(np.empty(32), *range(6)) self.assertLessEqual(compiled.get_regs_per_thread(), 57) - @require_context def test_set_registers_38(self): compiled = cuda.jit(max_registers=38)(func_with_lots_of_registers) compiled = compiled.specialize(np.empty(32), *range(6)) self.assertLessEqual(compiled.get_regs_per_thread(), 38) - @require_context def test_set_registers_eager(self): sig = void(float64[::1], int64, int64, int64, int64, int64, int64) compiled = cuda.jit(sig, max_registers=38)(func_with_lots_of_registers) From 7fca40f0cf817487087ba122d54a9c53c47eea9a Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Fri, 3 Dec 2021 14:50:11 +0000 Subject: [PATCH 2/3] CUDA: Fix linking linking PTX for lazy compilation (Issue #7607) When no signature was provided, the `link` kwarg was accidentally dropped in the wrapper function calling `jit` inside the decorator. --- numba/cuda/decorators.py | 3 ++- numba/cuda/tests/cudadrv/test_linker.py | 15 +++++++++++++-- 2 files changed, 15 insertions(+), 3 deletions(-) diff --git a/numba/cuda/decorators.py b/numba/cuda/decorators.py index f3d7f77c40d..e5c1dfaf2d3 100644 --- a/numba/cuda/decorators.py +++ b/numba/cuda/decorators.py @@ -107,7 +107,8 @@ def autojitwrapper(func): fastmath=fastmath) else: def autojitwrapper(func): - return jit(func, device=device, debug=debug, opt=opt, **kws) + return jit(func, device=device, debug=debug, opt=opt, + link=link, **kws) return autojitwrapper # func_or_sig is a function diff --git a/numba/cuda/tests/cudadrv/test_linker.py b/numba/cuda/tests/cudadrv/test_linker.py index 04f75c6f9d6..bf1e727fe3b 100644 --- a/numba/cuda/tests/cudadrv/test_linker.py +++ b/numba/cuda/tests/cudadrv/test_linker.py @@ -66,13 +66,18 @@ def test_linker_basic(self): linker = Linker.new() del linker - def test_linking(self): + def _test_linking(self, eager=False): global bar # must be a global; other it is recognized as a freevar bar = cuda.declare_device('bar', 'int32(int32)') link = os.path.join(os.path.dirname(__file__), 'data', 'jitlink.ptx') - @cuda.jit('void(int32[:], int32[:])', link=[link]) + if eager: + args = ['void(int32[:], int32[:])'] + else: + args = [] + + @cuda.jit(*args, link=[link]) def foo(x, y): i = cuda.grid(1) x[i] += bar(y[i]) @@ -84,6 +89,12 @@ def foo(x, y): self.assertTrue(A[0] == 123 + 2 * 321) + def test_linking_lazy_compile(self): + self._test_linking(eager=False) + + def test_linking_eager_compile(self): + self._test_linking(eager=True) + def test_try_to_link_nonexistent(self): with self.assertRaises(LinkerError) as e: @cuda.jit('void(int32[::1])', link=['nonexistent.a']) From 04ac1376d61ae594da004e7c302d4075126e8802 Mon Sep 17 00:00:00 2001 From: Graham Markall <535640+gmarkall@users.noreply.github.com> Date: Mon, 6 Dec 2021 10:48:50 +0000 Subject: [PATCH 3/3] Make eager a positional arg not a kwarg PR #7619 feedback. Co-authored-by: stuartarchibald --- numba/cuda/tests/cudadrv/test_linker.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/numba/cuda/tests/cudadrv/test_linker.py b/numba/cuda/tests/cudadrv/test_linker.py index bf1e727fe3b..88dd65b919d 100644 --- a/numba/cuda/tests/cudadrv/test_linker.py +++ b/numba/cuda/tests/cudadrv/test_linker.py @@ -66,7 +66,7 @@ def test_linker_basic(self): linker = Linker.new() del linker - def _test_linking(self, eager=False): + def _test_linking(self, eager): global bar # must be a global; other it is recognized as a freevar bar = cuda.declare_device('bar', 'int32(int32)')