About a month ago, one of my colleagues did a <a href="http://sandbox.arrayfire.com/quest-for-the-smallest-opencl-program/">post</a> on how to author the most concise OpenCL program using the C++ API provided by khronos. In today's post, we shall further modify that example to achieve the following two goals. 
<ol>
<li>Enable the kernel to work with different integral data types out of the box</li>
<li>Ensure that the kernels compile only once at run time per data type</li>
</ol>
To better understand the whole sample code required to achieve our goals, we shall split and explain it in three steps.
<h4>First Step</h4>
We can template the OpenCL kernels by passing a build option <code>-D T="typename"</code> to the kernel compilation step. To pass such options, we would need a construct that can give us a string literal that represents the corresponding integral type. Let us declare a struct with static method <code>getName</code> and add template specializations for the types we want our kernel to handle. For our example, let's add specializations for <code>int</code>, <code>float</code> and <code>unsigned int</code>. The entire code snippet with struct declaration and template specializations should look like the following.
<pre>
template<typename T> struct CLTypes {
	static const char * getName() { return "Unsupported"; }
};

template<> struct CLTypes<int> {
	static const char * getName() { return "int"; }
};

template<> struct CLTypes<float> {
	static const char * getName() { return "float"; }
};

template<> struct CLTypes<unsigned> {
	static const char * getName() { return "unsigned"; }
};
</pre>

<h4>Second Step</h4>
Abstract out the OpenCL related code into a function, <code>addVectors</code>. Given the following function signature for <code>addVectors</code>, <code>template<typename T> void addVectors(vector<T> &out, vector<T> &in1, vector<T> &in2, std::size_t nElems);</code>, the main program body will turn out as below.
<pre>
Context(CL_DEVICE_TYPE_GPU);

static const unsigned elements = 32;
vector<float> data(elements, 5);
vector<float> result(elements);
    
addVectors<float>(result, data, data, elements);

vector<int> data2(elements, 5);
vector<int> result2(elements);

addVectors<int>(result2, data2, data2, elements);
</pre>

<h4>Final and Crucial Step</h4>
We can template the OpenCL kernels by passing a compile option <code>-D T typename</code> to the kernel compilation step. The code written in Step 1 will help us out here. To understand how step 1 is helping step 3, look at the places where the variable <code>options</code> is being used in the below code snippet.

To ensure single runtime kernel compilation, we shall use the C++11 feature <a href="http://en.cppreference.com/w/cpp/thread/call_once">std::call_once</a>.
<pre>
template<typename T>
void addVectors(vector<T> &out, vector<T> &in1, vector<T> &in2, std::size_t nElems)
{
	static std::once_flag	compileFlag;

	static cl::Program		addProg;
	static cl::Kernel		addkernel;

	std::call_once(compileFlag, []() {
		std::string kern = "kernel void add(global const T * const a, global const T * const b, global T * restrict const c) { unsigned idx = get_global_id(0); c[idx] = a[idx] + b[idx]; }";
		
		std::ostringstream options;
		options<<"-D T="<< CLTypes<T>::getName();

		addProg		= cl::Program(kern, false);
		addProg.build(options.str().c_str());
		std::cout<<"vector addition kernel compiled for type: "<<CLTypes<T>::getName()<<std::endl;

		addkernel	= cl::Kernel(addProg, "add");
	});

	Buffer a(begin(in1), end(in1), true, false);
    Buffer b(begin(in2), end(in2), true, false);
    Buffer c(CL_MEM_READ_WRITE, nElems * sizeof(T));

	auto addOp = cl::make_kernel<Buffer, Buffer, Buffer>(addkernel);
		
    addOp(EnqueueArgs(nElems), a, b, c);

	cl::copy(c, begin(out), end(out));
}
</pre>

<h4>Conclusion</h4>
That is all that needs to be done for using the same kernel source with different data types and compiling them only once at run time per data type. The complete code sample is available <a href="https://gist.github.com/9prady9/848e7774c86a03febe4a">here</a>. This sample is merely an example on how you can parametrize kernel source with respect to types, more complicated algorithms would require further modifications to this sample.