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

Question

我有一个对象，比如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: %sn", 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 %dn", d_obj->n); // <-- Read access violation if the above line is commented out    
}

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

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

Answer 1

在windows和CUDA的任何最新版本(比如9.0或更新版本(下，统一内存(或托管内存-同义词(行为表示为：

在Windows上运行的应用程序(无论是TCC模式还是WDDM模式(将使用与6.x之前的体系结构相同的基本统一内存模型，即使它们运行在具有6.x或更高计算能力的硬件上。

稍后，文档指出，对于此类系统，在内核启动后，有必要发出cudaDeviceSynchronize()，然后CPU才能再次访问托管数据。

如果在windows上未能做到这一点，那么在尝试访问任何托管数据时，CPU代码中都会出现seg错误。通过编程，您可以在内核启动后使用文档中介绍的concurrentManagedAccess属性来检查是否需要这种类型的同步。您可以使用cudaDeviceGetAttribute()进行以下操作：

int cmm = 0;
int device_to_check = 0;
cudaDeviceGetAttribute(&cmm, cudaDevAttrConcurrentManagedAccess, device_to_check);
if (cmm) {
//cmm will be true/non-zero if it is safe to not use `cudaDeviceSynchronize()` after a kernel call
}
else {
//cmm is zero, this is the windows case...
}

一些可能的解决方案：

• 切换到Linux(假设您的GPU是cc6.x或更高版本(
• 使用主机固定("零拷贝"(内存，而不是托管内存。然而，对于大容量或大规模的数据访问，这可能会对性能产生影响

请注意，WSL也被认为是用于托管内存使用的windows平台。