CUDA和内核包装器和模板和编译错误

Question

我在内核包装函数中应用模板技术时遇到了问题。

这是我原始思想中的代码：

//----------------------------------------  
// cuda_demo.cuh
template<typename T> 
void kernel_wrapper(T param);

//----------------------------------------   
// cuda_demo.cu
#include <cuda.h>
#include <cuda_runtime.h>
#include "cuda_demo.cuh"

template<typename T>
__global__ void my_kernel(T param) { 
    // do something 
}

template<typename T>
void kernel_wrapper(T param) { 
    my_kernel<<<1,1>>>(param);
}

//---------------------------------------- 
// main.cpp
#include "cuda_demo.cuh"
int main() {
  int param = 10;
  kernel_wrapper(param);
  return 0;
}

很快我发现模板应该在头文件中实现（参见Why can templates only be implemented in the header file?）。

我从中得到两个解决方案，常见的是“在头文件中编写模板声明，然后在实现文件中实现该类（例如.tpp），并在头部末尾包含此实现文件” 。

所以我改变了代码：

//----------------------------------------  
// cuda_demo.cuh
template<typename T> 
void kernel_wrapper(T param);

#include "cuda_demo.cu"

//----------------------------------------   
// cuda_demo.cu
#include <cuda.h>
#include <cuda_runtime.h>

template<typename T>
__global__ void my_kernel(T param) { 
    // do something 
}

template<typename T>
void kernel_wrapper(T param) { 
    my_kernel<<<1,1>>>(param);
}

编译器给我以下错误：

error: expected primary-expression before < token
   my_kernel<<<1,1>>>(param);

当我将所有cuda代码放在“cuda_demo.cuh”中时，会发生同样的错误。

然后我尝试了第二个解决方案如下：

//----------------------------------------  
// cuda_demo.cuh
template<typename T> 
void kernel_wrapper(T param);

//----------------------------------------   
// cuda_demo.cu
#include <cuda.h>
#include <cuda_runtime.h>
#include "cuda_demo.cuh"

template<typename T>
__global__ void my_kernel(T param) { 
    // do something 
}

template<typename T>
void kernel_wrapper(T param) { 
    my_kernel<<<1,1>>>(param);
}

template void kernel_wrapper<int>(int param);

这个效果很好！但在我的项目中，'T'不是一个简单的类型，可能是递归的

Class_1<Class_2<Class_3<...>>>,

这意味着我无法事先弄清楚'T'的具体类型。

有人知道如何解决这个问题吗？

谢谢。

Answer 1

我找到了问题的本质。

所有cuda代码必须包含在.cu文件中，以便它们可以由nvcc编译。谢谢你的提醒。 @talonmies。

最近，我发现一些开源项目将cuda，C ++代码混合成.h或.cuh文件，然后包含来自.cpp文件和.cu文件的头文件。这让我相信cuda代码可以通过gcc编译。

但我终于发现，虽然许多.cpp文件包含cuda代码，但它们都没有在.cpp文件中调用cuda函数。并且cuda函数调用仅存在于.cu文件中。

他们是怎么做到的？答案是条件编译。这样，.cu文件中的cuda代码将由nvcc编译，但.cpp文件中的那些将被gcc忽略。

至于我原来的问题，最有效的解决方案是将模板cuda代码的所有实现写入头文件，并仅在.cu文件中调用内核包装器。

我花了很多时间在这个问题上，我希望我的经验可以帮助你。