test_boundscheck: fix test_no_cuda_boundscheck

This test used CUDA functionality, so it needs the SerialMixin to
prevent it running in a child process after the test runner already
initialized CUDA in the parent process. It is moved into its own class
to add the SerialMixin, to preserve the ability to run other tests from
the same class to run in parallel.

It also uses a CUDA kernel without a launch configuration, which will
soon (pending PR numba#5061) be an error, so we add a launch configuration
to it.

numba/tests/test_boundscheck.py

@@ -3,6 +3,7 @@
 import numpy as np
 
 from numba.compiler import compile_isolated, DEFAULT_FLAGS
+from numba.cuda.testing import SerialMixin
 from numba import typeof, config, cuda, njit
 from numba.types import float64
 from numba import unittest_support as unittest
@@ -104,6 +105,15 @@ def test_fancy_indexing_boundscheck(self):
         # Doesn't raise
         boundscheck(b)
 
+    def tearDown(self):
+        config.BOUNDSCHECK = self.old_boundscheck
+
+
+class TestNoCudaBoundsCheck(SerialMixin, unittest.TestCase):
+    def setUp(self):
+        self.old_boundscheck = config.BOUNDSCHECK
+        config.BOUNDSCHECK = None
+
     @unittest.skipIf(not cuda.is_available(), "NO CUDA")
     def test_no_cuda_boundscheck(self):
         with self.assertRaises(NotImplementedError):
@@ -121,7 +131,7 @@ def func2(x, a):
             # Out of bounds but doesn't raise (it does raise in the simulator,
             # so skip there)
             if not config.ENABLE_CUDASIM:
-                func2(x, a)
+                func2[1, 1](x, a)
 
     def tearDown(self):
         config.BOUNDSCHECK = self.old_boundscheck