gpu编程，图像部分缺失计算的问题

Question

使用gpu编程时，计算两张图的不同的代码如下，debug了两天了，因为是初学者，怎么也想不出是哪里的问题，导致图像一部分会缺失，不管裁剪多小的图，都会存在y方向的部分缺失，想知道这是什么问题？（ps 假设 cal difference没有错误）

__global__ void Cal_difference_Kernel(const uchar* R_max_tolerance, const uchar* R_min_tolerance, const uchar* I_max_tolerance, const uchar* I_min_tolerance, const uchar* src_image, const uchar* ref_model, uchar *dst, int rows, int cols) {

    
    int x = (threadIdx.x + blockIdx.x * blockDim.x);
    int y = (threadIdx.y + blockIdx.y * blockDim.y);

    if (x >= cols  || y >= rows) {
        return;
    }

    dst[y * cols + x] = cal_difference((int)R_max_tolerance[y*cols+x],(int)R_min_tolerance[y*cols+x],(int)I_max_tolerance[y*cols+x],(int)I_min_tolerance[y*cols+x],(int)src_image[y*cols+x], (int)ref_model[y*cols+x]); //设置dst为0或1 ，初始为0 


}


extern "C"
void Cal_difference_gpu(const Mat& ref,const Mat& src,const Mat& result) {

    auto t1 = std::chrono::system_clock::now();

    cv::cuda::GpuMat ref_model(ref);
    cv::cuda::GpuMat src_image(src);

    cv::cuda::GpuMat R_max_tolerance(ref.size(), CV_8UC1);
    cv::cuda::GpuMat R_min_tolerance(ref.size(), CV_8UC1);
    cv::cuda::GpuMat I_max_tolerance(ref.size(), CV_8UC1);
    cv::cuda::GpuMat I_min_tolerance(ref.size(), CV_8UC1);

    // ...这里省略了一些tolerance计算 

    cv::cuda::GpuMat Difference(ref.size(), CV_8UC1,Scalar(0));

    dim3 block(32, 32);
    dim3 grid((ref.rows -1 ) / 32 + 1 , (ref.cols - 1) / 32 +1 );  // 向上取整 
    Cal_difference_Kernel<<<grid, block>>>(R_max_tolerance.ptr<uchar>(),R_min_tolerance.data,I_max_tolerance.ptr<uchar>(),I_min_tolerance.ptr<uchar>(),src_image.ptr<uchar>(), ref_model.ptr<uchar>(), Difference.ptr<uchar>(), ref.rows, ref.cols);
    
    cudaDeviceSynchronize();

    Difference.download(result);

    //cudaFree((void*)differ_data);

}

缺失效果如下 ： （截了部分图）

Answer 1

引用 皆我百晓生 小程序回复内容作答：

根据提供的代码和缺失效果，问题可能出在以下几个方面：

1. 线程块大小和网格大小设置问题：在计算图像差异的核函数中，你设置了32x32的线程块大小和以线程块为单位计算网格大小。但你在计算网格大小时，使用了ref.rows和ref.cols而不是ref.cols和ref.rows。请确保将行和列的顺序正确设置。

正确的代码应该是：

dim3 grid((ref.cols - 1) / 32 + 1 , (ref.rows - 1) / 32 +1 );

1. 内存访问越界：在核函数中，你使用了threadIdx.x和threadIdx.y来计算图像中的位置索引。但是，你没有对x和y进行越界检查，导致可能访问到图像范围之外的内存区域。请确保在访问图像时进行正确的边界检查，以避免越界访问导致的数据缺失。

在修改核函数的内部if语句之前，可以添加以下代码来检查并调整x和y的取值范围：

x = max(min(x, cols - 1), 0);
y = max(min(y, rows - 1), 0);

并且在核函数中使用修正后的x和y的值：

dst[y * cols + x] = cal_difference(...);

这些修正可以帮助确保在计算图像差异时，不会访问超出图像范围的内存空间，从而解决图像部分缺失的问题。

如果问题仍然存在，请提供更多关于图像和计算差异函数(cal_difference)的详细信息，以便我能够提供更具体的帮助。