通过重载operator new分配CUDA统一内存导致非法访问

Question

我需要一个非常具体的结构，

基数是一个间隔对象数组。我正在使用 CUDA 样本中已经准备好的样本。

最重要的是，我创建了一个类来包装这个名为

ivector

（间隔向量）的数组 每个

ivector

代表一个盒子（或超级盒子）。

到目前为止一切顺利，使用这个 CUDA 统一内存 并继承他们提供的

Managed

类是可行的！

问题从这里开始。我的实验需要一组

ivector

，但我找不到让它工作的方法。

提供编译所需的每个代码有点困难，因为它非常具体。

假设我们的内核是这样的，并且我们的

ivector_gpu

对象有

operator[]

重载，例如：

#define DIMENSIONS 2

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 ivector_gpu: public Managed {
 public:
   __host__ ivector_gpu();
  __device__ __host__ ivector_gpu(const ivector_gpu &iv);
   __host__ ivector_gpu(int N);
  __device__ __host__ interval_gpu<double>& operator[](int i);
  __device__ __host__ ivector_gpu& operator=(ivector_gpu &x);
  __device__ __host__ int size() const;
 private:
  interval_gpu<double> * ivector;
  int dims;
};

inline __host__ ivector_gpu::ivector_gpu(){
  dims = DIMENSIONS;
  ivector = new interval_gpu<double>(DIMENSIONS);
}

inline __host__ ivector_gpu::ivector_gpu(int N){
  dims = N;
  ivector = new interval_gpu<double>(dims);
}

inline __host__ ivector_gpu::ivector_gpu(const ivector_gpu &iv){
  ivector = iv.ivector;
  dims = iv.dims;
  cudaMallocManaged(&ivector, dims);
  memcpy(ivector, iv.ivector, dims);
}

inline __device__ __host__ ivector_gpu& ivector_gpu::operator=(ivector_gpu &x){
  for(int i=0; i<size(); i++){
    ivector[i]=x[i];
  }
  return *this;
}

inline __device__ __host__ interval_gpu<double>& ivector_gpu::operator[](int i) {
  return ivector[i];
}

^ 我提醒您，如果我实例化 1 个对象，但当我想创建一个

ivector_gpu

s 数组时，这会起作用。假设

interval_gpu<T>

对象也像 NVIDIA 提供的那样按预期运行。

我的内核是这样的，我想访问第0个

interval_gpu

元素的第0个

ivector_gpu

元素。

__global__ void test(interval_gpu<double> a, ivector_gpu *&c){
    interval_gpu<double> first = interval_gpu<double>::empty();
    c[0][0] = first;

我的主要是这样的：

//create the array
ivector_gpu * v = new ivector_gpu[1];

//fill it with something
v[0][0] = interval_gpu<double>(0,10);
v[0][1] = interval_gpu<double>(5,10);

//let's print it for test purposes
std::cout << v[0][0].lower() << ' ' << v[0][0].upper()  << std::endl;
std::cout << v[0][1].lower() << ' ' << v[0][1].upper()  << std::endl;

// ^ so far so good, it compiles and works

//let's call the kernel
test<<<1,1>>>(t,s,v);               
CHECKED_CALL(cudaGetLastError());
CHECKED_CALL(cudaDeviceSynchronize()); 

内核抛出

interval.cu(89): ERROR: cudaDeviceSynchronize() returned an illegal memory access was encountered (err#700)

我假设我对指针做错了什么，或者它需要一个新的

cudaMallocManaged

指针，但我对此完全精疲力尽，试图让它工作几个小时。我无法理解 1 个对象有效但一组对象无效的原因。

更多信息，我正在尝试让它在 RTX 3060 TI 上运行

GPU Device 0: "Ampere" with compute capability 8.6

非常感谢任何帮助！

Answer 1

将

operator new[]

和

operator delete[]

的重载添加到

Managed

类应该允许您分配和释放这些对象的数组：

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

  void *operator new[](size_t len) {
    void *ptr;
    cudaMallocManaged(&ptr, len);
    cudaDeviceSynchronize();
    return ptr;
  }

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

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

话虽如此，自 CUDA 6 以来 UM 发生了变化。例如，我认为

cudaDeviceSynchronize()

调用仍然没有意义。人们可以在博文下方的评论中找到他们的理由：

基本上，编程模型假定任何启动的内核都可以访问附加到“全局”流的任何托管内存，即使该内存是在内核启动后分配的。这意味着如果你想立即在 CPU 上分配托管内存和访问权限，你必须确保所有内核都已同步，或者你必须在分配时附加到“主机”流（即做

cudaMallocManaged(&ptr, size, cudaMemAttachHost)

).如果需要从 GPU 访问数据，则后一种选择要求将数据附加到“全局”或特定流。

对于较新版本的 CUDA 与较新的 GPU 架构相结合的情况，其中的大部分内容 not 不再适用。引用CUDA初学者的统一内存（较新的博客文章）：

计算能力低于 6.0 的 CPU 和 GPU 不能同时访问托管内存。这是因为 Pre-Pascal GPU 缺乏硬件页面错误，因此无法保证一致性。在这些 GPU 上，在内核运行时来自 CPU 的访问将导致分段错误。 在 Pascal 和后来的 GPU 上，CPU 和 GPU 可以同时访问托管内存，因为它们都可以处理页面错误；但是，应用程序开发人员需要确保不存在由同时访问引起的竞争条件。

所以，虽然人们可能仍然希望在释放内存之前同步所有内核在某些内存上工作，但同步整个设备（而不是特定的流）有点不灵活。分配后的同步（在

operator new

和

operator new[]

中）现在似乎完全没有必要。

对于现代 C++ 接口，我建议使用由例如提供的内存资源。 推力或RMM代替。参见

thrust::cuda::universal_memory_resource

和

rmm::mr::managed_memory_resource

。

Thrust的优点是随CUDA Toolkit一起发货，但是Thrust内存资源的文档比较欠缺

在未来，Thrust 中的资源可能会被 libcu++ 中的new ones 取代。