我的CUDA学习之旅5——OTSU二值算法（最大类间方差法、大津法）CUDA实现

引言

最近在做医疗设备相关的项目，故在项目中大量用到了各类图像分割的算法，为了在图像中分割出特定目标，用到的算法可以有很多，比如阈值分割，多通道分割，边缘分割以及一些前沿的组合分割。而对大多数图像来说，分割的一大难点是将待识别的目标与背景分离，其中一种有效简单的方法就是二值化（并不都有效），本博客试着将二值化算法中的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;
}

实现结果

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

OpenCV实现结果图

CUDA实现后图像

关于本文以及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，也要定义数组，并把数组的头指针传给核函数