/*-------------------------------------------------------------------------
裡面比較重要的是
			InputArray;getGpuMat();PtrStepSzb;
			這麼幾個資料型別和函式
-------------------------------------------------------------------------*/


#include "cuda_main.h"
#include <iostream>
#include <opencv2/opencv.hpp>
#include "opencv2/cudaimgproc.hpp"
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudafilters.hpp"

#include <cuda_runtime.h>
#include <cufft.h>

#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 << <blocksPerGrid, threadsPerBlock >> >(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<cuda::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;
}