从 .cpp 文件调用模板化 CUDA 内核

Question

我知道，通常情况下，不能直接从 .cpp 文件调用 CUDA 内核。相反，如果需要这种能力，内核必须封装在 CPU 可调用函数中，该函数的接口进入 .h 文件，其实现与内核一起进入 .cu 文件。

但是，如果内核在其类型中被模板化，并且希望通过 CPU 包装器将模板化性传递给 .cpp 文件（因为模板接口必须在同一个文件中（. h) 作为其实现，因此会导致任何非 nvcc 编译器尝试访问该 .h 文件的问题。

有人知道绕过这个限制的方法吗？也许没有，正如（完全模板化的）CUDA Thrust 库只能从 .cu 文件（see here）直接调用这一事实所证明的那样？

Answer 1

你是对的。内核模板总是必须在.cu 文件中实例化。

对于足够简单的模板函数（例如，只有一个类型参数），重载函数有时可以满足您的需求。或者，您也可以为 .cpp 文件创建另一个模板。

kernel.cu

template <class T>
__global__ void kernel_axpy(T* x, T* y, int len) { ... }

void axpy(float* x, float* y, int len){ kernel_axpy<<<...>>>(x,y,len); }
void axpy(double* x, double* y, int len){ kernel_axpy<<<...>>>(x,y,len); }

axpy.h

extern void axpy(float* x, float* y, int len);
extern void axpy(double* x, double* y, int len);

template <class T> void cpp_axpy(T* x, T* y, int len) { std::cerr<<"Not implemented.\n"<<std::endl; }
template <> void cpp_axpy<float>(float* x, float* y, int len) { axpy(x,y,len); }
template <> void cpp_axpy<double>(double* x, double* y, int len) { axpy(x,y,len); }

main.cpp

#include "axpy.h"

...
{
    axpy(xx,yy,length);
    cpp_axpy<double>(xxx,yyy,lll);
}
...