CUDA统一内存高级问题——非法内存访问

Question

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

基础是一组间隔对象（我使用的是CUDA样本中已经准备好的）

最重要的是，我创建了一个类来包装这个名为 ivector（间隔向量）的数组 每个 ivector 代表一个盒子（或超级盒子）。

到目前为止一切顺利，使用这个 CUDA 统一内存 并继承他们提供的托管类是可行的！

问题从这里开始。我的实验需要一系列 ivectors，但我找不到让它工作的方法。

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

假设我们的内核是这样的，并且我们的 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_gpus 数组时，这会起作用。假设 interval_gpu 对象也像 NVIDIA 提供的那样按预期运行

我的内核是这样的，我想访问第 0 个 ivector_gpu 元素的第 0 个 interval_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()

调用。我建议使用由例如提供的内存资源。 推力或RMM代替。见

thrust::cuda::universal_memory_resource

和

rmm::mr::managed_memory_resource

。 Thrust 的优点是随CUDA Toolkit 一起发货，但是相比之下Thrust 内存资源的文档比较欠缺。将来，Thrust 中的资源可能会被 libcu++ 中的新资源取代。

Answer 2

Cuda中的继承类有一些限制，看看在这里输入链接描述