为OpenCL提供支持template

OpenCL本身（至少目前）并不支持模版。在许多场景下（例如移植CUDA代码），这会带来一些麻烦。

这里，我们介绍一种解决方案。主要思路是，利用OpenCL在运行时编译的特点，通过宏定义实现类似C++模型的功能。

首先我们要解决动态获取类型名称的问题

解决方法如下：

template<typename T>
struct TypeParseTraits {
static const char *name;
};
template<typename T>
const char * TypeParseTraits<T>::name = "Unsupported Type!";

#define REGISTER_PARSE_TYPE(T) template<> struct TypeParseTraits<T> {\
static const char* name; \
};  \
template<> const char * TypeParseTraits<T>::name = #T;

REGISTER_PARSE_TYPE(int);
REGISTER_PARSE_TYPE(float);
REGISTER_PARSE_TYPE(double);

对于注册过的类似，我们可以通过 TypeParseTraits::name得到相应的字符串。对于未注册的类型，默认返回“Unsupported”。（一个实例参见 boostCompute）

另外一种类似的方法如下：

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<double> {
static const char * getName() { return "double"; }
};

需要注意的是，无论哪各方法，我们可能会选择使用C++提供的内置函数功能。

typeid(T).name()

但 C++ 标准并不保证能返回期望的结果（例如int并不保证能返回我们期望的字符串”int”）（参见这里），因此，使用它这会存在潜在的移植性问题。

动态获得类型名称的基础上，我们可以通过宏替换 OpenCL kernel 定义中的”模型“类型。

static std::once_flag compiled;
std::call_once(compiled, []() {
std::ostringstream options;
options << "-D T=" << TypeParseTraits<T>::name;
prg = cl::Program(addVectors_src, false);
prg.build(options.str().c_str());
std::cout << "vector addition kernel compiled for type: " << TypeParseTraits<T>::name << std::endl;
kl = cl::Kernel(prg, "addVectors_kl");
});

为了保证每个kernel源码只会编译一次，这里我们使用了C++11的call_once功能（参见这里）

实现这个目的的另一种方法是使用利用常规的静态变量。

static bool compliled = []() {
std::ostringstream options;
options << "-D T=" << TypeParseTraits<T>::name;
prg = cl::Program(addVectors_src, false);
prg.build(options.str().c_str());
std::cout << "vector addition kernel compiled for type: " << TypeParseTraits<T>::name << std::endl;
kl = cl::Kernel(prg, "addVectors_kl");
return true;
};

这里，两种方法并无差别。但在其他情形下，call_once可能会更灵活一些，参见讨论。

完整的代码如下：

#define STRINGFY(src) #src

template<typename T>
struct TypeParseTraits {
static const char* name = "Unsupported";
};

#define REGISTER_PARSE_TYPE(X) template<> struct TypeParseTraits<X> { \
static const char* name = #X; \
};

REGISTER_PARSE_TYPE(int);
REGISTER_PARSE_TYPE(float);
REGISTER_PARSE_TYPE(double);

template<typename T>
void addVectors(vector<T> &out, vector<T> &in1, vector<T> &in2, size_t n) {
static std::once_flag compiled;
static cl::Program prg;
static cl::Kernel kl;

static const char * addVectors_src = STRINGFY(
kernel
void addVectors_kl(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::call_once(compiled, []() {
std::ostringstream options;
options << "-D T=" << TypeParseTraits<T>::name;
prg = cl::Program(addVectors_src, false);
prg.build(options.str().c_str());
std::cout << "vector addition kernel compiled for type: " << TypeParseTraits<T>::name << std::endl;
kl = cl::Kernel(prg, "addVectors_kl");
});

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

auto addVectors_kl = cl::make_kernel<Buffer, Buffer, Buffer>(kl);

addVectors_kl(EnqueueArgs(n), a, b, c);

cl::copy(c, begin(out), end(out));
}

一个简单的测试

void pseudotemplate_test() {
const int n = 10;

vector<int> iv1(n, 1);
vector<int> iv2(n, 10);
vector<int> iv3(n);

vector<double> dv1(n, 2.0);
vector<double> dv2(n, 3.0);
vector<double> dv3(n);

addVectors(iv3, iv1, iv2, iv1.size());
addVectors(dv3, dv1, dv2, dv1.size());

for (int i = 0; i < n; i++) {
cout << iv3[i] << ":\t" << iv1[i] + iv2[i] << endl;
}
cout << endl;

for (int i = 0; i < n; i++) {
cout << dv3[i] << ":\t" << dv1[i] + dv2[i] << endl;
}
cout << endl;
}

输出

11: 11
11: 11
11: 11
11: 11
11: 11
11: 11
11: 11
11: 11
11: 11
11: 11

5:  5
5:  5
5:  5
5:  5
5:  5
5:  5
5:  5
5:  5
5:  5
5:  5

关于STRINGFY，参见

1. stackoverflow上关于动态获取类型名称字符串的讨论

2. Templating and Caching OpenCL Kernels