CUDA代码的C++设计

Question

我有一段C++ CUDA代码,我必须在float中声明数据变量.我还必须重写声明数据变量为double的代码.

在CUDA中处理这种情况的好设计是什么？

我不希望有两组相同的代码,因为在将来任何更改我将不得不更改两组相同的代码.我还希望保持代码清洁,而不需要在代码中的#ifdeffloat和double之间进行更改.

任何人都可以建议任何好的(在维护和"易于阅读")设计？

Answer 1

CUDA支持类型模板,毫无疑问,它是实现内核代码的最有效方法,您需要在同一代码中处理多个类型.

作为一个简单的例子,考虑一个简单的BLAS AXPY类型内核:

template<typename Real>
__global__ void axpy(const Real *x, Real *y, const int n, const Real a)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for(; tid<n; tid += stride) {
        Real yval = y[tid];
        yval += a * x[tid];
        y[tid] = yval;
    }
}

这个模板化的内核可以实现双精度和单精度,而不会失去一般性:

template axpy<float>(const float *, float *, const int, const float);
template axpy<double>(const double *, double *, const int, const double);

推力模板库随CUDA工具包的所有最新版本一起提供,广泛使用此工具来实现类型不可知算法.