Opencl refactor

Author: cirosantilli

opencl/README.md

@@ -6,7 +6,7 @@
     1.  [Increment vector](inc_vector.c)
     1.  [Pass by value](pass_by_value.c)
     1.  [Work item built-ins](work_item_builtin.c)
-    1.  [Vector built-in](vector_builtin.c)
+    1.  [Vector type](vector_type.c)
 1.  Tools
     1.  [clinfo](clinfo.md)
     1.  [Benchmarks](benchmarks.md)

opencl/common.h

@@ -5,23 +5,75 @@
 #include <stdio.h>
 #include <stdlib.h>
 
+// http://stackoverflow.com/questions/28500496/opencl-function-found-deprecated-by-visual-studio
+#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
 #include <CL/cl.h>
 
-void common_check_build(
-	cl_int ret,
-    cl_device_id device,
-    cl_program program
+typedef struct {
+    cl_context context;
+    cl_device_id device;
+    cl_kernel kernel;
+    cl_program program;
+} Common;
+
+static void common_init(
+    Common *common,
+    const char *source
 ) {
 	char *err;
 	size_t err_len;
+    cl_int ret;
+    cl_platform_id platform;
+
+    clGetPlatformIDs(1, &platform, NULL);
+    clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &(common->device), NULL);
+    common->context = clCreateContext(NULL, 1, &(common->device), NULL, NULL, NULL);
+    common->program = clCreateProgramWithSource(common->context, 1, &source, NULL, NULL);
+    ret = clBuildProgram(common->program, 1, &(common->device), "", NULL, NULL);
     if (CL_SUCCESS != ret) {
-		clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &err_len);
+		clGetProgramBuildInfo(common->program, common->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &err_len);
 		err = malloc(err_len);
-		clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, err_len, err, NULL);
-		printf("Kernel build error:\n%s\n", err);
+		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->kernel = clCreateKernel(common->program, "main", NULL);
+}
+
+static char * common_read_file(const char *path) {
+    char *buffer;
+    FILE *f;
+    long length;
+
+    f = fopen(path, "r");
+    fseek (f, 0, SEEK_END);
+    length = ftell(f);
+    fseek(f, 0, SEEK_SET);
+    buffer = calloc(1, length + 1);
+    fread (buffer, 1, length, f);
+    fclose (f);
+    buffer[length] = '\0';
+    return buffer;
+}
+
+static void common_init_file(
+    Common *common,
+    const char *source_path
+) {
+    char *source;
+    source = common_read_file(source_path);
+    common_init(common, source);
+    free(source);
+}
+
+static void common_deinit(
+    Common *common
+) {
+    clReleaseProgram(common->program);
+    clReleaseKernel(common->kernel);
+    clReleaseContext(common->context);
+    clReleaseDevice(common->device);
 }
 
 #endif

opencl/getting-started.md

@@ -1,6 +1,8 @@
 # Getting started
 
-Tested in Ubuntu 15.10 NVIDIA 352.
+Tested in Ubuntu 15.10 NVIDIA 352, OpenCL 1.2.
+
+The day we do OpenCL 2.0, it will be put inside a subdirectory and clearly labeled.
 
 ## NVIDIA
 

opencl/inc.c

@@ -1,19 +1,23 @@
 /*
 Pass an int by reference and increment it.
 
-This is our OpenCL hello world. No error checking for simplicity.
+This is our OpenCL hello world, so we are not doing:
+
+- error checking
+- factoring out with anything else
 */
 
 #include <assert.h>
 #include <stdio.h>
 
+#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
 #include <CL/cl.h>
 
 int main(void) {
     const char *source =
         /* kernel pointer arguments must be __global, __constant, or __local. */
         /* https://www.khronos.org/registry/cl/sdk/2.1/docs/man/xhtml/restrictions.html */
-        "__kernel void increment(__global int *out) {\n"
+        "__kernel void main(__global int *out) {\n"
         "    out[0]++;\n"
         "}\n";
     cl_command_queue command_queue;
@@ -31,8 +35,7 @@ int main(void) {
     context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
     program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
 	clBuildProgram(program, 1, &device, "", NULL, NULL);
-    /* The name of the kernel function we want to call. */
-    kernel = clCreateKernel(program, "increment", NULL);
+    kernel = clCreateKernel(program, "main", NULL);
     buffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_int), &input, NULL);
     clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
     command_queue = clCreateCommandQueue(context, device, 0, NULL);
@@ -45,10 +48,11 @@ int main(void) {
     assert(input == 2);
 
     /* Cleanup. */
-    clReleaseKernel(kernel);
-    clReleaseProgram(program);
     clReleaseCommandQueue(command_queue);
-    clReleaseContext(context);
     clReleaseMemObject(buffer);
+    clReleaseProgram(program);
+    clReleaseKernel(kernel);
+    clReleaseContext(context);
+    clReleaseDevice(device);
     return EXIT_SUCCESS;
 }

opencl/inc_vector.c

@@ -17,27 +17,17 @@ int main(void) {
         "      out[get_global_id(0)]++;\n"
         "}\n";
     cl_command_queue command_queue;
-    cl_context context;
-    cl_device_id device;
-    cl_int input[] = {1, 2}, ret;
-    const size_t global_work_size = sizeof(input) / sizeof(cl_int);
-    cl_kernel kernel;
+    cl_int input[] = {1, 2};
     cl_mem buffer;
-    cl_platform_id platform;
-    cl_program program;
+    const size_t global_work_size = sizeof(input) / sizeof(cl_int);
+    Common common;
 
 	/* Run kernel. */
-    clGetPlatformIDs(1, &platform, NULL);
-    clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL);
-    context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
-    buffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(input), &input, NULL);
-    program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
-    ret = clBuildProgram(program, 1, &device, "", NULL, NULL);
-    common_check_build(ret, device, program);
-    kernel = clCreateKernel(program, "main", NULL);
-    clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
-    command_queue = clCreateCommandQueue(context, device, 0, NULL);
-    clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
+    common_init(&common, source);
+    buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(input), &input, NULL);
+    clSetKernelArg(common.kernel, 0, sizeof(cl_mem), &buffer);
+    command_queue = clCreateCommandQueue(common.context, common.device, 0, NULL);
+    clEnqueueNDRangeKernel(command_queue, common.kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
     clFlush(command_queue);
     clFinish(command_queue);
     clEnqueueReadBuffer(command_queue, buffer, CL_TRUE, 0, sizeof(input), &input, 0, NULL, NULL);
@@ -47,10 +37,8 @@ int main(void) {
     assert(input[1] == 3);
 
 	/* Cleanup. */
-    clReleaseKernel(kernel);
-    clReleaseProgram(program);
-    clReleaseCommandQueue(command_queue);
-    clReleaseContext(context);
     clReleaseMemObject(buffer);
+    clReleaseCommandQueue(command_queue);
+    common_deinit(&common);
     return EXIT_SUCCESS;
 }

opencl/pass_by_value.c

@@ -15,27 +15,17 @@ int main(void) {
         "    out[0] = in + 1;\n"
         "}\n";
     cl_command_queue command_queue;
-    cl_context context;
-    cl_device_id device;
-    cl_int input = 1, ret;
-    cl_kernel kernel;
+    cl_int input = 1;
     cl_mem buffer;
-    cl_platform_id platform;
-    cl_program program;
+    Common common;
 
 	/* Run kernel. */
-    clGetPlatformIDs(1, &platform, NULL);
-    clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL);
-    context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
-    buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, NULL);
-    program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
-    ret = clBuildProgram(program, 1, &device, "", NULL, NULL);
-    common_check_build(ret, device, program);
-    kernel = clCreateKernel(program, "main", NULL);
-    clSetKernelArg(kernel, 0, sizeof(cl_int), &input);
-    clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer);
-    command_queue = clCreateCommandQueue(context, device, 0, NULL);
-    clEnqueueTask(command_queue, kernel, 0, NULL, NULL);
+    common_init(&common, source);
+    buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, NULL);
+    clSetKernelArg(common.kernel, 0, sizeof(cl_int), &input);
+    clSetKernelArg(common.kernel, 1, sizeof(cl_mem), &buffer);
+    command_queue = clCreateCommandQueue(common.context, common.device, 0, NULL);
+    clEnqueueTask(command_queue, common.kernel, 0, NULL, NULL);
     clFlush(command_queue);
     clFinish(command_queue);
     clEnqueueReadBuffer(command_queue, buffer, CL_TRUE, 0, sizeof(cl_int), &input, 0, NULL, NULL);
@@ -44,10 +34,8 @@ int main(void) {
     assert(input == 2);
 
 	/* Cleanup. */
-    clReleaseKernel(kernel);
-    clReleaseProgram(program);
-    clReleaseCommandQueue(command_queue);
-    clReleaseContext(context);
     clReleaseMemObject(buffer);
+    clReleaseCommandQueue(command_queue);
+    common_deinit(&common);
     return EXIT_SUCCESS;
 }

opencl/vector_builtin.c

@@ -0,0 +1,45 @@
+/*
+Increment a vector type (int2) using a vector built-in type.
+
+http://stackoverflow.com/questions/13118228/how-to-pass-vector-paramater-to-opencl-kernel-in-c
+
+TODO: is using them faster than scalars?
+
+- http://stackoverflow.com/questions/20200203/using-own-vector-type-in-opencl-seems-to-be-faster
+*/
+
+#include "common.h"
+
+int main(void) {
+    const char *source =
+        "__kernel void main(__global int2 *out) {\n"
+        "      out[get_global_id(0)]++;\n"
+        "}\n";
+    cl_command_queue command_queue;
+    cl_int input[] = {0, 1, 2, 3};
+    cl_mem buffer;
+    Common common;
+    const size_t global_work_size = sizeof(input) / sizeof(cl_int2);
+
+	/* Run kernel. */
+    common_init(&common, source);
+    buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(input), &input, NULL);
+    clSetKernelArg(common.kernel, 0, sizeof(cl_mem), &buffer);
+    command_queue = clCreateCommandQueue(common.context, common.device, 0, NULL);
+    clEnqueueNDRangeKernel(command_queue, common.kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
+    clFlush(command_queue);
+    clFinish(command_queue);
+    clEnqueueReadBuffer(command_queue, buffer, CL_TRUE, 0, sizeof(input), &input, 0, NULL, NULL);
+
+	/* Assertions. */
+    assert(input[0] == 1);
+    assert(input[1] == 2);
+    assert(input[2] == 3);
+    assert(input[3] == 4);
+
+	/* Cleanup. */
+    clReleaseMemObject(buffer);
+    clReleaseCommandQueue(command_queue);
+    common_deinit(&common);
+    return EXIT_SUCCESS;
+}