SOCL: conditional compilation in OpenCL kernels depending on device type

SOCL now defines one of the following constants depending on the device the kernel is compiled on: SOCL_DEVICE_TYPE_CPU, SOCL_DEVICE_TYPE_GPU, SOCL_DEVICE_TYPE_ACCELERATOR, SOCL_DEVICE_TYPE_UNKNOWN. OpenCL kernels can use this value to perform conditional compilation.

socl/examples/basic/basic.c

@@ -37,11 +37,22 @@
 #define REALSIZE (SIZE * sizeof(TYPE))
 
 const char * kernel_src = "__kernel void add(__global float*s1, __global float*s2, __global float*d) { \
-   size_t x = get_global_id(0);\
-   size_t y = get_global_id(1);\
-   size_t w = get_global_size(0); \
-   int idx = y*w+x; \
-   d[idx] = s1[idx] + s2[idx];\
+   size_t x = get_global_id(0);\n\
+   size_t y = get_global_id(1);\n\
+   size_t w = get_global_size(0); \n\
+   int idx = y*w+x; \n\
+#ifdef SOCL_DEVICE_TYPE_GPU \n\
+   d[idx] = s1[idx] + s2[idx];\n\
+#endif \n\
+#ifdef SOCL_DEVICE_TYPE_CPU \n\
+   d[idx] = s1[idx] + 2* s2[idx];\n\
+#endif \n\
+#ifdef SOCL_DEVICE_TYPE_ACCELERATOR \n\
+   d[idx] = s1[idx] + 3 * s2[idx];\n\
+#endif \n\
+#ifdef SOCL_DEVICE_TYPE_UNKNOWN \n\
+   d[idx] = s1[idx] + 4 * s2[idx];\n\
+#endif \n\
 }";
 
 

socl/src/cl_buildprogram.c

@@ -32,7 +32,17 @@ static void soclBuildProgram_task(void *data) {
 
    DEBUG_MSG("[Worker %d] Building program...\n", wid);
 
-   err = clBuildProgram(d->program->cl_programs[range], 1, &device, d->options, NULL, NULL);
+   cl_device_type dev_type;
+   clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(cl_device_type), &dev_type, NULL);
+   char * dev_type_str = (dev_type == CL_DEVICE_TYPE_CPU ? "CPU" :
+                          dev_type == CL_DEVICE_TYPE_GPU ? "GPU" :
+                          dev_type == CL_DEVICE_TYPE_ACCELERATOR ? "ACCELERATOR" : "UNKNOWN");
+
+   char opts[4096];
+   sprintf(opts, "-DSOCL_DEVICE_TYPE_%s %s",
+    dev_type_str, (d->options != NULL ? d->options : ""));
+
+   err = clBuildProgram(d->program->cl_programs[range], 1, &device, opts, NULL, NULL);
    if (err != CL_SUCCESS) {
       size_t len;
       clGetProgramBuildInfo(d->program->cl_programs[range], device, CL_PROGRAM_BUILD_LOG, 0, NULL, &len);