ImageProcessing 에 사용된 기술(Adaptive Binarization)을 GPU로 처리해보자

# include "cuda_runtime.h"
# include "device_launch_parameter.h"

# include <stdio.h>
// gpu에 할당할 메모리를 저장할 변수를 선언한다.
// 하나는 호스트에서 받아올 용도 다른 하나는 호스트로 넘겨줄 용도
unsigned char* g_tempBuffer[2] = { nullptr };

// gpu에 할당된 메모리를 0으로 세팅하는 함수를 정의한다.
// 영상처리 할 것이므로 width와 height를 인자로 받는다.
int ImageProc_InitializeMemory(int width, int height)
{
	cudaError cudaStatus;
	cudaStatus = cudaMemset(g_tempBuffer[0], 0,
		sizeof(unsigned char)*width*height);
	// 첫번째 메모리 세팅에 실패하면 -1을 반환한다.
	if (cudaStatus != cudaSuccess)
		return -1;

	cudaStatus = cudaMemset(g_tempBuffer[1], 0,
		sizeof(unsigned char)*width*height);
	// 두번째 메모리 세팅에 실패하면 -2을 반환한다.
	if (cudaStatus != cudaSuccess)
		return -2;

	return 1;
}

// gpu에 메모리를 할당하는 함수를 정의한다.
// 해당 기능은 외부에서 쓰이므로 extern "C" 선언을 해준다.
extern "C"
int ImageProc_AllocGPUMemory(int width, int height)
{
	cudaError cudaStatus;
	cudaStatus = cudaMalloc((void**)&g_tempBuffer[0], 
		sizeof(unsigned char)*width*height);
	// 첫번째 메모리 할당에 실패하면 -1을 반환한다.
	if (cudaStatus != cudaSuccess)
		return -1;

	cudaStatus = cudaMalloc((void**)&g_tempBuffer[1],
		sizeof(unsigned char)*width*height);
	// 두번째 메모리 할당에 실패하면 -2을 반환한다.
	if (cudaStatus != cudaSuccess)
		return -2;

	// 할당한 메모리를 0 으로 세팅한다.
	return ImageProc_InitializeMemory(width, height);
}

// gpu에 할당한 메모리를 반환하는 기능을 정의한다.
int ImageProc_DeAllocGPUMemory(void)
{
	cudaError cudaStatus;
	cudaStatus = cudaFree(g_tempBuffer[0]);
	// 첫번째 메모리 반환에 실패하면 -1을 반환한다.
	if (cudaStatus != cudaSuccess)
		return -1;

	cudaStatus = cudaFree(g_tempBuffer[1]);
	// 두번째 메모리 반환에 실패하면 -2을 반환한다.
	if (cudaStatus != cudaSuccess)
		return -2;

	return 1;
}

...
//CUDAImageProc.cu 에 정의한 함수를 쓰기위해 선언한다.
extern "C"
{
    int ImageProc_AllocGPUMemory(int width, int height);
	int ImageProc_DeAllocGPUMemory();
}

class ImageProc
{
    ...
public:
	bool AllocateGPUMemory(int width, int height);
	bool DeAllocateGPUMemory(void);
}

bool ImageProc::AllocateGPUMemory(int width, int height)
{
	if (ImageProc_AllocGPUMemory(width, height))
		return true;
	else
		return false;
}

bool ImageProc::DeAllocateGPUMemory(void)
{
	if (ImageProc_DeAllocGPUMemory())
		return true;
	else
		return false;
}

extern "C"
int ImageProc_AdaptiveBinarization(unsigned char* image_gray,
	int width, int height, int ksize)
{
	// GPU에 메모리 할당하고 0으로 세팅한다.
	if (ImageProc_AllocGPUMemory(width, height) < 0)
		return -1;

	// GPU 메모리에 호스트 데이터 복사한다.
	cudaError cudaStatus;
	cudaStatus = cudaMemcpy(g_tempBuffer[0],image_gray,
		sizeof(unsigned char)*width*height, cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess)
		return -2;

	// 커널 함수를 실행한다.
	// GridDim.x, GridDim.y, BlockDim.x, BlockDim.y 를 정의한다.
	dim3 Db = dim3(8,8);
	dim3 Dg = dim3((width + Db.x - 1) / Db.x, (height + Db.y - 1) / Db.y);
	Kernel_AdaptiveBinarization<<< Dg, Db >>> (g_tempBuffer[0],
		g_tempBuffer[1],width,height,ksize);

	// 커널 함수 실행이 제대로 되었는지 확인한다.
	cudaStatus = cudaGetLastError();
	if (cudaStatus != cudaSuccess)
		return -3;


	// 커널 함수 모두 종료를 확인한다.
	cudaStatus = cudaDeviceSynchronize();
	if (cudaStatus != cudaSuccess)
		return -4;

	// 결과를 호스트 메모리로 복사한다.
	cudaStatus = cudaMemcpy(image_gray, g_tempBuffer[1],
		sizeof(unsigned char)*width*height, cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess)
		return -5;

	return 1;
	
}

__global__ void Kernel_AdaptiveBinarization(unsigned char* image_gray,
	unsigned char* output_image, int width, int height, int ksize)
{
	if (ksize == 1 || ksize % 2 == 0) return;
	int neighbor = ksize / 2;
	// blockIdx, blockDim, threadIdx 로 좌표를 찾는다.
	// blockDim 은 block 사이즈 이므로 아래와 같이 좌표를 구할 수 있다.
	int i = blockIdx.x*blockDim.x+threadIdx.x;
	int j = blockIdx.y*blockDim.y+threadIdx.y;

	float avg = 0.f;
	float cnt = 0.f;

	for (int x = -neighbor; x <= neighbor; x++)
	{
		for (int y = -neighbor; y <= neighbor; y++)
		{
			if (i + x < 0 || i + x >= width || j + y < 0 || j + y >= height)
				continue;
			avg += image_gray[width*(j + y) + (i + x)];
			cnt += 1.f;
		}
	}
	avg = avg / cnt;
	if (image_gray[width*j + i] > avg)
		output_image[width*j + i] = 255;
	else
		output_image[width*j + i] = 0;
}

extern "C"
{
	...
	int ImageProc_AdaptiveBinarization(image_gray,width,height,ksize);
}
...
// ImageProc 에서 정의할 함수도 선언한다. .. 생략.

bool ImageProc::GPU_AdaptiveBinarization(unsigned char* image_gray,
	int width, int height, int ksize)
{
	printf("GPU_AdaptiveBinarization\n");
	int res = ImageProc_AdaptiveBinarization(image_gray, width, height, ksize);
	if (res < 0)
	{
		printf("GPU_AdaptiveBinarization failed error code is %d \n,",res);
		return false;
	}
	else
		return true;
}

class CImageProcessingDoc: public CDocument
{
public:
	// ImageProc 객체를 참조하는 변수를 선언
	ImageProc* obj_ImageProc;
}

CImageProcessingDoc::CImageProcessingDoc()
{
	// TODO: 여기에 일회성 생성 코드를 추가합니다.
	...
	obj_ImageProc = nullptr;
}

CImageProcessingDoc::~CImageProcessingDoc()
{
	...
	if (obj_ImageProc)
		delete obj_ImageProc;
}

BOOL CImageProcessingDoc::OnNewDocument()
{
	if (!CDocument::OnNewDocument())
		return FALSE;

	if (obj_ImageProc)
		delete obj_ImageProc;
	obj_ImageProc = new ImageProc;

	return TRUE;
}

void CImageProcessingDoc::OnImageprocessinggpuAdaptivebinarization()
{
	// TODO: 여기에 명령 처리기 코드를 추가합니다.
	printf("OnImageprocessing GPU Adaptivebinarization\n");

	QueryPerformanceFrequency(&Frequency);
	QueryPerformanceCounter(&BeginTime);

	obj_ImageProc->GPU_AdaptiveBinarization(m_Images[cur_index].image_gray,
		m_Images[cur_index].width, m_Images[cur_index].height, 87);

	QueryPerformanceCounter(&Endtime);
	int elapsed = Endtime.QuadPart - BeginTime.QuadPart;
	double duringtime = (double)elapsed / (double)Frequency.QuadPart;
	printf("GPU AdaptiveBinarization time : %f\n", duringtime);

	// GPU Memory DeAllocation
	obj_ImageProc->DeAllocateGPUMemory();


	CImageProcessingView* pView = (CImageProcessingView*)((CMainFrame*)(AfxGetApp()->m_pMainWnd))->GetActiveView();

	pView->SetDrawImage(m_Images[cur_index].image_color, m_Images[cur_index].image_gray,
		m_Images[cur_index].width, m_Images[cur_index].height, 1);

	pView->OnInitialUpdate();
}