mocking object type using cuda.core objects

Author: isVoid

numba_cuda/numba/cuda/codegen.py

@@ -6,6 +6,8 @@
 from numba.cuda.cudadrv.libs import get_cudalib
 from numba.cuda.cudadrv.linkable_code import LinkableCode
 
+from cuda.core.experimental import ObjectCode
+
 import os
 import subprocess
 import tempfile
@@ -250,14 +252,15 @@ def get_cufunc(self):
             return cufunc
 
         cubin = self.get_cubin(cc=device.compute_capability)
-        module = ctx.create_module_image(cubin)
+
+        # just a mock, https://github.com/NVIDIA/numba-cuda/pull/133 will
+        # formalize the object code interface
+        obj_code = ObjectCode.from_cubin(cubin)
+        cufunc = obj_code.get_kernel(self._entry_name)
 
         # Init
         for init_fn in self._init_functions:
-            init_fn(module)
-
-        # Load
-        cufunc = module.get_function(self._entry_name)
+            init_fn(obj_code)
 
         # Populate caches
         self._cufunc_cache[device.id] = cufunc

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

@@ -4,7 +4,11 @@
 from numba.cuda.cudadrv.linkable_code import CUSource
 from numba.cuda.testing import CUDATestCase
 
-from cuda.bindings.driver import cuModuleGetGlobal, cuMemcpyHtoD
+from cuda.bindings.driver import (
+    cuModuleGetGlobal,
+    cuMemcpyHtoD,
+    cuLibraryGetModule
+)
 
 
 class TestModuleInitCallback(CUDATestCase):
@@ -21,11 +25,13 @@ def setUp(self):
 }
 """
 
-        def set_fourty_two(mod):
+        def set_fourty_two(obj):
             # Initialize 42 to global variable `num`
-            res, dptr, size = cuModuleGetGlobal(
-                mod.handle.value, "num".encode()
-            )
+            culib = obj._handle
+            res, mod = cuLibraryGetModule(culib)
+            self.assertEqual(res, 0)
+
+            res, dptr, size = cuModuleGetGlobal(mod, "num".encode())
             self.assertEqual(res, 0)
             self.assertEqual(size, 4)