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?）。

我从中得到了两种解决方案，常见的一种是 "to write the template declaration in a header file, then implement the class in an implementation file (for example .tpp), and include this implementation file at the end of the header"。

所以我更改代码：

//----------------------------------------  
// 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代码，但是其中none调用了.cpp文件中的cuda函数。并且cuda函数调用只存在于.cu文件中

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

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

我在这个问题上花了很多时间，希望我的经验能对你有所帮助。

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

CUDA & Kernel Wrapper & Template & Compile Error

templates

cuda

compiler-errors

wrapper