removing kernel finalizers

Author: isVoid

numba_cuda/numba/cuda/dispatcher.py

@@ -1,5 +1,4 @@
 import numpy as np
-import weakref
 import os
 import re
 import sys
@@ -1025,7 +1024,6 @@ def compile(self, sig):
                 raise RuntimeError("Compilation disabled")
 
             kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
-            weakref.finalize(kernel, _kernel_finalize_callback, kernel)
             # We call bind to force codegen, so that there is a cubin to cache
             kernel.bind()
             self._cache.save_overload(sig, kernel)
@@ -1151,18 +1149,3 @@ def _reduce_states(self):
         """
         return dict(py_func=self.py_func,
                     targetoptions=self.targetoptions)
-
-
-def _kernel_finalize_callback(kernel):
-    module = kernel.library.get_cufunc().module
-    try:
-        if driver.USE_NV_BINDING:
-            key = module.handle
-        else:
-            key = module.handle.value
-    except ReferenceError:
-        return
-
-    ctx = cuda.current_context()
-    if key in ctx.modules:
-        del ctx.modules[key]

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

@@ -1,14 +1,23 @@
-import gc
+import unittest
 
 import numpy as np
 
-from numba import cuda
+from numba import cuda, config
 from numba.cuda.cudadrv.linkable_code import CUSource
 from numba.cuda.testing import CUDATestCase
 
 from cuda.bindings.driver import cuModuleGetGlobal, cuMemcpyHtoD
 
 
+def wipe_all_modules_in_context():
+    ctx = cuda.current_context()
+    ctx.modules.clear()
+
+
+@unittest.skipIf(
+    config.CUDA_USE_NVIDIA_BINDING,
+    "NV binding support superceded by cuda.bindings."
+)
 class TestModuleCallbacksBasic(CUDATestCase):
 
     def test_basic(self):
@@ -33,13 +42,9 @@ def kernel():
         self.assertEqual(counter, 1)
         kernel[1, 1]() # cached
         self.assertEqual(counter, 1)
-        breakpoint()
-        del kernel
-        gc.collect()
-        cuda.current_context().deallocations.clear()
+
+        wipe_all_modules_in_context()
         self.assertEqual(counter, 0)
-        # We don't have a way to explicitly evict kernel and its modules at
-        # the moment.
 
     def test_different_argtypes(self):
         counter = 0
@@ -66,11 +71,8 @@ def kernel(arg):
         kernel[1, 1](3.14)  # (float64)->() : module 2
         self.assertEqual(counter, 2)
 
-        # del kernel
-        # gc.collect()
-        # cuda.current_context().deallocations.clear()
-        # self.assertEqual(counter, 0) # We don't have a way to explicitly
-        # evict kernel and its modules at the moment.
+        wipe_all_modules_in_context()
+        self.assertEqual(counter, 0)
 
     def test_two_kernels(self):
         counter = 0
@@ -98,11 +100,8 @@ def kernel2():
         kernel2[1, 1]()
         self.assertEqual(counter, 2)
 
-        # del kernel
-        # gc.collect()
-        # cuda.current_context().deallocations.clear()
-        # self.assertEqual(counter, 0) # We don't have a way to explicitly
-        # evict kernel and its modules at the moment.
+        wipe_all_modules_in_context()
+        self.assertEqual(counter, 0)
 
 
 class TestModuleCallbacks(CUDATestCase):
@@ -137,10 +136,6 @@ def teardown(mod, stream):
         self.lib = CUSource(
             module, setup_callback=set_forty_two, teardown_callback=teardown)
 
-    def tearDown(self):
-        super().tearDown()
-        del self.lib
-
     def test_decldevice_arg(self):
         get_num = cuda.declare_device("get_num", "int32()", link=[self.lib])