如何正确实现Cuda/C++中从主机和设备代码调用其成员的类

Question

我有一个主机类TestClass，它有一个指向类TestTable的指针作为成员，该类的数据存储在GPU上的浮点数组中。TestClass调用访问TestTable内部数据的内核，以及来自TestClass的方法GetValue()。

在阅读了大量内容并尝试了几个选项后，我觉得我的所有选项最终都会归结为相同的内存访问错误。这些选项涉及为哪些方法和类使用哪些类型说明符，以及如何(以及在哪里(初始化TestTable。因此，我对Cuda/C++如何工作的理解可能不足以正确地实现它。应该如何正确设置我的代码？

这是我的main.cu:的最低版本的内容

#include <iostream>
#include <cuda_runtime.h>
#define CUDA_CHECK cuda_check(__FILE__,__LINE__)
inline void cuda_check(std::string file, int line)
{
cudaError_t e = cudaGetLastError();
if (e != cudaSuccess) {
std::cout << std::endl
<< file << ", line " << line << ": "
<< cudaGetErrorString(e) << " (" << e << ")" << std::endl;
exit(1);
}
}
class TestTable {
float* vector_;
int num_cells_;
public:
void Init() {
num_cells_ = 1e4;
cudaMallocManaged(&vector_, num_cells_*sizeof(float));
CUDA_CHECK;
}
void Free() {
cudaFree(vector_);
}
__device__
bool UpdateValue(int global_index, float val) {
int index = global_index % num_cells_;
vector_[index] = val;
return false;
}
};
class TestClass {
private:
float value_;
TestTable* test_table_;
public:
TestClass() : value_(1.) {
// test_table_ = new TestTable;
cudaMallocManaged(&test_table_, sizeof(TestTable));
test_table_->Init();
CUDA_CHECK;
}
~TestClass() {
test_table_->Free();
cudaFree(test_table_);
CUDA_CHECK;
}
__host__ __device__
float GetValue() {
return value_;
}
__host__
void RunKernel();
};
__global__
void test_kernel(TestClass* test_class, TestTable* test_table) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < 1e6; i += stride) {
const float val = test_class->GetValue();
test_table->UpdateValue(i, val);
}
}
__host__
void TestClass::RunKernel() {
test_kernel<<<1,1>>>(this, test_table_);
cudaDeviceSynchronize(); CUDA_CHECK;
}
int main(int argc, char *argv[]) {
TestClass* test_class = new TestClass();
std::cout << "TestClass successfully constructed" << std::endl;
test_class->RunKernel();
std::cout << "Kernel successfully run" << std::endl;
delete test_class;
std::cout << "TestClass successfully destroyed" << std::endl;
return 0;
}

我得到的错误是line 88: an illegal memory access was encountered (700)。

我认为错误在于其中一个问题：

• TestTable没有使用new正确创建，这可能很糟糕。但是，在TestClass()中取消对test_table_ = new TestTable;的注释并不能解决此问题
• test_kernel中的GetValue()未返回有效的浮点变量。如果我用任意的浮点值(例如1.f(替换它，程序运行时不会出错。然而，在我的代码的真实(非最小(版本中，GetValue()会在代码库的不同点进行一系列计算，因此硬编码不是一种选择
• 我从不将TestClass复制到GPU，而是从内核调用它的一个成员函数。我知道这一定会引起麻烦，但我觉得知道在哪里以及如何复制它并不直观。如果我只在内核中调用GetValue()而不重用它的结果，那么就没有错误，所以我的程序似乎可以在不将类复制到GPU的情况下调用GetValue()

我无法应用于我的特定问题的可能相关问题：

• 从cuda内核中访问类数据成员-如何设计正确的主机/设备交互？-这个看起来很相似，但不知怎么的，我没能把它翻译成我的用例
• 在不同的CUDA内核中访问Class Member-在这里，我不确定我有两个类的事实如何"；相互作用"；相互影响会影响解决方案
• CUDA和Classes——这个问题对我来说似乎更通用

非常感谢您的帮助！

Answer 1

这里的问题与如何分配TestClass:有关

TestClass* test_class = new TestClass();

test_class现在是指向主机内存的普通指针。如果您有意在设备代码中使用该指针：

void TestClass::RunKernel() {
test_kernel<<<1,1>>>(this, test_table_);
^^^^

和：

void test_kernel(TestClass* test_class, TestTable* test_table) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < 1e6; i += stride) {
const float val = test_class->GetValue();
^^^^^^^^^^

那行不通。在CUDA中，在设备代码中取消引用主机指针通常是一个基本问题。

我们可以通过使用带有托管分配器的放置new来修复这个问题，用于顶级类：

//TestClass* test_class = new TestClass();
TestClass* test_class;
cudaMallocManaged(&test_class, sizeof(TestClass));
new(test_class) TestClass();

当我们这样做时，也有必要更改解除定位器。正如注释中所示，您还应该确保在取消分配之前调用了析构函数：

// delete test_class;
test_class->~TestClass();
cudaFree(test_class);

当我做出这些更改时，你的代码运行时不会出现运行时错误