CUDA实现自适应阈值化异常：代码问题排查与修复咨询

问题描述 #

运行以下基于CUDA实现的自适应阈值化代码后，输出图像不符合预期（见下方原始图像与异常输出图），且调整blockDim尺寸过大或过小时功能无法正常生效，需修复代码中的问题以使其正常工作。

原始图像：

异常输出图像：

原始代码 #

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>

using namespace cv;
using namespace std;

// CUDA kernel function
__global__ void adaptiveThresholdCUDA(const uchar* image, uchar* binary, const int* iimage, int nl, int nc, int halfSize, int blockSize, int threshold) {
    // Calculate global indices
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // Check boundary conditions
    if (j >= nl - halfSize - 1 || i >= nc - halfSize - 1)
        return;

    // Get row addresses
    const uchar* data = image + j * nc;
    uchar* binaryRow = binary + j * nc;
    const int* idata1 = iimage + (j - halfSize) * nc;
    const int* idata2 = iimage + (j + halfSize + 1) * nc;

    // Calculate sum
    int sum = (idata2[i + halfSize + 1] - idata2[i - halfSize] - idata1[i + halfSize + 1] + idata1[i - halfSize]) / (blockSize * blockSize);

    // Apply adaptive threshold
    if (data[i] < (sum - threshold))
        binaryRow[i] = 0;
    else
        binaryRow[i] = 255;
}

int main()
{
    Mat image = imread("image/test.jpg", 0);
    if (!image.data) return 0;

    resize(image, image, Size(), 1.0, 1.0);
    namedWindow("Original Image");
    imshow("Original Image", image);

    /*          Function for Adaptive Thresholding          */
    int blockSize = 16; // Neighborhood size
    int threshold = 10; // Pixel comparison threshold

    Mat binary = image.clone();

    int nl = binary.rows; // Number of lines
    int nc = binary.cols; // Total number of elements per line

    Mat iimage;
    integral(image, iimage, CV_32S);

    // Allocate memory on device (GPU)
    uchar* imageDevice;
    uchar* binaryDevice;
    int* iimageDevice;
    cudaMalloc((void**)&imageDevice, sizeof(uchar) * nl * nc);
    cudaMalloc((void**)&binaryDevice, sizeof(uchar) * nl * nc);
    cudaMalloc((void**)&iimageDevice, sizeof(int) * nl * nc);

    // Copy input data from host to device
    cudaMemcpy(imageDevice, image.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
    cudaMemcpy(binaryDevice, binary.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
    cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * nl * nc, cudaMemcpyHostToDevice);

    // Define grid and block dimensions
    dim3 blockDim(16, 16);  // You may need to adjust the block dimensions
    dim3 gridDim((nc + blockDim.x - 1) / blockDim.x, (nl + blockDim.y - 1) / blockDim.y);

    // Launch the CUDA kernel
    adaptiveThresholdCUDA <<<gridDim, blockDim>>> (imageDevice, binaryDevice, iimageDevice, nl, nc, blockSize / 2, blockSize, threshold);

}

需修复的问题点 #

1. 边界条件检查错误 #

当前边界判断逻辑会错误丢弃有效像素，且未处理j < halfSize或i < halfSize的越界访问情况。修复为：

if (j < halfSize || j >= nl - halfSize || i < halfSize || i >= nc - halfSize)
    return;

2. 积分图索引与内存分配错误 #

OpenCV的integral函数返回的积分图尺寸为(nl+1)×(nc+1)，而非原图像的nl×nc，当前代码的索引计算和内存分配均不符合该规则：

• 内核中修正区域和计算逻辑：

int sum = (iimage[(j + halfSize + 1) * (nc + 1) + (i + halfSize + 1)] 
           - iimage[(j + halfSize + 1) * (nc + 1) + (i - halfSize)] 
           - iimage[(j - halfSize) * (nc + 1) + (i + halfSize + 1)] 
           + iimage[(j - halfSize) * (nc + 1) + (i - halfSize)]) 
           / (blockSize * blockSize);

• 主机端修正积分图的设备内存分配与拷贝：

cudaMalloc((void**)&iimageDevice, sizeof(int) * (nl + 1) * (nc + 1));
cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * (nl + 1) * (nc + 1), cudaMemcpyHostToDevice);

3. 缺少设备到主机的数据拷贝与结果展示 #

内核执行后未将GPU处理结果拷贝回主机，也没有显示输出图像。添加以下代码：

// 拷贝结果回主机
cudaMemcpy(binary.data, binaryDevice, sizeof(uchar) * nl * nc, cudaMemcpyDeviceToHost);
// 显示结果
namedWindow("Adaptive Threshold Result");
imshow("Adaptive Threshold Result", binary);
waitKey(0);

4. 未添加CUDA错误检查 #

所有CUDA操作（内存分配、数据拷贝、内核启动）均无错误检查，一旦出错无法排查。需为每个CUDA操作添加检查，例如：

// 内存分配错误检查
cudaError_t err = cudaMalloc((void**)&imageDevice, sizeof(uchar) * nl * nc);
if (err != cudaSuccess) {
    cerr << "cudaMalloc failed: " << cudaGetErrorString(err) << endl;
    return -1;
}
// 内核启动后检查
adaptiveThresholdCUDA <<<gridDim, blockDim>>> (...);
err = cudaGetLastError();
if (err != cudaSuccess) {
    cerr << "Kernel launch failed: " << cudaGetErrorString(err) << endl;
    return -1;
}
cudaDeviceSynchronize(); // 等待内核执行完成

5. BlockSize奇偶性处理问题 #

自适应阈值的邻域尺寸通常为奇数，当前blockSize=16（偶数）会导致邻域计算偏差。建议改为奇数，或添加强制奇数的逻辑：

int blockSize = 15; // 改为奇数
int halfSize = blockSize / 2;
// 或强制转为奇数
if (blockSize % 2 == 0) blockSize++;
halfSize = blockSize / 2;

6. 边界像素处理优化 #

当前代码直接跳过边界像素，导致输出图像边缘有未处理区域。可选择将边界像素统一设为固定值（如0或255），或对原图像做镜像填充后再处理，避免边缘缺失。

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