#include <stdlib.h>
#include <stdio.h>
#include <windows.h>
#include <opencv/cv.h>
#include <opencv/highgui.h>
#include <opencv2/opencv.hpp>
#include <opencv2/gpu/gpu.hpp>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
using namespace std;
#define width 800
#define height 600
__global__ static void erosion(uchar* d_src, uchar* d_dst, int len)
{
unsigned int i = blockIdx.x*blockDim.x+threadIdx.x;
unsigned int j = blockIdx.y*blockDim.y+threadIdx.y;
int wth = (len-1)/2;
if( i>0 && i < width && j>0 && j < height )
{
d_dst[j*width+i] = d_src[j*width+i];
}
if(i > wth && i < width - wth && j > wth && j < height - wth)
{
for(int w = -wth; w < wth+1; w++)
{
for(int h = -wth; h < wth+1; h++)
{
if(d_src[(j+h)*width+i+w] < d_dst[j*width+i])
d_dst[j*width+i] = d_src[(j+h)*width+i+w];
}
}
}
}
__global__ static void dilation(uchar* d_src, uchar* d_dst, int len)
{
unsigned int i = blockIdx.x*blockDim.x+threadIdx.x;
unsigned int j = blockIdx.y*blockDim.y+threadIdx.y;
int wth = (len-1)/2;
if( i>0 && i < width && j>0 && j < height )
{
d_dst[j*width+i] = d_src[j*width+i];
}
if(i > wth && i < width - wth && j > wth && j < height - wth)
{
for(int w = -wth; w < wth+1; w++)
{
for(int h = -wth; h < wth+1; h++)
{
if(d_src[(j+h)*width+i+w] > d_dst[j*width+i])
d_dst[j*width+i] = d_src[(j+h)*width+i+w];
}
}
}
}
void preprocess(uchar* d_src, uchar* d_tmp1, uchar* d_tmp2, uchar* d_dst)
{
dim3 dimBlock(16, 16);
dim3 dimGrid( (width+dimBlock.x-1)/dimBlock.x, (height+dimBlock.y-1)/dimBlock.y );
erosion<<<dimGrid, dimBlock, 0>>>(d_src, d_tmp1, 3);
dilation<<<dimGrid, dimBlock, 0>>>(d_tmp1, d_tmp2, 5);
erosion<<<dimGrid, dimBlock, 0>>>(d_tmp2, d_dst, 3);
cudaThreadSynchronize();
}
int main()
{
IplImage* src = cvLoadImage("F:\\VS2010 projects\\source\\002.bmp", 1);
IplImage* srcCopy = cvCreateImage(cvGetSize(src), IPL_DEPTH_8U, 1);
//转换成单通道
cvCvtColor(src, srcCopy, CV_RGB2GRAY);
unsigned char* d_img_src;
unsigned char* d_img_tmp1;
unsigned char* d_img_tmp2;
unsigned char* d_img_dst;
cudaMalloc((void**)&d_img_src, width*height);
cudaMalloc((void**)&d_img_tmp1, width*height);
cudaMalloc((void**)&d_img_tmp2, width*height);
cudaMalloc((void**)&d_img_dst, width*height);
cudaMemcpy(d_img_src, srcCopy->imageData, width*height, cudaMemcpyHostToDevice);
preprocess(d_img_src, d_img_tmp1, d_img_tmp2, d_img_dst);
cudaMemcpy( srcCopy->imageData, d_img_dst, width*height, cudaMemcpyDeviceToHost);
cvNamedWindow("test",CV_WINDOW_AUTOSIZE);
cvShowImage("test",srcCopy);
cvSaveImage("preprocess.bmp", srcCopy);
cvWaitKey(0);
cvDestroyAllWindows();
cvReleaseImage(&src);
cudaFree(d_img_src);
cudaFree(d_img_tmp1);
cudaFree(d_img_tmp2);
cudaFree(d_img_dst);
return 0;
}
CUDA用于图像处理中的形态学开闭运算
最新推荐文章于 2023-08-16 21:07:07 发布