test: enable fail-on-warn and clean up resulting failures (#529)

Author: cpcloud

numba_cuda/numba/cuda/core/typed_passes.py

@@ -70,12 +70,9 @@ def fallback_context(state, msg):
             # this emits a warning containing the error message body in the
             # case of fallback from npm to objmode
             loop_lift = "" if state.flags.enable_looplift else "OUT"
-            msg_rewrite = (
-                "\nCompilation is falling back to object mode "
-                "WITH%s looplifting enabled because %s" % (loop_lift, msg)
-            )
             warnings.warn_explicit(
-                "%s due to: %s" % (msg_rewrite, e),
+                "Compilation is falling back to object mode "
+                f"WITH{loop_lift} looplifting enabled because {msg} due to: {e}",
                 errors.NumbaWarning,
                 state.func_id.filename,
                 state.func_id.firstlineno,

numba_cuda/numba/cuda/core/untyped_passes.py

@@ -66,12 +66,9 @@ def fallback_context(state, msg):
             # this emits a warning containing the error message body in the
             # case of fallback from npm to objmode
             loop_lift = "" if state.flags.enable_looplift else "OUT"
-            msg_rewrite = (
-                "\nCompilation is falling back to object mode "
-                "WITH%s looplifting enabled because %s" % (loop_lift, msg)
-            )
             warnings.warn_explicit(
-                "%s due to: %s" % (msg_rewrite, e),
+                "Compilation is falling back to object mode "
+                f"WITH{loop_lift} looplifting enabled because {msg} due to: {e}",
                 errors.NumbaWarning,
                 state.func_id.filename,
                 state.func_id.firstlineno,

numba_cuda/numba/cuda/cudadrv/driver.py

@@ -1495,7 +1495,7 @@ def _alloc_finalizer(memory_manager, ptr, alloc_key, size):
 
     def core():
         if allocations:
-            del allocations[alloc_key]
+            allocations.pop(alloc_key, None)
         deallocations.add_item(driver.cuMemFree, ptr, size)
 
     return core

numba_cuda/numba/cuda/memory_management/memsys.cu

@@ -89,10 +89,10 @@ extern "C" __global__ void NRT_MemSys_print(void)
   if (TheMSys != nullptr)
   {
     printf("TheMSys->stats.enabled %d\n", TheMSys->stats.enabled);
-    printf("TheMSys->stats.alloc %lu\n", TheMSys->stats.alloc.load());
-    printf("TheMSys->stats.free %lu\n", TheMSys->stats.free.load());
-    printf("TheMSys->stats.mi_alloc %lu\n", TheMSys->stats.mi_alloc.load());
-    printf("TheMSys->stats.mi_free %lu\n", TheMSys->stats.mi_free.load());
+    printf("TheMSys->stats.alloc %zu\n", TheMSys->stats.alloc.load());
+    printf("TheMSys->stats.free %zu\n", TheMSys->stats.free.load());
+    printf("TheMSys->stats.mi_alloc %zu\n", TheMSys->stats.mi_alloc.load());
+    printf("TheMSys->stats.mi_free %zu\n", TheMSys->stats.mi_free.load());
   } else {
     printf("TheMsys is null.\n");
   }

numba_cuda/numba/cuda/simulator/api.py

@@ -22,7 +22,6 @@
 from .kernel import FakeCUDAKernel
 from numba.cuda.core import config
 from numba.cuda.core.sigutils import is_signature
-from warnings import warn
 from ..args import In, Out, InOut  # noqa: F401
 
 
@@ -111,8 +110,8 @@ def synchronize(self):
         pass
 
     def elapsed_time(self, event):
-        warn("Simulator timings are bogus")
-        return 0.0
+        """This is here to preserve the API; the output is meaningless."""
+        return -1.0
 
 
 event = Event

numba_cuda/numba/cuda/simulator/cudadrv/devicearray.py

@@ -8,6 +8,8 @@
 
 from contextlib import contextmanager
 from numba.cuda.np.numpy_support import numpy_version
+from numba.cuda.np import numpy_support
+from numba.cuda import types
 
 import numpy as np
 
@@ -57,8 +59,8 @@ def __wrap_if_fake(self, item):
             return item
 
     def __getattr__(self, attrname):
-        if attrname in dir(self._item._ary):  # For e.g. array size.
-            return self.__wrap_if_fake(getattr(self._item._ary, attrname))
+        if (value := getattr(self._item, attrname, None)) is not None:
+            return self.__wrap_if_fake(value)
         else:
             return self.__wrap_if_fake(self._item.__getitem__(attrname))
 
@@ -109,6 +111,23 @@ def __init__(self, ary, stream=0):
         self._ary = ary
         self.stream = stream
 
+    @property
+    def _numba_type_(self):
+        """
+        Magic attribute expected by Numba to get the numba type that
+        represents this object.
+        """
+        broadcast = 0 in self.strides
+        if self.is_c_contiguous() and not broadcast:
+            layout = "C"
+        elif self.is_f_contiguous() and not broadcast:
+            layout = "F"
+        else:
+            layout = "A"
+
+        dtype = numpy_support.from_dtype(self._ary.dtype)
+        return types.Array(dtype, self._ary.ndim, layout)
+
     @property
     def alloc_size(self):
         return self._ary.nbytes

numba_cuda/numba/cuda/testing.py

@@ -187,44 +187,52 @@ def assertFileCheckMatches(
 
 def skip_on_cudasim(reason):
     """Skip this test if running on the CUDA simulator"""
+    assert isinstance(reason, str)
     return unittest.skipIf(config.ENABLE_CUDASIM, reason)
 
 
-def skip_on_standalone_numba_cuda(reason):
-    """Skip this test if running on standalone numba_cuda"""
-    return unittest.skipIf(not HAS_NUMBA, reason)
+skip_on_standalone_numba_cuda = unittest.skipUnless(
+    HAS_NUMBA, "requires base numba install"
+)
 
 
 def skip_unless_cudasim(reason):
     """Skip this test if running on CUDA hardware"""
+    assert isinstance(reason, str)
     return unittest.skipUnless(config.ENABLE_CUDASIM, reason)
 
 
 def skip_unless_conda_cudatoolkit(reason):
     """Skip test if the CUDA toolkit was not installed by Conda"""
+    assert isinstance(reason, str)
     return unittest.skipUnless(get_conda_ctk_libdir() is not None, reason)
 
 
 def skip_if_external_memmgr(reason):
     """Skip test if an EMM Plugin is in use"""
+    assert isinstance(reason, str)
     return unittest.skipIf(config.CUDA_MEMORY_MANAGER != "default", reason)
 
 
 def skip_under_cuda_memcheck(reason):
+    assert isinstance(reason, str)
     return unittest.skipIf(os.environ.get("CUDA_MEMCHECK") is not None, reason)
 
 
 def skip_without_nvdisasm(reason):
+    assert isinstance(reason, str)
     nvdisasm_path = shutil.which("nvdisasm")
     return unittest.skipIf(nvdisasm_path is None, reason)
 
 
 def skip_with_nvdisasm(reason):
+    assert isinstance(reason, str)
     nvdisasm_path = shutil.which("nvdisasm")
     return unittest.skipIf(nvdisasm_path is not None, reason)
 
 
 def skip_on_arm(reason):
+    assert isinstance(reason, str)
     cpu = platform.processor()
     is_arm = cpu.startswith("arm") or cpu.startswith("aarch")
     return unittest.skipIf(is_arm, reason)
@@ -236,6 +244,7 @@ def skip_on_wsl2(reason):
     Detection is based on the kernel release string, which typically contains
     "microsoft-standard-WSL2" on WSL2 systems.
     """
+    assert isinstance(reason, str)
     rel = platform.release().lower()
     is_wsl2 = ("microsoft-standard-wsl2" in rel) or ("wsl2" in rel)
     return unittest.skipIf(is_wsl2, reason)
@@ -269,6 +278,7 @@ def skip_if_curand_kernel_missing(fn):
 
 def skip_if_mvc_enabled(reason):
     """Skip a test if Minor Version Compatibility is enabled"""
+    assert isinstance(reason, str)
     return unittest.skipIf(
         config.CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY, reason
     )
@@ -321,6 +331,7 @@ def skip_if_cudadevrt_missing(fn):
 
 
 def skip_if_nvjitlink_missing(reason):
+    assert isinstance(reason, str)
     return unittest.skipIf(not driver._have_nvjitlink(), reason)
 
 

numba_cuda/numba/cuda/tests/cudadrv/test_cuda_libraries.py

@@ -9,7 +9,7 @@ class TestIsFP16Supported(CUDATestCase):
     def test_is_fp16_supported(self):
         self.assertTrue(cuda.is_float16_supported())
 
-    @skip_on_cudasim
+    @skip_on_cudasim("fp16 not available in sim")
     @skip_unless_cc_53
     def test_device_supports_float16(self):
         self.assertTrue(cuda.get_current_device().supports_float16)

numba_cuda/numba/cuda/tests/cudadrv/test_is_fp16.py

@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 import numpy as np
-import warnings
+import pytest
 from numba.cuda.testing import unittest
 from numba.cuda.testing import (
     skip_on_cudasim,
@@ -12,7 +12,6 @@
 from numba.cuda.testing import CUDATestCase, test_data_dir
 from numba.cuda.cudadrv.driver import CudaAPIError, _Linker, LinkerError
 from numba.cuda import require_context
-from numba.cuda.tests.support import ignore_internal_warnings
 from numba import cuda
 from numba.cuda import void, float64, int64, int32, float32
 from numba.cuda.typing.typeof import typeof
@@ -172,8 +171,7 @@ def test_linking_cu_log_warning(self):
 
         link = str(test_data_dir / "warn.cu")
 
-        with warnings.catch_warnings(record=True) as w:
-            ignore_internal_warnings()
+        with pytest.warns(UserWarning) as w:
 
             @cuda.jit("void(int32)", link=[link])
             def kernel(x):