OpenCL kernel run tests

test/helpers.py

@@ -27,3 +27,5 @@ def cuda_compile(prg, options, f, check):
   status = f.compile(prog, len(options), to_char_p_p(options))
   if status != 0: raise RuntimeError(f"compile failed: {get_bytes(prog, f.getLogSize, f.getLog, check)}")
   return get_bytes(prog, f.getCodeSize, f.getCode, check)
+
+def ctype_buffer(dtype, sz: int, data): return (dtype * sz)(*data)

test/test_opencl.py

@@ -1,7 +1,7 @@
 from typing import Optional
 import ctypes
 import unittest
-from helpers import to_char_p_p
+from helpers import ctype_buffer, to_char_p_p
 import gpuctypes.opencl as cl
 
 def check(status, info:Optional[str]=None):
@@ -69,5 +69,33 @@ def test_create_program(self):
     assert binary_sizes[0] > 0
     assert len(binary_pointers[0]) > 0
 
+  def test_run_program(self):
+    program = cl_compile(self.context, self.device_array, """
+    __kernel void vector_add(__global int *A, __global const int *B, __global const int *C) {
+      int i = get_global_id(0);
+      A[i] = B[i] + C[i];
+    }""")
+    status = ctypes.c_int32()
+    kernel = cl.clCreateKernel(program, b"vector_add", ctypes.byref(status))
+    check(status.value)
+    assert kernel is not None
+
+    A = ctype_buffer(ctypes.c_int32, 5, ())
+    B = ctype_buffer(ctypes.c_int32, 5, (1,2,3,4,5))
+    C = ctype_buffer(ctypes.c_int32, 5, (5,4,3,2,1))
+
+    bufA = cl.clCreateBuffer(self.context, cl.CL_MEM_WRITE_ONLY, ctypes.sizeof(A), None, ctypes.byref(status))
+    bufB = cl.clCreateBuffer(self.context, cl.CL_MEM_READ_ONLY | cl.CL_MEM_COPY_HOST_PTR, ctypes.sizeof(B), ctypes.byref(B), ctypes.byref(status))
+    bufC = cl.clCreateBuffer(self.context, cl.CL_MEM_READ_ONLY | cl.CL_MEM_COPY_HOST_PTR, ctypes.sizeof(C), ctypes.byref(C), ctypes.byref(status))
+    cl.clSetKernelArg(kernel, 0, ctypes.sizeof(bufA), ctypes.byref(bufA))
+    cl.clSetKernelArg(kernel, 1, ctypes.sizeof(bufB), ctypes.byref(bufB))
+    cl.clSetKernelArg(kernel, 2, ctypes.sizeof(bufC), ctypes.byref(bufC))
+
+    global_size = ctypes.c_size_t(5)
+    cl.clEnqueueNDRangeKernel(self.queue, kernel, 1, None, ctypes.byref(global_size), None, 0, None, None)
+    cl.clFinish(self.queue)
+    cl.clEnqueueReadBuffer(self.queue, bufA, cl.CL_TRUE, 0, ctypes.sizeof(A), ctypes.byref(A), 0, None, None)
+    self.assertEqual(list(A), [6, 6, 6, 6, 6])
+
 if __name__ == '__main__':
   unittest.main()