@@ -4,6 +4,7 @@
 from itertools import product
 from numba import vectorize
 from numba import cuda, int32, float32, float64
+from numba.cuda.cudadrv.driver import CudaAPIError, driver
 from numba.cuda.testing import skip_on_cudasim
 from numba.cuda.testing import CUDATestCase
 import unittest
@@ -227,6 +228,56 @@ def bar(x):
 
         self.assertEqual(bar.__name__, 'bar')
 
+    def test_no_transfer_for_device_data(self):
+        # Initialize test data on the device prior to banning host <-> device
+        # transfer
+
+        noise = np.random.randn(1, 3, 64, 64).astype(np.float32)
+        noise = cuda.to_device(noise)
+
+        # A mock of a CUDA function that always raises a CudaAPIError
+
+        def raising_transfer(*args, **kwargs):
+            raise CudaAPIError(999, 'Transfer not allowed')
+
+        # Use the mock for transfers between the host and device
+
+        old_HtoD = getattr(driver, 'cuMemcpyHtoD', None)
+        old_DtoH = getattr(driver, 'cuMemcpyDtoH', None)
+
+        setattr(driver, 'cuMemcpyHtoD', raising_transfer)
+        setattr(driver, 'cuMemcpyDtoH', raising_transfer)
+
+        # Ensure that the mock functions are working as expected
+
+        with self.assertRaisesRegex(CudaAPIError, "Transfer not allowed"):
+            noise.copy_to_host()
+
+        with self.assertRaisesRegex(CudaAPIError, "Transfer not allowed"):
+            cuda.to_device([1])
+
+        try:
+            # Check that defining and calling a ufunc with data on the device
+            # induces no transfers
+
+            @vectorize(['float32(float32)'], target='cuda')
+            def func(noise):
+                return noise + 1.0
+
+            func(noise)
+        finally:
+            # Replace our mocks with the original implementations. If there was
+            # no original implementation, simply remove ours.
+
+            if old_HtoD is not None:
+                setattr(driver, 'cuMemcpyHtoD', old_HtoD)
+            else:
+                del driver.cuMemcpyHtoD
+            if old_DtoH is not None:
+                setattr(driver, 'cuMemcpyDtoH', old_DtoH)
+            else:
+                del driver.cuMemcpyDtoH
+
 
 if __name__ == '__main__':
     unittest.main()
@@ -105,7 +105,10 @@ def _fill_argtypes(self):
         """
         for i, ary in enumerate(self.arrays):
             if ary is not None:
-                self.argtypes[i] = np.asarray(ary).dtype
+                dtype = getattr(ary, 'dtype')
+                if dtype is None:
+                    dtype = np.asarray(ary).dtype
+                self.argtypes[i] = dtype
 
     def _resolve_signature(self):
         """Resolve signature.