时间戳

首页 > 解决方案 > 如何正确实现其成员同时从 Cuda/C++ 中的主机和设备代码调用的类？

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

问题描述

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

在阅读了很多内容并尝试了几个关于哪些类型说明符用于哪些方法和类以及如何（和在哪里）初始化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).

我认为错误在于以下问题之一：

我无法应用于我的具体问题的可能相关问题：

非常感谢任何帮助！

标签： c++oopcuda

解决方案

这里的问题与您如何分配有关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();

当我们这样做时，还需要更改释放器。正如评论中所指出的，您还应该确保在 de-allocation 之前调用析构函数：

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

当我进行这些更改时，您的代码运行时不会出现运行时错误。

推荐阅读