CUDA 和模板：需要专业化声明吗？

Question

我有一个模板化的包装函数，它调用__global__像这样的 .cu 文件中定义的内核 ( )

template<typename T, class M> 
__global__ void compute_kernel(T* input, T* output, n) {
    M m;
    // compute stuff using m
};

template<typename T, class M> 
void compute(T* input, T* output, int n) {
    // ... compute blocks, threads, etc.
    compute_kernel<T,M> <<<dim_grid, dim_block>>>(input, output, n);
    // ...
};

和一个包含在主机代码中的头文件，它只有声明

template<typename T, class M> 
void compute(T* input, T* output, int n);

但是，compute()从带有任意模板参数的主机调用，编译失败，undefined reference to 'void reduce(...)'并且仅当我在.cu文件末尾添加专门化声明时，代码才会编译：

template void
compute<int, Method1<int> >(int* input, int* output, int n);

template void
compute<float, Method1<float> >(float* input, float* output, int n);

template void
compute<int, Method2<int> >(int* input, int* output, int n);

template void
compute<float, Method2<float> >(float* input, float* output, int n);

那么，是否有必要专门化每个模板化函数以使其可从主机调用？（这是一个相当大的缺点）

感谢您的意见！

Answer 1

这是一个C++ FAQ，不限于 CUDA。

如果您在 .cpp 或 .cu 文件中有模板实现，那么当您编译该翻译单元时，编译器不可能知道您需要什么模板参数排列。因此，当您链接时，您会收到错误。

您可以将实现放在头文件中（在这种情况下，您需要在 .cu 文件中实例化，因为它包含 CUDA），或者您必须显式实例化所有所需的排列。如果您必须执行其中许多操作，那么您可以使用宏来实例化所有排列。