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 runtime</li>
</ol>
The whole sample code required to achieve our goals will be split and explained in three steps.

<h4>First Step</h4>
Declare a struct with static method <code>getName</code> and add template specializations for the types you want your kernel to handle. For our example, let's add specializations for <code>int</code>, <code>float</code> and <code>unsigned int</code>. The corresponding code snippet will look as below.
<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 single <code>addVectors</code> function call. If <code>addVectors</code> has the following function signature <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>
There are two things we need to make sure during this stage.
<ul>
<li>Kernel should work out of the box with different data types</li>
<li> </li>
</ul>

The first goal can be achieved 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 time kernel compilation, we shall use the C++11 feature <a href="http://en.cppreference.com/w/cpp/thread/call_once">std::call_once</a>. It is a fancy and secure way to do the task at hand. It is fancy because similar execution pattern can be achieved using a static boolean in a single threaded environment. It is secure because the code encapsulated by <code>std::call_once</code> works fine even in multi-threaded environment as per the documentation.
<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>
The complete code sample is available <a href="https://gist.github.com/9prady9/848e7774c86a03febe4a">here</a>.