Attempt binary shader

Author: cirosantilli

c/stdio_h.c

@@ -75,16 +75,16 @@ Same as `file_size`, but takes the path instead of a `FILE*`.
 */
 long file_size(char *path) {
     FILE *fp;
-    long retur_value;
+    long return_value;
     fp = fopen(path, "r");
     if (fp == NULL) {
         return -1L;
     }
-    retur_value = fget_file_size(fp);
+    return_value = fget_file_size(fp);
     if (fclose(fp) == EOF) {
         return -1L;
     }
-    return retur_value;
+    return return_value;
 }
 
 /*

opencl/binary_shader.c

@@ -0,0 +1,57 @@
+/*
+*/
+
+#include "common.h"
+
+int main(void) {
+    const char *source =
+        "__kernel void kmain(__global int *out) {\n"
+        "    out[get_global_id(0)]++;\n"
+        "}\n"
+    ;
+    cl_int input[] = {1, 2};
+    cl_kernel kernel;
+    cl_mem buffer;
+    cl_program program;
+    Common common;
+    const size_t global_work_size = sizeof(input) / sizeof(input[0]);
+
+	/* Run kernel. */
+    common_init(&common, source);
+
+    unsigned char *binary;
+    size_t binary_size;
+    FILE *f;
+
+    clGetProgramInfo(common.program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binary_size, NULL);
+    binary = malloc(binary_size);
+    clGetProgramInfo(common.program, CL_PROGRAM_BINARIES, binary_size, binary, NULL);
+
+    /* Not mandatory, but fun to reverse engineer their format later on. */
+    f = fopen("a.bin.tmp", "w");
+    fwrite(binary, binary_size, 1, f);
+    fclose(f);
+
+    program = clCreateProgramWithBinary(common.context, 1, &common.device, &binary_size, (const unsigned char **)&binary, NULL, NULL);
+    kernel = clCreateKernel(program, "kmain", NULL);
+    free(binary);
+
+
+    buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(input), input, NULL);
+    clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
+    clEnqueueNDRangeKernel(common.command_queue, kernel, 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] == 2);
+    assert(input[1] == 3);
+
+	/* Cleanup. */
+    clReleaseKernel(kernel);
+    clReleaseProgram(program);
+    clReleaseMemObject(buffer);
+    common_deinit(&common);
+    return EXIT_SUCCESS;
+}

opencl/common.h

@@ -58,7 +58,7 @@ void common_create_program(
         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);
+        fprintf(stderr, "error: clCreateProgramWithSource:\n%s\n", err);
         free(err);
         exit(EXIT_FAILURE);
     }

opencl/inc_vector.c

@@ -1,8 +1,8 @@
 /*
 Increment a vector, one value per work item.
 
-It is useless to do this on a GPU, not enough work per IO,
-it is just a clEnqueueNDRangeKernel + get_global_id hello world.
+It is useless to do this on a GPU, not enough work / IO,
+it's just a clEnqueueNDRangeKernel + get_global_id hello world.
 
 - http://stackoverflow.com/questions/15194798/vector-step-addition-slower-on-cuda
 - http://stackoverflow.com/questions/22005405/how-to-add-up-the-elements-of-an-array-in-gpu-any-function-similar-to-cublasdas
@@ -11,47 +11,31 @@ it is just a clEnqueueNDRangeKernel + get_global_id hello world.
 
 #include "common.h"
 
-int main(int argc, char **argv) {
+int main(void) {
     const char *source =
-        "__kernel void kmain(__global int *io) {\n"
-        "    io[get_global_id(0)]++;\n"
+        "__kernel void kmain(__global int *out) {\n"
+        "    out[get_global_id(0)]++;\n"
         "}\n";
-    cl_int *io, *expected_output;
+    cl_int input[] = {1, 2};
     cl_mem buffer;
     Common common;
-    size_t i, n, io_sizeof;
+    const size_t global_work_size = sizeof(input) / sizeof(input[0]);
 
-    if (argc > 1) {
-        n = strtoul(argv[1], NULL, 10);
-    } else {
-        n = 2;
-    }
-
-    /* Initialize data. */
-    io_sizeof = n * sizeof(*io);
-    io = malloc(io_sizeof);
-    expected_output = malloc(n * sizeof(*expected_output));
-    for (i = 0; i < n; ++i) {
-        io[i] = i;
-        expected_output[i] = i + 1;
-    }
-
-    /* Run kernel. */
+	/* Run kernel. */
     common_init(&common, source);
-    buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, io_sizeof, io, NULL);
+    buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(input), input, NULL);
     clSetKernelArg(common.kernel, 0, sizeof(buffer), &buffer);
-    clEnqueueNDRangeKernel(common.command_queue, common.kernel, 1, NULL, &n, NULL, 0, NULL, NULL);
+    clEnqueueNDRangeKernel(common.command_queue, common.kernel, 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, io_sizeof, io, 0, NULL, NULL);
+    clEnqueueReadBuffer(common.command_queue, buffer, CL_TRUE, 0, sizeof(input), input, 0, NULL, NULL);
 
-    /* Assertions. */
-    common_vec_assert_eq_i(io, expected_output, n);
+	/* Assertions. */
+    assert(input[0] == 2);
+    assert(input[1] == 3);
 
-    /* Cleanup. */
+	/* Cleanup. */
     clReleaseMemObject(buffer);
     common_deinit(&common);
-    free(io);
-    free(expected_output);
     return EXIT_SUCCESS;
 }