c++ - 为什么 CUDA 在访问类成员时崩溃？

Question

抱歉，如果这是一个新问题，但我找不到它。我一直在编写一个可以使用 CDUA 加速的程序。但是，我的方法在很大程度上依赖于将类对象(完全，而不仅仅是它们的成员)传递给内核来运行。为确定这是否可行，我编写了一个小测试程序。

class bulkArray {
public:
    double* value;
    int xSize;

    bulkArray(int xSize) {
        value = new double[xSize];
        this->xSize = xSize;
    }
};

__global__ void addArrays(bulkArray *a, bulkArray *b, bulkArray *c, int N) {
    int id = blockIdx.x*blockDim.x + threadIdx.x;

    if (id < N)
        c->value[id] = a->value[id] + b->value[id];
}

int main() {
    int N = 50000000;

    bulkArray *a;
    bulkArray *b;
    bulkArray *c;

    a = new bulkArray(N);
    b = new bulkArray(N);
    c = new bulkArray(N);

    // allocate unified memory.
    cudaMallocManaged(&a, sizeof(a));
    cudaMallocManaged(&b, sizeof(b));
    cudaMallocManaged(&c, sizeof(c));

    // init vectors on host.
    for (int i = 0; i < N; i++) {
        // CRASHING HERE.
        a->value[i] = sin(i) * cos(i);
        b->value[i] = sin(i) * cos(i);
    }

    int blockSize = 1024;
    int gridSize = (int)ceil((float)N / blockSize);

    addArrays << <gridSize, blockSize >> > (a, b, c, N);

    // sum up vector c.
    double sum = 0;
    for (int i = 0; i < N; i++) {
        sum += c->value[i];
    }

    cout << "Final result: " << sum << endl;

    cudaFree(a);
    cudaFree(b);
    cudaFree(c);

    return 0;
}

不知何故，将 3 个对象加载到统一内存中是可行的，但是当我尝试访问它们进行修改时失败了。我查看了其他指南和示例，但无法完全确定出了什么问题。对于发布的代码量，我深表歉意，但我忍不住担心我正在尝试的事情是不可能的。

Answer 1

您可能对托管内存或 cudaMallocManaged 的工作方式有一些误解。在某些方面，cudaMallocManaged 在概念上类似于 C 库 malloc，只是它分配托管内存。有关在简单 CUDA 应用程序中使用托管内存的更详细介绍，您可能需要引用 this blog .

具体来说，您的代码存在一些问题:

1. 当对具有嵌入式指针的对象数组使用托管分配时，如果您打算在设备上使用它，则所有级别的分配都需要替换为托管分配。因此，如果您想在设备代码中访问该成员/字段，则在构造函数中使用 new 将不起作用。我们可以在那里替换 cudaMallocManaged。

2. 我们不会使用 new 分配一个指针，然后使用 cudaMallocManaged 重新分配相同的指针。

3. C(或 C++)中的 sizeof 函数在指针上使用时将返回该指针 的大小，而不是它指向的任何内容的大小.因此，这不是为对象数组分配资源的明智方法。

4. CUDA 内核启动是异步的，因此在内核启动后，如果您想使用托管数据，则需要创建某种同步。

以下是为解决上述问题而进行的一组最小修改。我省略了适当的 CUDA 错误检查，但我强烈建议在开发 CUDA 代码时这样做。如果您遇到困难，我还建议使用 cuda-memcheck 运行此代码:

$ cat t51.cu
#include <iostream>
using namespace std;
class bulkArray {
public:
    double* value;
    int xSize;

    void init(int xSize) {
        cudaMallocManaged(&value, xSize*sizeof(double));
        this->xSize = xSize;
    }
};

__global__ void addArrays(bulkArray *a, bulkArray *b, bulkArray *c, int N) {
    int id = blockIdx.x*blockDim.x + threadIdx.x;

    if (id < N)
        c->value[id] = a->value[id] + b->value[id];
}

int main() {
    int N = 50000;

    bulkArray *a;
    bulkArray *b;
    bulkArray *c;

    // allocate unified memory.
    cudaMallocManaged(&a, sizeof(bulkArray));
    cudaMallocManaged(&b, sizeof(bulkArray));
    cudaMallocManaged(&c, sizeof(bulkArray));
    a->init(N);
    b->init(N);
    c->init(N);
    // init vectors on host.
    for (int i = 0; i < N; i++) {
        a->value[i] = sin(i) * cos(i);
        b->value[i] = sin(i) * cos(i);
    }

    int blockSize = 1024;
    int gridSize = (int)ceil((float)N / blockSize);

    addArrays << <gridSize, blockSize >> > (a, b, c, N);
    cudaDeviceSynchronize();
    // sum up vector c.
    double sum = 0;
    double sum2 = 0;
    for (int i = 0; i < N; i++) {
        sum += c->value[i];
        sum2 += a->value[i] + b->value[i];
    }

    cout << "Final result: " << sum << " should be: " << sum2 <<  endl;

    cudaFree(a);
    cudaFree(b);
    cudaFree(c);

    return 0;
}
$ nvcc -arch=sm_35 -o t51 t51.cu
$ cuda-memcheck ./t51
========= CUDA-MEMCHECK
Final result: 0.624013 should be: 0.624013
========= ERROR SUMMARY: 0 errors
$