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));//传入灰度图至GPUcudaMemcpy(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, ndim3 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，也要定义数组，并把数组的头指针传给核函数