Skip to content

Instantly share code, notes, and snippets.

@chick
Created December 30, 2016 22:33
Show Gist options
  • Save chick/ff938a3f2210084b7ca7ee056cc6bff6 to your computer and use it in GitHub Desktop.
Save chick/ff938a3f2210084b7ca7ee056cc6bff6 to your computer and use it in GitHub Desktop.
Non-pencil model OpenCL kernel for 7pt operator
// <file: control.c>
#include <stdio.h>
#include <time.h>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
int control(cl_command_queue queue, cl_kernel kernel_0, cl_mem out, cl_mem mesh) {
int error_code = 0;
size_t global_134217728 [1] = {134217728};
size_t local_512 [1] = {512};
error_code |= clSetKernelArg(kernel_0, 0, 8, & out);
error_code |= clSetKernelArg(kernel_0, 1, 8, & mesh);
error_code |= clEnqueueNDRangeKernel(queue, kernel_0, 1, NULL, global_134217728, local_512, 0, NULL, NULL);
clFinish(queue);
if (error_code != 0) printf("error code %d\n", error_code);
return error_code;
};
)))
INFO:ctree.c.nodes:compilation command: gcc -shared -fPIC -std=c99 -O2 -framework OpenCL -o /var/folders/ls/0xl3wjy949b2b9j36yn36v3c0000gn/T/tmpEfglgb/snowflake/5598646763549957670/_3980259355556366124/_7985492147856592190/OpenCLCompiler/control.so /var/folders/ls/0xl3wjy949b2b9j36yn36v3c0000gn/T/tmpEfglgb/snowflake/5598646763549957670/_3980259355556366124/_7985492147856592190/OpenCLCompiler/control.c -framework OpenCL
INFO:ctree.nodes:Recreating source
INFO:ctree.nodes:file for generated OpenCL: /var/folders/ls/0xl3wjy949b2b9j36yn36v3c0000gn/T/tmpEfglgb/snowflake/5598646763549957670/_3980259355556366124/_7985492147856592190/OpenCLCompiler/kernel_0.cl
INFO:ctree.nodes:generated OpenCL code: (((
// <file: kernel_0.cl>
#define encode514_514_514(x0, x1, x2) (264196 * (long) (x0) + 514 * (long) (x1) + 1 * (long) (x2))
#define encode514_514(x0, x1) (514 * (long) (x0) + 1 * (long) (x1))
__kernel void kernel_0(__global float* out, __global float* mesh) {
size_t global_id = get_global_id(0);
size_t index_0;
size_t index_1;
size_t index_2;
index_0 = (((global_id / 512) / 4096) * 8 + ((global_id % 512) / 64)) * 1 + 1;
index_1 = ((((global_id / 512) % 4096) / 64) * 8 + (((global_id % 512) % 64) / 8)) * 1 + 1;
index_2 = (((((global_id / 512) % 4096) % 64) / 1) * 8 + ((((global_id % 512) % 64) % 8) / 1)) * 1 + 1;
out[encode514_514_514(index_0, index_1, index_2)] = mesh[encode514_514_514(index_0 + -1, index_1 + 0, index_2 + 0)] + mesh[encode514_514_514(index_0 + 0, index_1 + -1, index_2 + 0)] + mesh[encode514_514_514(index_0 + 0, index_1 + 0, index_2 + -1)] + mesh[encode514_514_514(index_0 + 0, index_1 + 0, index_2 + 0)] * -6 + mesh[encode514_514_514(index_0 + 0, index_1 + 0, index_2 + 1)] + mesh[encode514_514_514(index_0 + 0, index_1 + 1, index_2 + 0)] + mesh[encode514_514_514(index_0 + 1, index_1 + 0, index_2 + 0)];
};
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment