使用统一内存时 CUDA 中出现意外的读取访问冲突错误

如何解决使用统一内存时 CUDA 中出现意外的读取访问冲突错误

我有一个对象说 d_obj，它在统一内存上有一些成员，在设备内存上有一些明确的成员。然后我调用一个 CUDA 内核来获取对象并使用它。我想立即让 cpu 在内核调用后立即对统一内存上的成员做一些事情，但失败了。在这里，我使用短代码重现了我的问题：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"    
#include <stdio.h>

#define CHECK_CUDA(call)                                            \
{                                                                   \
const cudaError_t error = call;                                     \
if (error != cudaSuccess)                                           \
{                                                                   \
printf("ERROR:: File: %s,Line: %d,",__FILE__,__LINE__);         \
printf("code: %d,reason: %s\n",error,cudaGetErrorString(error)); \
exit(EXIT_FAILURE);                                                 \
}                                                                   \
}

class MyClass
{
public:
    MyClass(int n_) : n(n_) { }
    void allocateMeOnDevice() {
        CHECK_CUDA(cudamalloc((void**)&vec,n * sizeof(float)));
    }
    int n;
    float* vec;
};

__global__ void kernel(MyClass* obj) {
    for (int i = 0; i < obj->n; i++) {
        obj->vec[i] = 1;
    }
}

int main() {
    
    int n = 1000;

    MyClass h_obj(n); 

    MyClass* d_obj;
    CHECK_CUDA(cudamallocManaged((void**)&d_obj,sizeof(MyClass)));
    CHECK_CUDA(cudamemcpy(d_obj,&h_obj,sizeof(MyClass),cudamemcpyHostToDevice));
    d_obj->allocateMeOnDevice();

    kernel << <1,1 >> > (d_obj);

    //CHECK_CUDA(cudaDeviceSynchronize()); 
    printf("** d_obj->n is %d\n",d_obj->n); // <-- Read access violation if the above line is commented out    

}

是否不能同时从主机和设备访问统一内存上的某些内容？我想知道这个问题是否有任何解决方法？

操作系统：Windows 10/CUDA 11.2/设备：GeForce RTX 3090

解决方法

在 Windows 和任何最新版本的 CUDA（例如 9.0 或更高版本）下，unified memory（或托管内存 - 同义词）行为表示为：

在 Windows 上运行的应用程序（无论是在 TCC 还是 WDDM 模式下）将使用与 6.x 之前的架构一样的基本统一内存模型，即使它们运行在具有 6.x 或更高计算能力的硬件上也是如此。

后来，the documentation 表示对于此类系统，必须在内核启动后发出 cudaDeviceSynchronize()，然后 CPU 才能再次访问托管数据。

如果您未能在 Windows 上执行此操作，您将在尝试访问任何托管数据的 CPU 代码中遇到段错误。

一些可能的解决方法：

切换到 Linux（假设您的 GPU 是 cc6.x 或更高版本）
使用主机固定 ("zero-copy") 内存，而不是托管内存。但是，对于批量或大规模数据访问，这可能会对性能产生影响。