binary_shader

Author: cirosantilli

opencl/README.md

@@ -13,7 +13,8 @@
     1.  [Write buffer](write_buffer.c)
     1.  [Multi-kernel](multi_kernel.c)
     1.  [Benchmark examples](benchmark-examples.md)
-    1.  [Interactive](interactive)
+    1.  [Binary shader](binary_shader)
+    1.  [Interactive](interactive/)
 1.  Tools
     1.  [clinfo](clinfo.md)
     1.  [Benchmarks](benchmarks.md)

opencl/binary_shader.c

@@ -1,42 +1,78 @@
-/*
+/*Let's play with clCreateProgramWithBinary.
+
+Compile the "inc.cl" CL C shader to binary, save the binary to a file, and load the shader from the binary:
+
+    ./prog
+
+Ignore the CL C shader, load binary from a file, and run it:
+
+    ./prog 0
+
+Use another shader instead:
+
+    ./prog 1 dec.cl
+
+The name "binary" is a possible lie: NVIDIA Linux driver 375.39
+just dumps PTX human readable assembly, which we can modify without
+reverse engineering. Hurray!
 */
 
+#define BIN_PATH __FILE__ ".bin.tmp"
+
 #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};
+int main(int argc, char **argv) {
+    Common common;
+    FILE *f;
+    char *binary, *source_path;
+    cl_int input[] = {1, 2}, errcode_ret, binary_status;
     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;
+    int use_cache;
+    long lenght;
     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);
+    if (argc > 1) {
+        use_cache = !strcmp(argv[1], "0");
+    } else {
+        use_cache = 0;
+    }
+    if (argc > 2) {
+        source_path = argv[2];
+    } else {
+        source_path = "inc.cl";
+    }
 
-    /* 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);
+    /* Get the binary, and create a kernel with it. */
+    if (use_cache) {
+        common_init(&common, NULL);
+        binary = common_read_file(BIN_PATH, &lenght);
+        binary_size = lenght;
+    } else {
+        common_init_file(&common, source_path);
+        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);
+        f = fopen(BIN_PATH, "w");
+        fwrite(binary, binary_size, 1, f);
+        fclose(f);
+    }
+    program = clCreateProgramWithBinary(
+        common.context, 1, &common.device, &binary_size,
+        (const unsigned char **)&binary, &binary_status, &errcode_ret
+    );
+    assert(NULL != program);
+    common_assert_success(binary_status);
+    common_assert_success(errcode_ret);
     free(binary);
+    common_build_program(&common, NULL, &program);
+    kernel = clCreateKernel(program, "kmain", &errcode_ret);
+    assert(NULL != kernel);
+    common_assert_success(errcode_ret);
 
-
+    /* Run the kernel created from the 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);

opencl/common.h

@@ -23,7 +23,11 @@ typedef struct {
     cl_program program;
 } Common;
 
-char* common_read_file(const char *path) {
+void common_assert_success(cl_int ret) {
+    assert(ret == CL_SUCCESS);
+}
+
+char* common_read_file(const char *path, long *length_out) {
     char *buffer;
     FILE *f;
     long length;
@@ -33,37 +37,55 @@ char* common_read_file(const char *path) {
     fseek(f, 0, SEEK_END);
     length = ftell(f);
     fseek(f, 0, SEEK_SET);
-    buffer = malloc(length + 1);
+    buffer = malloc(length);
     if (fread(buffer, 1, length, f) < (size_t)length) {
         return NULL;
     }
     fclose(f);
-    buffer[length] = '\0';
+    if (NULL != length_out) {
+        *length_out = length;
+    }
     return buffer;
 }
 
-void common_create_program(
+char* common_read_file_null(const char *path) {
+    char *f;
+    long length;
+    f = common_read_file(path, &length);
+    f = realloc(f, length + 1);
+    f[length] = '\0';
+    return f;
+}
+
+void common_build_program(
     Common *common,
-    const char *source,
     const char *options,
     cl_program *program
 ) {
     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: clCreateProgramWithSource:\n%s\n", err);
+        fprintf(stderr, "error: clBuildProgram:\n%s\n", err);
         free(err);
         exit(EXIT_FAILURE);
     }
 }
 
+void common_create_program(
+    Common *common,
+    const char *source,
+    const char *options,
+    cl_program *program
+) {
+    *program = clCreateProgramWithSource(common->context, 1, &source, NULL, NULL);
+    common_build_program(common, options, program);
+}
+
 void common_create_kernel(
     Common *common,
     const char *source,
@@ -85,7 +107,7 @@ void common_create_kernel_file(
     const char *options
 ) {
     char *source;
-    source = common_read_file(source_path);
+    source = common_read_file_null(source_path);
     common_create_kernel(common, source, options);
     free(source);
 }
@@ -117,7 +139,7 @@ void common_init_file_options(
     const char *options
 ) {
     char *source;
-    source = common_read_file(source_path);
+    source = common_read_file_null(source_path);
     common_init_options(common, source, options);
     free(source);
 }

opencl/inc.cl

@@ -0,0 +1,3 @@
+__kernel void kmain(__global int *out) {
+    out[get_global_id(0)]++;
+}

opencl/inc_vector.c

@@ -15,7 +15,8 @@ int main(void) {
     const char *source =
         "__kernel void kmain(__global int *out) {\n"
         "    out[get_global_id(0)]++;\n"
-        "}\n";
+        "}\n"
+    ;
     cl_int input[] = {1, 2};
     cl_mem buffer;
     Common common;