重载operator new分配CUDA统一内存导致数组非法访问问题

问题背景 #

基于CUDA示例库的interval_gpu<double>封装ivector_gpu类，继承Managed类通过重载operator new分配CUDA统一内存，单个ivector_gpu对象可正常工作，但创建ivector_gpu数组并在核函数中访问时，触发非法内存访问错误（err#700）。

相关代码 #

类定义代码 #

#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];
}

核函数代码 #

__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)

运行环境 #

RTX 3060 Ti（GPU Device 0: "Ampere" with compute capability 8.6）

问题原因 #

1. 数组分配未重载operator new[]：Managed类仅重载了单个对象的operator new，创建数组new ivector_gpu[1]时调用默认operator new[]，分配的是主机内存，GPU无法访问。
2. ivector内存分配错误：构造函数中用new interval_gpu<double>(DIMENSIONS)分配单个对象而非数组，且未使用统一内存，GPU访问该指针时触发非法访问。
3. 拷贝构造函数逻辑错误：先赋值ivector = iv.ivector再覆盖指针，且memcpy长度为dims字节，实际应为dims * sizeof(interval_gpu<double>)，拷贝不完整。
4. 核函数参数错误：核函数接收ivector_gpu *&c（主机端指针引用），GPU无法直接处理主机端引用，需改为值传递。

修复方案 #

1. 给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);
  }
};

2. 修复ivector_gpu构造函数，确保ivector用统一内存分配数组 #

inline __host__ ivector_gpu::ivector_gpu(){
  dims = DIMENSIONS;
  // 分配dims个interval_gpu对象的统一内存
  cudaMallocManaged(&ivector, dims * sizeof(interval_gpu<double>));
  // 原地构造每个interval_gpu对象（调用默认构造）
  for(int i=0; i<dims; i++){
    new(&ivector[i]) interval_gpu<double>();
  }
}

inline __host__ ivector_gpu::ivector_gpu(int N){
  dims = N;
  cudaMallocManaged(&ivector, dims * sizeof(interval_gpu<double>));
  for(int i=0; i<dims; i++){
    new(&ivector[i]) interval_gpu<double>();
  }
}

3. 修正拷贝构造函数的逻辑与拷贝长度 #

inline __host__ ivector_gpu::ivector_gpu(const ivector_gpu &iv){
  dims = iv.dims;
  // 分配对应大小的统一内存
  cudaMallocManaged(&ivector, dims * sizeof(interval_gpu<double>));
  // 拷贝完整的interval_gpu数组
  memcpy(ivector, iv.ivector, dims * sizeof(interval_gpu<double>));
}

4. 修改核函数参数为值传递 #

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

5. 修正主函数核函数调用（注意t,s需定义或移除） #

确保t,s是已初始化的interval_gpu<double>对象，否则删除多余参数：

// 示例：若t,s未定义，改为无多余参数的调用
// test<<<1,1>>>(v);               
test<<<1,1>>>(t,s,v);               
CHECKED_CALL(cudaGetLastError());
CHECKED_CALL(cudaDeviceSynchronize()); 

内容的提问来源于stack exchange，提问作者knasiotis