logo标签列表

调用CUDA模板核函数时出现问题

c++templatescuda
7

我一直在尝试创建模板内核，但在程序中调用时遇到了一些问题。我有一个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调用之前，代码编译得很好。

谢谢。

我不知道你可以使用CUDA和C++（！）。微不足道的建议：尝试在<T>和<<<之间放置一个空格，以防将它们一起运行导致解析问题。 - Rup
编译器告诉你错误在哪一行了吗？CUDA模板代码中有很多<s，因此缩小错误所在的行数范围会很有用。 - Matt Bond
@Rup：可以的，甚至可以将对象作为参数传递给核函数（前提是你已经将感兴趣的数据复制到设备内存）。我也会尝试你的建议。 @Bomadeno：错误出现在执行核函数调用的那一行。 - Renan
谢谢 - 听起来我只是过时了。维基百科说C++支持是在CUDA 2.x中新增的。 - Rup
2个回答
5
因此，看起来您确实有一个奇怪的#include，导致代码被错误的编译器编译。通过使用.cu.h作为cuda头文件的标识，区分gpu头文件和cpu头文件。确保只有NVCC编译.cu和.cu.h文件。Cuda文件不应包含在cpp文件中。内核和内核调用应在.cu或.cu.h文件中，并且这些文件不应在cpps中任何地方包含。
由于您的.cu正在被包含在由主机编译器编译的标头中，因此主机编译器最终会命中标记<<< - 它无法识别。它可能确实理解标记<<，因此它消耗了那个标记，留下一个意外的<。
以下是另一种应该有效的方法（我没有尝试过，但它类似于我们使用的代码）
（注意，这可能有效，但也可能不是解决问题的正确方法。我的老板不喜欢它作为解决方案，而更喜欢添加每个变化的实现）
根本问题似乎是主机代码和设备代码之间缺乏区别。在我的解决方案中，我将在类型和实现细节上对Matrix.h进行模板化。
我将在类型和实现细节上对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);
         ...
     }
 }

当我们从主机代码中使用矩阵时，我们在主机上对模板进行求和实现，不需要看到任何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，我们可以使用加速的GPU实现sum：

 #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);

希望这对您有用！
我打算尝试一下。但同时，MatrixKernel.h中的__global__关键字没有引起编译器的投诉似乎有些奇怪（这只能意味着是NVCC在编译，对吗？）另外一件事：如果你说的确实是问题，那么我应该在哪里实现sum方法？如果我不在"Matrix.h"中写'#include "Matrix.cu"'，就会出现链接错误，因为模板必须在同一个文件中声明和定义... - Renan
我认为你是对的，我必须放弃将Matrix实现为模板类的想法，因为没有办法让它以那种方式工作。如果我在头文件中包含一个.cu文件，那么每个包含Matrix头文件的其他文件也会包含这个.cu文件，甚至是.cpp文件，这将不可避免地导致编译错误。使用模板内核是可以的，但是由于我刚才解释的原因，使调用它们的C++方法也成为模板是不可能的。毕竟有点复杂... - Renan
因为 MatrixKernel.h 只被 cu 文件包含，所以只有 nvcc 包括它。如果您在主机 cpp 文件中包含了 MatrixKernel.h，我怀疑它会崩溃。我将 cuda 特定头文件命名为 .cu.h，以澄清它仅供 .cu 文件包含使用。我正在努力回答“如何做”部分的问题——尝试找到一个优雅的解决方案 :) - Matt Bond
1
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接