调用模板CUDA内核时出现问题

Question

我一直在尝试创建模板内核,但是在我的程序中调用它们时遇到了一些麻烦.我有一个Matrix<T>模板类,并在其中定义了一些方法

Matrix.h:

template <typename T> class Matrix {
    ...
    void sum(Matrix<T>& m1, Matrix<T>& m2, Matrix<T>& sum);
    ...
}

#include "Matrix.cu"

Matrix.cu:

#include "MatrixKernel.h"

template<typename T> void Matrix<T>::sum(const Matrix<T>& m, Matrix<T>& sum) {
    ...
    sumKernel<T><<<dimGrid, dimBlock>>>(Matrix<T> m1, Matrix<T> m2, Matrix<T> sum)
    ...
}

MatrixKernel.h:

template<typename T> __global__ void sumKernel(const Matrix<T> m1, const Matrix<T> m2, Matrix<T> sum) {
...
}

问题是,当我从sum中调用sumKernel时,编译器会给我以下错误:

error C2059: syntax error : '<'

有人知道发生了什么吗？在我包含sumKernel调用之前,代码编译得很好.

谢谢.

Answer 1

所以,似乎你确实有一个奇怪的#include,导致代码被错误的编译器编译.使用.cu.h对cuda标头区分gpu标头和cpu标头.确保只有 NVCC编译.cu和.cu.h文件.Cuda文件永远不应该包含在cpp文件中.内核和内核调用应该在一个.cu或多个.cu.h文件中,并且这些文件不应该包含在cpps中的任何位置.

因为您.cu被包含在主机编译器正在编译的头中,所以主机编译器最终命中令牌<<<- 它无法识别.它可能确实理解了令牌,<<因此消耗了令牌,留下了意外<.

这是另一种应该工作的方式(没有尝试过,但它与我们使用的代码类似)

(注意,这可能有用,但它也可能不是解决问题的正确方法.我的老板不喜欢它作为解决方案,并且更愿意为每个变体添加一个实现)

潜在的问题似乎是主机和设备代码之间缺乏区别.我将细节留在我的解决方案中 - 例如将结果复制到设备和从设备复制,总结实现等.

我试图解决的问题是,给定一个结构,你如何模板化它以便在主机和设备上使用？

我将Matrix.h在类型和实现细节上进行模板化.

 template <typename T, typename Implementation<T> > class Matrix {
     void sum(Matrix<T>& m1, Matrix<T>& m2, Matrix<T>& sum)
     {
         Implementation.sumImp(m1, m2, sum);
     }
 }

主机实现,HostMatrixSum.h将在cpu上做的事情:

 #include "Matrix.h"

 template <typename T> struct HostMatrixSum
 {
     void sumImp(Matrix<T>& m1, Matrix<T>& m2, Matrix<T>& sum)
     {
         ...
     }
 }

虽然GpuMatrixSum.cu.h将上传矩阵,但总结并恢复结果:

 #include "Matrix.h"

 template <typename T> struct GpuMatrixSum
 {   
     template<typename T> __global__ void sumKernel(const Matrix<T> m1, const Matrix<T> m2, Matrix<T> sum)
     {
         ...
     }

     void sumImp(Matrix<T>& m1, Matrix<T>& m2, Matrix<T>& sum)
     {
         ...
         sumKernel<T> <<< dimGrid, dimBlock >>> (m1,m2);
         ...
     }
 }

然后,当我们从主机代码中使用Matrix时,我们在主机总和实现上模板化,并且永远不需要查看任何cuda细节:

 #include "Matrix.h"
 #include "HostMatrixSum.h"

 Matrix<int, HostMatrixSum> m1 = Matrix<int>(...);
 Matrix<int, HostMatrixSum> m2 = Matrix<int>(...);
 Matrix<int, HostMatrixSum> result;
 Matrix.sum(m1,m2,result);

如果我们正在处理gpu,我们可以使用sum的加速gpu实现:

 #include "Matrix.h"
 #include "GpuMatrixSum.cu.h"

 Matrix<int, GpuMatrixSum> m1 = Matrix<int>(...);
 Matrix<int, GpuMatrixSum> m2 = Matrix<int>(...);
 Matrix<int, GpuMatrixSum> result;
 Matrix.sum(m1,m2,result);

希望对你有用!