为 OpenCL 提供支模版支持

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