cl multi kernel

Author: cirosantilli

opencl/README.md

@@ -11,6 +11,7 @@
     1.  [clinfo](clinfo.c)
     1.  [Preprocessor](matmul.c)
     1.  [Write buffer](write_buffer.c)
+    1.  [Multi-kernel](multi_kernel.c)
     1.  [Benchmark examples](benchmark-examples.md)
         1.  [Matrix multiplication](matmul.c)
 1.  Tools

opencl/common.h

@@ -42,26 +42,35 @@ char* common_read_file(const char *path) {
     return buffer;
 }
 
-void common_create_kernel(
+void common_create_program(
     Common *common,
     const char *source,
-    const char *options
+    const char *options,
+    cl_program *program
 ) {
-    cl_int ret;
     char *err;
+    cl_int ret;
     size_t err_len;
 
+    *program = clCreateProgramWithSource(common->context, 1, &source, NULL, NULL);
+    ret = clBuildProgram(*program, 1, &(common->device), options, NULL, NULL);
+    if (CL_SUCCESS != ret) {
+        clGetProgramBuildInfo(*program, common->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &err_len);
+        err = malloc(err_len);
+        clGetProgramBuildInfo(*program, common->device, CL_PROGRAM_BUILD_LOG, err_len, err, NULL);
+        fprintf(stderr, "error: kernel build:\n%s\n", err);
+        free(err);
+        exit(EXIT_FAILURE);
+    }
+}
+
+void common_create_kernel(
+    Common *common,
+    const char *source,
+    const char *options
+) {
     if (NULL != source) {
-        common->program = clCreateProgramWithSource(common->context, 1, &source, NULL, NULL);
-        ret = clBuildProgram(common->program, 1, &(common->device), options, NULL, NULL);
-        if (CL_SUCCESS != ret) {
-            clGetProgramBuildInfo(common->program, common->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &err_len);
-            err = malloc(err_len);
-            clGetProgramBuildInfo(common->program, common->device, CL_PROGRAM_BUILD_LOG, err_len, err, NULL);
-            fprintf(stderr, "error: kernel build:\n%s\n", err);
-            free(err);
-            exit(EXIT_FAILURE);
-        }
+        common_create_program(common, source, options, &common->program);
         common->kernel = clCreateKernel(common->program, "main", NULL);
     } else {
         common->kernel = NULL;

opencl/multi_kernel.c

@@ -0,0 +1,46 @@
+/*
+We can pass data through multiple kernels with multiple
+clEnqueueNDRangeKernel calls on a single command queue.
+*/
+
+#include "common.h"
+
+int main(void) {
+    const char *source =
+        "__kernel void add(__global int *out) {\n"
+        "    out[get_global_id(0)]++;\n"
+        "}\n"
+        "__kernel void mul(__global int *out) {\n"
+        "    out[get_global_id(0)] *= 2;\n"
+        "}\n"
+    ;
+    cl_int input[] = {1, 2};
+    cl_kernel kernel_add, kernel_mul;
+    cl_mem buffer;
+    cl_program program;
+    Common common;
+    const size_t global_work_size = sizeof(input) / sizeof(input[0]);
+
+	/* Run kernel. */
+    common_init(&common, NULL);
+    common_create_program(&common, source, NULL, &program);
+    kernel_add = clCreateKernel(program, "add", NULL);
+    kernel_mul = clCreateKernel(program, "mul", NULL);
+    buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(input), input, NULL);
+    clSetKernelArg(kernel_add, 0, sizeof(buffer), &buffer);
+    clSetKernelArg(kernel_mul, 0, sizeof(buffer), &buffer);
+    clEnqueueNDRangeKernel(common.command_queue, kernel_add, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
+    clEnqueueNDRangeKernel(common.command_queue, kernel_mul, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
+    clFlush(common.command_queue);
+    clFinish(common.command_queue);
+    clEnqueueReadBuffer(common.command_queue, buffer, CL_TRUE, 0, sizeof(input), input, 0, NULL, NULL);
+
+	/* Assertions. */
+    assert(input[0] == 4);
+    assert(input[1] == 6);
+
+	/* Cleanup. */
+    clReleaseMemObject(buffer);
+    common_deinit(&common);
+    return EXIT_SUCCESS;
+}