Cuda：混合 c++ 和 cuda 代码

Question

我的问题如下： 我想将 cuda 代码添加到已经存在的 c++ 库中，并尽可能地重用我现有的代码。 为了使用多态性，我使用模板类和模板内核。因此，一切都在 .cpp、.h 和 .cuh 文件中实现。不涉及 .cu 文件，因此不使用 nvcc 并且 c++ 编译器会阻塞 >> 内核调用语法。

我已经看过了 [How to separate the kernel file CUDA with the main .cpp file 和 [How to call a CUDA file from a C++ header file? 但我找不到任何可以解决我的问题的设计。

涉及的文件：

main.cpp

实例化一堆我已经存在的类，将它们传递给组成它们的 CudaPrepare 类，并负责准备要传递给只有原始类型的 cuda 代码的数据。

#include "CudaPrepare.h"
#include "CudaSpecificType1.h"
#include "A.h" //already existing classes 
#include "B.h" //already existing classes

void main()
{
    A a(...);
    B b(...);
    CudaSpecificType1 cudaType(...);
    CudaPrepare<CudaSpecificType> cudaPrepare(a, b, cudaType);
    cudaPrepare.run();

}

CudaSpecificType1.cuh

class CudaSpecificType1
{
protected:
/*
a few members
*/
public:
CudaSpecificType1(...) : /*initializations*/ {}
float polymorphicFunction(/*args*/); 
};

CudaPrepare.h

#include "A.h" //already existing classes 
#include "B.h" //already existing classes

template<typename T>
class CudaPrepare
{
protected:
const A& a;
const B& b;
const T& t;
public:
CudaPrepare(const A& a, const B& b, const T& t): A(a), B(b), T(t) {/*some initialization stuff*/}
void run() const
{
/*
data preparation : various discretizations,  sticking to primitive type only, casting to single precision etc...
*/

CudaClass<T> cudaClass(t, /*all the prepared data here*/);
cudaClass.run();

}
};

CudaClass.cuh

template <typename T>
__global__ void kernel(const T t, /*other args*/, float* results)
{
int threadId = ...;
results[threadId] = t.polymorphicFunction(...);

}



template<typename T>
class CudaClass
{
protected:
const T& t;
/*
all the prepared data with primitive types
*/
public:
CudaClass(const T& t, ...) : t(t) /*other initialization*/ {}
void run() const
{
/*
grid size calculation, cuda memory allocation, data transfer to device...
*/
//kernel invocation
kernel<T><<</*grid & block size*/>>>(/*args*/);
/*
clean up with cudaFree(...);
*/
}
};

c++ 编译器在内核调用时按预期给出错误。 CudaClass::run() 无法移动到 .cu 文件，因为该类是模板化的。 我唯一能想到的是引入一个 .cu 文件替换 main.cpp / 或包含一个将从 main.cpp 调用的存根， 但随后 nvcc 无法处理某些 c++11 功能。特别是 A.h 和 B.h 包含很多枚举类...

Answer 1

我尝试了 Cuda 7.0（之前是 6.5）。可悲的是，似乎仍然不支持（至少）以下 c++11 功能：

1. 枚举类

2. 最终关键字

3. 基于范围的 for 循环

然而，正如 Robert Crovella 所建议的，显式模板实例化解决了这个问题。

CudaClass.cuh 必须一分为二：

CudaClass.cuh

template <typename T>
__global__ void kernel(const T t, /*other args*/, float* results)
{
int threadId = ...;
results[threadId] = t.polymorphicFunction(...);

}



template<typename T>
class CudaClass
{
protected:
const T& t;
/*
all the prepared data with primitive types
*/
public:
CudaClass(const T& t, ...) : t(t) /*other initialization*/ {}

void run() const;

};

CudaClass.cu

#include "CudaClass.cuh"



//explicit instantiation, so that the kernel invocation can be in a .cu file
template class CudaClass<CudaSpecificType1>;
/*
other explicit instantiations for various types
*/



template<typename T>
void run() const
{
/*
grid size calculation, cuda memory allocation, data transfer to device...
*/
//kernel invocation
kernel<T><<</*grid & block size*/>>>(/*args*/);
/*
clean up with cudaFree(...);
*/
}