我的CUDA学习之旅2——图像形态学腐蚀、膨胀CUDA实现

引言

由于上两周杂七杂八的事情比较多加上自己写的算法有些问题，一直改bug。。。。没时间继续写博客，今天开始补上博客。从这篇文章起，开始将一些较为典型的OpenCV算法通过CUDA进行实现，本文实现的为图像处理中最为常见的形态学腐蚀以及膨胀，由于本文目的在于算法移植后的验证，故在图片的选择上用小图像作为输入的示例图像，有不当之处欢迎评论或私信~

任务要求

输入一张图片，将其转为灰度图后，通过CUDA在GPU中对图片实现形态学腐蚀、膨胀操作，最后将结果输出至CPU并进行显示，要求输出图与用OpenCV实现后的结果一致。

实现思路

关于腐蚀与膨胀的算法原理网上已有完备的的资料，在这里不再复述，具体原理可见图像的腐蚀原理 
由于是对经典算法的移植，故在thread以及block的设计上不能单单针对某一张图片，而是要通用，同时为了尽可能提高运算速度，将其设计为32*32的1024个thread大小的block（本人显卡Nvidia GeForce 755 M），block数量则是根据传入图片的大小动态变化。

实现环境

VS2013 + CUDA7.5 + Opencv2.4.13

实现代码

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

//腐蚀
__global__ void erodeInCuda(unsigned char *dataIn, unsigned char *dataOut, Size erodeElement, int imgWidth, int imgHeight)
{
    //Grid中x方向上的索引
    int xIndex = threadIdx.x + blockIdx.x * blockDim.x;
    //Grid中y方向上的索引
    int yIndex = threadIdx.y + blockIdx.y * blockDim.y;

    int elementWidth = erodeElement.width;
    int elementHeight = erodeElement.height;
    int halfEW = elementWidth / 2;
    int halfEH = elementHeight / 2;
    //初始化输出图
    dataOut[yIndex * imgWidth + xIndex] = dataIn[yIndex * imgWidth + xIndex];;
    //防止越界
    if (xIndex > halfEW && xIndex < imgWidth - halfEW && yIndex > halfEH && yIndex < imgHeight - halfEH)
    {
        for (int i = -halfEH; i < halfEH + 1; i++)
        {
            for (int j = -halfEW; j < halfEW + 1; j++)
            {
                if (dataIn[(i + yIndex) * imgWidth + xIndex + j] < dataOut[yIndex * imgWidth + xIndex])
                {
                    dataOut[yIndex * imgWidth + xIndex] = dataIn[(i + yIndex) * imgWidth + xIndex + j];
                }
            }
        }
    }
}

//膨胀
__global__ void dilateInCuda(unsigned char *dataIn, unsigned char *dataOut, Size dilateElement, int imgWidth, int imgHeight)
{
    //Grid中x方向上的索引
    int xIndex = threadIdx.x + blockIdx.x * blockDim.x;
    //Grid中y方向上的索引
    int yIndex = threadIdx.y + blockIdx.y * blockDim.y;

    int elementWidth = dilateElement.width;
    int elementHeight = dilateElement.height;
    int halfEW = elementWidth / 2;
    int halfEH = elementHeight / 2;
    //初始化输出图
    dataOut[yIndex * imgWidth + xIndex] = dataIn[yIndex * imgWidth + xIndex];;
    //防止越界
    if (xIndex > halfEW && xIndex < imgWidth - halfEW && yIndex > halfEH && yIndex < imgHeight - halfEH)
    {
        for (int i = -halfEH; i < halfEH + 1; i++)
        {
            for (int j = -halfEW; j < halfEW + 1; j++)
            {
                if (dataIn[(i + yIndex) * imgWidth + xIndex + j] > dataOut[yIndex * imgWidth + xIndex])
                {
                    dataOut[yIndex * imgWidth + xIndex] = dataIn[(i + yIndex) * imgWidth + xIndex + j];
                }
            }
        }
    }
}


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


    unsigned char *d_in;//输入图片在GPU内的内存
    unsigned char *d_out1;//腐蚀后输出图片在GPU内的内存
    unsigned char *d_out2;//膨胀后输出图片在GPU内的内存

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

    Mat dstImg1(imgHeight, imgWidth, CV_8UC1, Scalar(0));//腐蚀后输出图片在CPU内的内存
    Mat dstImg2(imgHeight, imgWidth, CV_8UC1, Scalar(0));//膨胀后输出图片在CPU内的内存
    //在GPU中开辟内存
    cudaMalloc((void**)&d_in, imgWidth * imgHeight * sizeof(unsigned char));
    cudaMalloc((void**)&d_out1, imgWidth * imgHeight * sizeof(unsigned char));
    cudaMalloc((void**)&d_out2, imgWidth * imgHeight * sizeof(unsigned char));
    //将输入图片传入GPU
    cudaMemcpy(d_in, grayImg.data, imgWidth * imgHeight * sizeof(unsigned char), cudaMemcpyHostToDevice);
    //定义block中thread的分布
    dim3 threadsPerBlock(32, 32);
    //根据输入图片的宽高定义block的大小
    dim3 blocksPerGrid((imgWidth + threadsPerBlock.x - 1) / threadsPerBlock.x, (imgHeight + threadsPerBlock.y - 1) / threadsPerBlock.y);
    //算子大小
    Size Element(3, 5);
    //CUDA腐蚀
    erodeInCuda << <blocksPerGrid, threadsPerBlock >> >(d_in, d_out1, Element, imgWidth, imgHeight);
    //将结果传回CPU
    cudaMemcpy(dstImg1.data, d_out1, imgWidth * imgHeight * sizeof(unsigned char), cudaMemcpyDeviceToHost);
    //CPU内腐蚀（OpenCV实现）
    Mat erodeImg;
    Mat element = getStructuringElement(MORPH_RECT, Size(3, 5));
    erode(grayImg, erodeImg, element);
    //CUDA膨胀
    dilateInCuda << <blocksPerGrid, threadsPerBlock >> >(d_in, d_out2, Element, imgWidth, imgHeight);
    //将结果传回CPU
    cudaMemcpy(dstImg2.data, d_out2, imgWidth * imgHeight * sizeof(unsigned char), cudaMemcpyDeviceToHost);
    //CPU内膨胀（OpenCV实现）
    Mat dilateImg;
    dilate(grayImg, dilateImg, element);

    return 0;
}

原灰度图 

腐蚀后图片 

膨胀后图片 

通过比对发现CUDA输出结果与OpenCV输出结果一致~