Consider the following example of using the preprocessor for constants:
__kernel void exponentor(__global int* data, const uint exponent) { int tid = get_global_id(0); int base = data[tid]; for (int i = 1; i < exponent; ++i) { data[tid] *= base; } }
The number of iterations for the inner for
loop is determined
at run time, after the kernel is issued for execution. However, you can
use the OpenCL™ dynamic compilation feature to ensure the exponent is
known at kernel compile time, which is done during the host run time.
In this case, the kernel appears as follows:
__kernel void exponentor(__global int* data) { int tid = get_global_id(0); int base = data[tid]; for (int i = 1; i < EXPONENT; ++i) { data[tid] *= base; } }
Capitalization indicates that exponent became a preprocessor macro.
The original version of the host code passes exponent_val
through kernel arguments as follows:
clSetKernelArg(kernel, 1, exponent_val);
The updated version uses a compilation step:
sprintf(buildOptions, “-DEXPONENT=%u”, exponent_val); clBuildProgram(program, <...>, buildOptions, <...>);
Thus, the value of exponent is passed during preprocessing of the kernel code. Besides saving stack space used by the kernel, this also enables the compiler to perform optimizations, such as loop unrolling or elimination. This technique is often useful for transferring parameters like image dimensions to video-processing kernels, where the value is only known at host run time, but does not change once it is defined.
NOTE: This approach requires recompiling
the program every time the value of exponent_val
changes.
Prefer the variable-passing approach if you expect to change this value
often.