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

Posted 2023-03-23

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

【英文标题】：Unexpected read access violation error in CUDA when working with unified memory

【发布时间】：2021-07-10 21:07:35

【问题描述】：

我有一个对象说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上不能同时从主机和设备访问统一内存。在启动内核之后，在从主机代码访问统一内存之前，cudaDeviceSynchronize() 调用是强制性的。没有解决方法可以让您在 Windows 上同时从主机和设备访问统一内存。一种可能的解决方法是切换到 linux。这在 CUDA 编程指南的统一内存部分中有介绍。

我希望在 Windows 上有一个解决方法，因为在这种情况下切换到 Linux 不是一种选择。无论如何感谢您的帮助。至少现在我知道这是我必须克服的限制。

从主机和设备访问 Windows 内存资源的另一种可能的解决方法是使用零复制技术，即主机固定内存。例如，使用here，一般矩阵乘法使用普通设备内存，但通信是通过固定内存进行的。但是，从 CUDA 设备代码对固定资源执行大规模操作可能会令人失望，就性能而言。

【参考方案1】：

在 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") 内存，而不是托管内存。但是，对于批量或大规模数据访问，这可能会对性能产生影响。