add test for MemoryPointer inputs
@@ -7,6 +7,7 @@ import unittest
import cupy
from cupy import testing
from cupy.cuda import compiler
+from cupy.cuda import memory
_test_source1 = r'''
@@ -316,6 +317,16 @@ __global__ void my_func(double* input, int N) {
}
'''
+test_cast = r'''
+extern "C" __global__ void my_func(void* input, int N) {
+ double* arr = (double*)(input);
+ unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
+ if (x < N) {
+ arr[x] = 3.0 * arr[x] - 8.0;
+ }
+}
+'''
+
if 'CUPY_CACHE_DIR' in os.environ:
_old_cache_dir = os.environ['CUPY_CACHE_DIR']
_is_cache_env_var_set = True
@@ -713,6 +724,19 @@ class TestRaw(unittest.TestCase):
match='named symbol not found'):
mod.get_function('my_sqrt<double>')
+ def test_raw_pointer(self):
+ mod = cupy.RawModule(code=test_cast, backend=self.backend)
+ ker = mod.get_function('my_func')
+
+ a = cupy.ones((100,), dtype=cupy.float64)
+ arr = memory.alloc(100 * a.dtype.itemsize)
+ arr.copy_from(a.data, 100 * a.dtype.itemsize) # one-initialize
+ b = cupy.ndarray((100,), cupy.float64, memptr=arr)
+
+ ker((1,), (100,), (arr, 100))
+ a = 3. * a - 8.
+ assert (a == b).all()
+
@testing.multi_gpu(2)
def test_context_switch_RawKernel(self):
# run test_basic() on another device
