About a month ago, one of my colleagues has posted a <a href="http://sandbox.arrayfire.com/quest-for-the-smallest-opencl-program/">blog</a> on how to author the most concise OpenCL program using the C++ API provided by khronos. The sample code provided in that post definitely addresses the question at hand, however applying the new tricks that we have learned in real world applications is a completely different problem. 

In today's post, we will be looking at one such situation which is most often encountered by library/middleware programmers. While programming using OpenCL, we often find ourselves in situations where the application needs to compile the kernels once and reuse the built binary later for different data types or different data sets. To reuse the same kernels for different data types, we would have to do some extra steps while building the kernel source itself. We will also take the help of C++ template programming to certain extent though the OpenCL compilers themselves doesn't support templates. Let's dive into the details now, we will modify the code sample from <a href="http://sandbox.arrayfire.com/quest-for-the-smallest-opencl-program/">this</a> blog to serve our purpose.

<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> Compile the kernels only once</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="">here</a>.