c++ - 如何将 cudaMallocManaged 与带有构造函数的对象一起使用

Question

我正在尝试将 CUDA 托管内存与我通过构造函数创建的对象结合使用。

struct A {
    A(float x) : x(x) {}
    float x;
}

__global__ void myKernel(A *a) {
    printf("%f", a->x);
}

int main() {
    A *a;
    cudaMallocManaged(&a, sizeof(A));
    a->x = 42.f;        // This works
    // a = new a(42.f); // This obviously doesn't because a doesn't point
                        // to managed memory now.
    myKernel<<<1,1>>>(a);
}

使用构造函数而不是直接初始化 a->x = ... 对于 A 更复杂的类来说会很方便。

我当然可以使用构造函数创建对象 a 并使用“正常”cudaMalloc 和 cudaMemcpy，如下所示。

A a(42.f);
A *d_a;
cudaMalloc(&d_a, sizeof(A));
cudaMemcpy(d_a, &a, sizeof(A), cudaMemcpyHostToDevice);

是否有可能使用托管内存并仍然使用其构造函数初始化对象？

Answer 1

正如马克·哈里斯 (Mark Harris) 所讨论的那样 here ，您可以为您的类重载 new 和 delete 运算符，以便通过托管内存 API 分配和释放实例，而不是通过主机标准库内存分配。这样做意味着您无法在设备代码中为以此方式定义的任何类调用 new。

为此，一个非常方便的设计模式(再次完全归功于 Mark Harris 博客的想法)，是定义一个仅包含合适的 new 和 delete 的类运算符，然后在定义将使用统一内存的自己的类时继承它。像这样:

class Managed {
public:
  void *operator new(size_t len) {
    void *ptr;
    cudaMallocManaged(&ptr, len);
    cudaDeviceSynchronize();
    return ptr;
  }

  void operator delete(void *ptr) {
    cudaDeviceSynchronize();
    cudaFree(ptr);
  }
};

class A : public Managed
{
public:
    A(float x) : x(x) {}
    float x;
}

// ....

A *a = new A(42.f);

请注意，理想情况下您还应该为数据类定义赋值运算符，以便复制构造和其他有用的 C++ 习惯用法能够正常工作。