drop to ctypes ptr in nrt when using nv binding

Author: brandon-b-miller

numba_cuda/numba/cuda/cudadrv/driver.py

@@ -2553,7 +2553,6 @@ def launch_kernel(cufunc_handle,
                   hstream,
                   args,
                   cooperative=False):
-
     param_ptrs = [addressof(arg) for arg in args]
     params = (c_void_p * len(param_ptrs))(*param_ptrs)
 

numba_cuda/numba/cuda/runtime/nrt.py

@@ -11,6 +11,7 @@
 from numba.cuda.utils import _readenv
 
 
+
 # Check environment variable or config for NRT statistics enablement
 NRT_STATS = (
     _readenv("NUMBA_CUDA_NRT_STATS", bool, False) or
@@ -117,6 +118,13 @@ def _single_thread_launch(self, module, stream, name, params=()):
         if stream is None:
             stream = cuda.default_stream()
 
+        if config.CUDA_USE_NVIDIA_BINDING:
+            from numba.cuda.cudadrv.drvapi import cu_device_ptr
+            from cuda.cuda import CUdeviceptr
+            params = tuple(
+                cu_device_ptr.from_address(ptr.getPtr()) if isinstance(ptr, CUdeviceptr) else ptr for ptr in params
+            )
+
         func = module.get_function(name)
         launch_kernel(
             func.handle,
@@ -294,6 +302,7 @@ def set_memsys_to_module(self, module, stream=None):
             raise RuntimeError(
                 "Please allocate NRT Memsys first before setting to module.")
 
+
         self._single_thread_launch(
             module,
             stream,