#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda.h>
#include "BitmapStruc.h"

#include <stdio.h>
#include <time.h>

#define DIM_BLOCK_LIN	256
#define DIM_GRID_LIN	1024
#define DIM_BLOCK_2D	16

#define ITERATION_COUNT	100

#define RUN_ON_GPU		1

//------------------------------------------------------------------------------
// Kernel Definíciók

// Kép simításra
__global__ void smoothKernel(uchar4* out, uchar4* in, int w, int h, int odd);

//------------------------------------------------------------------------------
// CPU-n futó segédfüggvények

// ugyanaz a számítás csak a CPU-n.
__host__ void smoothOnCPU();

//------------------------------------------------------------------------------
// main
int main()
{
	int i;
	unsigned int time;

	//	kép beolvasásáa
	BitmapStruct image("Lena.bmp");

	//	egy kics időmérés
	time = clock();

	//	tömbök a GPU-nál
	float* dev_src;
	float* dev_dst;

	//	grafikus memóriaterületek foglalása
	cudaMalloc((void**)&dev_src, 4 * image.x * image.y * sizeof(unsigned char));
	cudaMalloc((void**)&dev_dst, 4 * image.x * image.y * sizeof(unsigned char));

	//	kiinduló kép másolás a grafikus memóriába
	cudaMemcpy(dev_src, image.pixels, image.image_size(), cudaMemcpyHostToDevice);

	//	csak kiírjuk az időt
	printf("Kep beolvasva: %d s\n", clock() - time);

	//	blokkméret előkészítése
	dim3 grid((image.x + DIM_BLOCK_2D - 1) / DIM_BLOCK_2D, (image.y + DIM_BLOCK_2D - 1) / DIM_BLOCK_2D);
	dim3 block(DIM_BLOCK_2D, DIM_BLOCK_2D);

	//	iteratív simítás.
	for (i = 0; i<ITERATION_COUNT; i++) {
		if (i % 2) {
			//	ping
			//	páratlan cilkusban a 2. tömbből olvasunk és az eredményt az 1.-be írjuk
			smoothKernel << <grid, block >> >((uchar4*)dev_src, (uchar4*)dev_dst, image.x, image.y, 1);
		}
		else {
			//	pong
			//	páros cilkusban az 1. tömbből olvasunk és az eredményt a 2.-ba írjuk
			smoothKernel << <grid, block >> >((uchar4*)dev_dst, (uchar4*)dev_src, image.x, image.y, 0);
		}
	}

	//	Bevárjuk, hogy a kernelek lefussanak
	//	csak az időmérés miatt kell. a CPU futása aszinkrom megy elindítja a kerneleket
	//	  és megy tovább amíg nem kell az általuk szolgáltatott adat.
	cudaThreadSynchronize();

	//	Időmérés a számítás végén
	printf("Iterativ simitas kesz: %d s\n", clock() - time);

	//	és visszamásoljuk a központi memóriába
	cudaMemcpy(image.pixels, dev_src, image.image_size(), cudaMemcpyDeviceToHost);

	//	utolsó időmérés
	printf("Kep visszalakitva: %d s\n", clock() - time);

	//	program vége, memóriafelszabadítás
	cudaFree(dev_src);
	cudaFree(dev_dst);

	//	és az eredmény megjelenítése
	image.displayImage();

	return 0;
}

//------------------------------------------------------------------------------
// Kernel megvalósítások

// a simító kernel
__global__ void smoothKernel(uchar4* out, uchar4* in, int w, int h, int odd) {
	//	koordináták kiszámítása
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	float4 pixel;
	uchar4 temp;
	int i, j;
	int tx, ty;
	int index = w * y + x;

	//	ha a szél "lelóg" a képről akkor leálítjuk a szálat
	if (x >= w || y >= h)
		return;

	//	pixel értékeinek előkészítése
	pixel.x = 0;
	pixel.y = 0;
	pixel.z = 0;
	pixel.w = 0;

	//	bejárjuk a 3x3-as szomszédságot
	for (i = -1; i<2; i++) {
		for (j = -1; j<2; j++) {
			//	páros lépésben a egyik képről olvasunk, páratlanban a másikból.
			//	de ezt más a föprogramban kezeltük. :)
			tx = x + i;
			ty = y + j;

			if (tx < 0)
				tx = 0;

			if (ty < 0)
				ty = 0;

			if (tx >= w)
				tx = w-1;

			if (ty >= h)
				ty = h-1;

			//	Képelem beolvasása (1 db 32 bires olvasás)
			temp = in[ty*w + tx];

			pixel.x += temp.x;
			pixel.y += temp.y;
			pixel.z += temp.z;
			pixel.w += temp.w;
		}
	}

	//	osztás a maszk méretével
	pixel.x /= 9;
	pixel.y /= 9;
	pixel.z /= 9;
	pixel.w /= 9;

	//	vissza kell konvertálni az eredményt uchar4-re
	temp.x = pixel.x;
	temp.y = pixel.y;
	temp.z = pixel.z;
	temp.w = pixel.w;

	//	eredmény kiírása
	out[y*w + x] = temp;

	return;
}