opencv 에서 cuda 모듈 의 데이터 구조 간이 용법

/*-------------------------------------------------------------------------
        
			InputArray;getGpuMat();PtrStepSzb;
			           
-------------------------------------------------------------------------*/


#include "cuda_main.h"
#include 
#include 
#include "opencv2/cudaimgproc.hpp"
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudafilters.hpp"

#include 
#include 

#include "opencv2/core/cuda_stream_accessor.hpp"
#include "opencv2/core/cuda_types.hpp"

using namespace std;
using namespace cv;
using namespace cv::cuda;


//  
__global__ void erodeInCuda1(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 erodeInCuda(const PtrStepSzb dataIn, PtrStepSzb 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.ptr(yIndex)[xIndex] = dataIn.ptr(yIndex)[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.ptr(i + yIndex)[xIndex + j] < dataOut.ptr(yIndex)[xIndex])
				{
					dataOut.ptr(yIndex)[xIndex] = dataIn.ptr(i + yIndex)[xIndex + j];
				}
			}
		}
	}
}


void xyz(const PtrStepSzb src_image, PtrStepSzb dst_image)
{
	int imgWidth = src_image.cols;
	int imgHeight = src_image.rows;
	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 << > >(src_image, dst_image, Element, imgWidth, imgHeight);
}

bool inputarray_test(InputArray arra)
{
	GpuMat src = arra.getGpuMat();

	Mat dst;
	src.download(dst);
	if (dst.empty())
		return false;
	return true;
}


int test_Gpu_mat()
{
	cuda::setDevice(0);

	cv::Mat src_host = cv::imread("F00005252.bmp", CV_LOAD_IMAGE_GRAYSCALE);  //           
	int imgWidth = src_host.cols;
	int imgHeight = src_host.rows;

	cv::cuda::GpuMat dst, src;

	//src.create(src_host.size(), CV_8UC1);
	src.upload(src_host);

	//   
	cv::cuda::threshold(src, dst, 128.0, 255.0, CV_THRESH_BINARY);

	//    
	//Ptr<:filter> m_medianFilter = cuda::createMedianFilter(CV_8UC1, 3, 128);
	//m_medianFilter->apply(src, dst);

	
	xyz(src, dst);
	bool ret = inputarray_test(src);

	cv::Mat result_host;
	dst.download(result_host);

	if (result_host.empty())
		ret = false;

	if (ret)
		cout << "GpuMat    ！" << endl;

	return 0;
}