Created
December 30, 2016 22:33
-
-
Save chick/ff938a3f2210084b7ca7ee056cc6bff6 to your computer and use it in GitHub Desktop.
Non-pencil model OpenCL kernel for 7pt operator
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
// <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