CUDA精进之路(五):图像处理——OTSU二值算法(最大类间方差法、大津法)

引言


最近在做医疗设备相关的项目,故在项目中大量用到了各类图像分割的算法,为了在图像中分割出特定目标,用到的算法可以有很多,比如阈值分割,多通道分割,边缘分割以及一些前沿的组合分割。而对大多数图像来说,分割的一大难点是将待识别的目标与背景分离,其中一种有效简单的方法就是二值化(并不都有效),本博客试着将二值化算法中的OTSU算法进行cuda改写。

任务要求


输入一张8bit的灰度图,通过CUDA在GPU中对图片实现otsu二值化,最后将结果输出至CPU并进行显示,要求输出图与用CPU内实现后的结果一致。

实现思路


关于OTSU(大津法)二值算法的具体实现思路,具体可见此博文 最大类间方差法(大津法OTSU) 
该文对最终的方差计算公式进行了一定的变形,减小了总体的计算量。 
通过对OTSU算法的阅读,可以发现在遍历计算最大类间方差时,程序存在着一定的时序性,故解决方案为通过并行计算出所有需要的数据,通过数组进行保存,在时序性算法部分(这里指最大值寻找)仍然利用串行的手法,如此完成算法的改写。 
实现过程: 
1.统计图像灰度直方图hist数组 
2.计算图像最大类间方差 
3.根据计算出的最佳阈值进行二值化操作

实现环境


VS2013 + CUDA7.5 + Opencv2.4.13

实现代码

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <cuda.h>
#include <device_functions.h>
#include <iostream>
#include <string.h>
#include <opencv2\opencv.hpp>
using namespace std;
using namespace cv;

/*
计算最大类间方差串行程序
只能在CPU端调用,需要将hist数组传出才可计算
计算量变大时(大图像)速度较慢
*/
__host__ int otsuThresh(int* hist, int imgHeight, int imgWidth)
{
    float sum = 0;
    for (int i = 0; i < 256; i++)
    {
        sum += i * hist[i];
    }
    float w0 = 0, u0 = 0;
    float u = sum / (imgHeight * imgWidth);
    float val = 0, maxval = 0;
    float s = 0, n = 0;
    int thresh = 0;
    for (int i = 0; i < 256; i++)
    {
        s += hist[i] * i;
        n += hist[i];
        w0 = n / (imgHeight * imgWidth);
        u0 = s / n;
        val = (u - u0) * (u - u0) * w0 / (1 - w0);
        if (val > maxval)
        {
            maxval = val;
            thresh = i;
        }
    }
    return thresh;
}

//灰度直方图统计
__global__ void imhistInCuda(unsigned char* dataIn, int* hist, int imgHeight, int imgWidth)
{
    int xIndex = threadIdx.x + blockIdx.x * blockDim.x;
    int yIndex = threadIdx.y + blockIdx.y * blockDim.y;

    if (xIndex < imgWidth && yIndex < imgHeight)
    {
        atomicAdd(&hist[dataIn[yIndex * imgWidth + xIndex]], 1);
    }
}

//计算最大类间方差CUDA改编程序
__global__ void OTSUthresh(const int* hist, float* sum, float* s, float* n, float* val, int imgHeight, int imgWidth, int* OtsuThresh)
{
    if (blockIdx.x == 0)
    {
        int index = threadIdx.x;
        atomicAdd(&sum[0], hist[index] * index);
    }
    else
    {
        int index = threadIdx.x;
        if (index < blockIdx.x)
        {
            atomicAdd(&s[blockIdx.x - 1], hist[index] * index);
            atomicAdd(&n[blockIdx.x - 1], hist[index]);
        }
    }
    __syncthreads(); //所有线程同步
    if (blockIdx.x > 0)
    {
        int index = blockIdx.x - 1;
        float u = sum[0] / (imgHeight * imgWidth);
        float w0 = n[index] / (imgHeight * imgWidth);
        float u0 = s[index] / n[index];
        if (w0 == 1)
        {
            val[index] = 0;
        }
        else
        {
            val[index] = (u - u0) * (u - u0) * w0 / (1 - w0);
        }
    }
    __syncthreads(); //所有线程同步
    if (threadIdx.x == 0 && blockIdx.x == 0)
    {
        float maxval = 0;
        for (int i = 0; i < 256; i++)
        {
            if (val[i] > maxval)
            {
                maxval = val[i];
                OtsuThresh[0] = i;
                OtsuThresh[1] = val[i];
            }
        }
    }
}

//阈值化
__global__ void otsuInCuda(unsigned char* dataIn, unsigned char* dataOut, int imgHeight, int imgWidth, int* hThresh)
{
    int xIndex = threadIdx.x + blockIdx.x * blockDim.x;
    int yIndex = threadIdx.y + blockIdx.y * blockDim.y;

    if (xIndex < imgWidth && yIndex < imgHeight)
    {
        if (dataIn[yIndex * imgWidth + xIndex] > hThresh[0])
        {
            dataOut[yIndex * imgWidth + xIndex] = 255;
        }
    }
}

int main()
{
    //传入灰度图
    Mat srcImg = imread("1.jpg", 0);

    int imgHeight = srcImg.rows;
    int imgWidth = srcImg.cols;

    //opencv实现OTSU二值化
    Mat dstImg1;
    threshold(srcImg, dstImg1, 0, 255, THRESH_OTSU);

    //CUDA改编
    Mat dstImg2(imgHeight, imgWidth, CV_8UC1, Scalar(0));

    //在GPU端开辟内存
    unsigned char* d_in;
    int* d_hist;

    cudaMalloc((void**)&d_in, imgHeight * imgWidth * sizeof(unsigned char));
    cudaMalloc((void**)&d_hist, 256 * sizeof(int));

    //传入灰度图至GPU
    cudaMemcpy(d_in, srcImg.data, imgHeight * imgWidth * sizeof(unsigned char), cudaMemcpyHostToDevice);

    dim3 threadsPerBlock1(32, 32);
    dim3 blocksPerGrid1((imgWidth + 32 - 1) / 32, (imgHeight + 32 - 1) / 32);

    imhistInCuda << <blocksPerGrid1, threadsPerBlock1 >> >(d_in, d_hist, imgHeight, imgWidth);

    float* d_sum;
    float* d_s;
    float* d_n;
    float *d_val;
    int* d_t;

    cudaMalloc((void**)&d_sum, sizeof(float));
    cudaMalloc((void**)&d_s, 256 * sizeof(float));
    cudaMalloc((void**)&d_n, 256 * sizeof(float));
    cudaMalloc((void**)&d_val, 256 * sizeof(float));
    cudaMalloc((void**)&d_t, 2 * sizeof(int));

    //定义最大类间方差计算并行规格,其中257为1 + 256,
    //第1个block用来计算图像灰度的sum,后256个block用于计算256个灰度对应的s, n
    dim3 threadsPerBlock2(256, 1);
    dim3 blocksPerGrid2(257, 1);

    OTSUthresh << <blocksPerGrid2, threadsPerBlock2 >> >(d_hist, d_sum, d_s, d_n, d_val, imgHeight, imgWidth, d_t);

    unsigned char* d_out;

    cudaMalloc((void**)&d_out, imgHeight * imgWidth * sizeof(unsigned char));

    otsuInCuda << <blocksPerGrid1, threadsPerBlock1 >> >(d_in, d_out, imgHeight, imgWidth, d_t);

    //输出结果图像
    cudaMemcpy(dstImg2.data, d_out, imgHeight * imgWidth * sizeof(unsigned char), cudaMemcpyDeviceToHost);

    ////调试用输出
    //int th[2] = { 0, 0 };
    //float n[256];
    //memset(n, 0, sizeof(n));
    //cudaMemcpy(th, d_t, 2 * sizeof(int), cudaMemcpyDeviceToHost);
    //cudaMemcpy(n, d_n, 256 * sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(d_in);
    cudaFree(d_out);
    cudaFree(d_hist);
    cudaFree(d_sum);
    cudaFree(d_s);
    cudaFree(d_n);
    cudaFree(d_val);
    cudaFree(d_t);

    imwrite("result1.jpg", dstImg1);
    imwrite("result2.jpg", dstImg2);

    return 0;
}


实现结果


贴上我大酷玩演唱会图,进行验证~ 

原图 

CUDA精进之路(五):图像处理——OTSU二值算法(最大类间方差法、大津法)
 

OpenCV实现结果图 

CUDA精进之路(五):图像处理——OTSU二值算法(最大类间方差法、大津法)

CUDA实现后图像 

CUDA精进之路(五):图像处理——OTSU二值算法(最大类间方差法、大津法)

关于本文以及CUDA的一些思考
在本文的OTSU算法中,其实在改编的过程中,一直怀疑把那段计算最大类间方差的串行代码(代码中host部分)改成部分并行部分串行并没有起到提速的作用,而事实上我自己做了测速,发现确实基本没什么区别,分析了下原因在于: 
1.计算量不大,GPU加深没发挥真正的作用 
2.改写的过程涉及时序的部分只能采用串行,而串行的效率GPU反而低于CPU 
3.算法还未优化至最佳 
而从刚开始接触CUDA到现在陆陆续续也写了不少CUDA的代码了,给刚入坑的朋友提几点不成熟的建议~: 
1.CUDA不是万能的,很多时候一些复杂的算法无法改写,尤其是涉及时序性的 
2.加速的关键在于对于GPU内存的使用规划以及原算法的性能 
3.有成熟的库函数能用的时候可不用CUDA,因为提速效果不是很明显(可能是我显卡渣的原因==),例如上面的OTSU算法,OpenCV实现20ms左右,CUDA实现10ms左右 
4.CUDA的时间花费大部分都在GPU传至CPU上(一般占总时间50%以上),所以在进行编码的时候能不传数据出来就尽量不要传,尽量在GPU中完成所有算法的实现,争取一进一出~ 
5.CUDA中传变量只能通过数组的形式,所以就算你的变量数量为1,也要定义数组,并把数组的头指针传给核函数
 

上一篇:CUDA精进之路(五):图像处理——OTSU二值算法(最大类间方差法、大津法)


下一篇:net core 中间件详解及项目实战