CUDA C++：使用调用模板内核的模板函数

Question

我有一个具有模板功能的类。该函数调用模板内核。我正在 Linux 机器上的 Nsight 中进行开发。在此过程中，我遇到了以下两个相互冲突的要求：

1 - 实现模板函数时，定义必须出现在 *.h（或 *.cu.h）文件中，因为在需要模板之前不会生成代码。

2 - 内核代码必须出现在 *.cu 中，因为编译器无法识别头文件中的 >> 标记。

我认为可能有一种方法可以通过一点编译器巫术来绕过第二个。

当我设置模板成员函数在 *.cu.h 文件中的系统时，我得到以下编译器错误：

错误：“

错误：'>' 标记之前的预期主表达式

这似乎表明它正在解析 > 标记，而不是识别 >> 标记。

代码相关部分的结构概要如下：

在 MyClass.cu.h 中：

#include "MyKernels.cu.h"

class MyClass{
    template <typename T> void myFunction(T* param1, int param2);
};

template <typename T> void myFunction(T* param1, int param2){
    blocks = 16;
    blockSize = 512;
    myKernel<<<blocks, bockSize>>>(d_param1, param2);
}

在 MyKernels.cu.h 中：

#ifndef MYKERNELS_H_
#define MYKERNELS_H_

template <typename T>
extern __global__ void myKernel(T* param1, int param2);
#endif

在 MyKernels.cu 中：

#include "MyKernels.cu.h"

template<typename T>
__global__ void myKernel(T* param1, int param2){
    //Do stuff
}

2015 年 7 月 31 日编辑： 为了使我要完成的工作的结构更加清晰，我编写了一个小型演示项目。它在 github 上公开发布，网址如下：

https://github.com/nvparrish/CudaTemplateProblem

Answer 1

包装函数声明需要在头文件中。函数定义没有。

这是我的想法：

$ cat MyClass.cuh
template <typename T> void kernel_wrapper(T*, int);
class MyClass{
  public:
    template <typename T> void myFunction(T* param1, int param2);
};

template <typename T> void MyClass::myFunction(T* param1, int param2){
    kernel_wrapper(param1, param2);
}
$ cat MyKernels.cu
#include "MyClass.cuh"
#define nTPB 256

template <typename T>
__global__ void myKernel(T* param1, int param2){

  int i = threadIdx.x+blockDim.x*blockIdx.x;
  if (i < param2){
    param1[i] += (T)param2;
  }
}

template <typename T>
void kernel_wrapper(T* param1, int param2){
  myKernel<<<(param2+nTPB-1)/nTPB,nTPB>>>(param1, param2);
  cudaDeviceSynchronize();
}

template void MyClass::myFunction(float *, int);
template void MyClass::myFunction(int *, int);

$ cat mymain.cpp
#include "MyClass.cuh"

int main(){

  MyClass A;
  float *fdata;
  int *idata, size;
  A.myFunction(fdata, size);
  A.myFunction(idata, size);
}

$ nvcc -c MyKernels.cu
$ g++ -o test mymain.cpp MyKernels.o -L/usr/local/cuda/lib64 -lcudart
$

注意强制模板实例化。如果您希望在一个编译单元（内核定义所属的 .cu 文件）中发生模板特化，这将是必要的，因此它可以在另一个编译单元（.cpp 文件，它不理解 cuda 语法）中使用。