自定义CUDA后端Vec类析构时设备内存释放失败问题求助

问题根源分析 #

你的代码存在两个核心错误，直接导致了析构阶段的非法内存访问报错：

1. 共享内存越界访问
   你定义的BLOCK_SIZE为512，但调用vec_dot内核时指定的block大小是1024：

   vec_dot<<<100,1024>>>(A.data_ptr, B.data_ptr, A.size, out_device);
   

   而内核中声明的共享内存数组大小仅为BLOCK_SIZE（512）：

   __shared__ T sdata[BLOCK_SIZE];
   

   当线程索引threadIdx.x超过511时，访问sdata[threadIdx.x]会触发共享内存越界，引发非法内存访问错误。该错误会污染CUDA上下文，导致后续cudaFree调用返回异常。

2. 归约阶段错误访问全局内存
   内核归约循环中，错误地直接访问全局内存的vecA和vecB，而非使用已加载到共享内存的sdata：

   if (thread_index < s) sdata[thread_index] += vecA[thread_index + s] * vecB[thread_index + s];
   

   归约操作应基于共享内存中的中间结果完成，而非重新计算全局内存的乘积，这不仅逻辑错误，还会引发不必要的全局内存访问开销。

修复方案 #

1. 统一Block Size定义 #

将内核调用的block大小与共享内存数组大小统一，同时计算合理的grid size以避免冗余block：

// 替换原内核调用代码
dim3 grid((A.size + BLOCK_SIZE - 1)/BLOCK_SIZE);
dim3 block(BLOCK_SIZE);
vec_dot<T><<<grid, block>>>(A.data_ptr, B.data_ptr, A.size, out_device);

2. 修复归约阶段的共享内存访问 #

修改内核归约逻辑，正确使用共享内存完成累加：

for (unsigned int s=blockDim.x/2; s>0; s>>=1) {
    __syncthreads();
    if (thread_index < s) sdata[thread_index] += sdata[thread_index + s];
}

注意：每次归约前必须调用__syncthreads()，确保所有线程完成上一轮的共享内存写入操作。

3. 修正内核参数类型 #

将内核中的size参数从int改为size_t，避免大向量场景下的类型溢出：

__global__ void vec_dot(const T* vecA, const T* vecB, size_t size, T* out)

4. 完善类的拷贝控制（可选但关键） #

Vec类未定义拷贝构造函数和赋值运算符，默认浅拷贝会导致多个对象持有同一设备指针，引发重复释放问题。添加以下代码禁用拷贝或实现移动语义：

// 禁用拷贝构造与赋值（无需拷贝场景）
Vec(const Vec&) = delete;
Vec& operator=(const Vec&) = delete;

// 或实现移动构造与赋值（需要所有权转移场景）
Vec(Vec&& other) noexcept : size(other.size), data_ptr(other.data_ptr) {
    other.data_ptr = nullptr;
    other.size = 0;
}

Vec& operator=(Vec&& other) noexcept {
    if (this != &other) {
        if (data_ptr != nullptr) cudaFree(data_ptr);
        size = other.size;
        data_ptr = other.data_ptr;
        other.data_ptr = nullptr;
        other.size = 0;
    }
    return *this;
}

修复后的核心代码示例 #

修复后的点积内核 #

template <typename T>
__global__ void vec_dot(const T* vecA, const T* vecB, size_t size, T* out) {
    static_assert(std::is_arithmetic<T>::value, "T must be a numeric type");

    __shared__ T sdata[BLOCK_SIZE];

    int thread_index = threadIdx.x;
    sdata[thread_index] = 0;
    size_t idx = threadIdx.x + blockDim.x * blockIdx.x;

    while (idx < size) {
        sdata[thread_index] += vecA[idx] * vecB[idx]; 
        idx += blockDim.x * gridDim.x;   
    }

    __syncthreads();
    for (unsigned int s=blockDim.x/2; s>0; s>>=1) {
        if (thread_index < s) sdata[thread_index] += sdata[thread_index + s];
        __syncthreads();
    }
    if (thread_index == 0) atomicAdd(out, sdata[0]);
}

修复后的点积运算符 #

template <typename T> 
T operator%(const Vec<T>& A, const Vec<T>& B) {
    if (A.size != B.size) {
        throw std::invalid_argument("Vectors are not the same size.");
    }

    T* out_device; 
    cudaError_t err = cudaMalloc(&out_device, sizeof(T));
    if (err != cudaSuccess) {
        std::cerr << "Failed to allocate output device memory: " << cudaGetErrorString(err) << std::endl;
        std::terminate();
    }

    err = cudaMemset(out_device, 0, sizeof(T));
    if (err != cudaSuccess) {
        std::cerr << "Failed to initialize output memory: " << cudaGetErrorString(err) << std::endl;
        cudaFree(out_device);
        std::terminate();
    }

    dim3 grid((A.size + BLOCK_SIZE - 1)/BLOCK_SIZE);
    dim3 block(BLOCK_SIZE);
    vec_dot<T><<<grid, block>>>(A.data_ptr, B.data_ptr, A.size, out_device);

    err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "Kernel launch failed: " << cudaGetErrorString(err) << std::endl;
        cudaFree(out_device);
        std::terminate();
    }

    cudaError_t sync_err = cudaDeviceSynchronize();
    if (sync_err != cudaSuccess) {
        std::cerr << "Device synchronization failed: " << cudaGetErrorString(sync_err) << std::endl;
        cudaFree(out_device);
        std::terminate();
    }

    T out_host; 
    err = cudaMemcpy(&out_host, out_device, sizeof(T), cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        std::cerr << "Failed to copy result to host: " << cudaGetErrorString(err) << std::endl;
        cudaFree(out_device);
        std::terminate();
    }

    cudaFree(out_device);
    return out_host; 
}

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