#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

texture<float4, 2, cudaReadModeElementType> tex_in;
texture<float4, 2, cudaReadModeElementType> tex_out;

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

// byte <-> float konverziókra
__global__ void byte2FloatKernel(float* dst, unsigned char* src, int count);
__global__ void float2ByteKernel(unsigned char* dst, float* src, int count);

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

//------------------------------------------------------------------------------
// CPU-n futó segédfüggvények
__host__ void convByte2FloatGPU(float* dev_dst, unsigned char* dev_src, int count);
__host__ void convFloat2ByteGPU(unsigned char* dev_dst, float* dev_src, int count);

// 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
	unsigned char* dev_byte_ptr;
	float* dev_src;
	float* dev_dst;

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

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

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

	//	kép konvertálás float formátumra
	convByte2FloatGPU(dev_src, dev_byte_ptr, 4*image.x*image.y);

	//	textúrák előkészítése
	cudaChannelFormatDesc chDesc = cudaCreateChannelDesc<float4>();
	cudaBindTexture2D(NULL, tex_in, dev_src, chDesc, image.x, image.y, 4*image.x*sizeof(float));
	cudaBindTexture2D(NULL, tex_out, dev_dst, chDesc, image.x, image.y, 4*image.x*sizeof(float));

	//	még egy idő kiírás
	printf("Kep konvertalva es elokeszitve: %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<0; i++) {
		if(i%2) {
			//	ping
			//	páratlan cilkusban a 2. textúrából olvasunk és az eredményt az 1.-be írjuk
			smoothKernel<<<grid, block>>>((float4*)dev_src, image.x, image.y, 1);
		} else {
			//	pong
			//	páros cilkusban az 1. textúrából olvasunk és az eredményt a 2.-ba írjuk
			smoothKernel<<<grid, block>>>((float4*)dev_dst, 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();

	printf("Iterativ simitas kesz: %d s\n", clock()-time);

	//	ha kész van az eredmény, akkor viszakonvertáljuk byte-tömbbé
	convFloat2ByteGPU(dev_byte_ptr, dev_src, 4*image.x*image.y);

	//	és visszamásoljuk a központi memóriába
	cudaMemcpy(image.pixels, dev_byte_ptr, 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_byte_ptr);
	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(float4* out, 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;
	float4 temp;
	int i, j;

	//	ha a szél "lelóg" a képről akkor leálítjuk
	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 textúrából olvasunk, páratlanban a másikból
			//	a kilógó pixelekkel nem foglalkozunk. a textúrázó lekezeli.
			if(odd) {
				temp = tex2D(tex_out, x+i, y+j);
			} else {
				temp = tex2D(tex_in, x+i, y+j);
			}
			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;

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

	return;
}

//	byte -> float konverzió
__global__ void byte2FloatKernel(float* dst, unsigned char* src, int count) {
	//	kiszámítjuk az első konvertálandó elem pozícióját.
	int tid = blockDim.x * blockIdx.x + threadIdx.x;

	//	majd ciklikusan
	while(tid < count) {
		//	konvertálunk (csak egy másolás)
		dst[tid] = src[tid];
		//	és lépünk a többi szál által konvertált értékeken túlra
		tid += blockDim.x * gridDim.x;
	}

	return;
}

//	float -> byte konverzió
__global__ void float2ByteKernel(unsigned char* dst, float* src, int count) {
	int tid = blockDim.x * blockIdx.x + threadIdx.x;

	while(tid < count) {
		dst[tid] = src[tid];
		tid += blockDim.x * gridDim.x;
	}

	return;
}

//------------------------------------------------------------------------------
// CPU segédfüggvény megvalósítások

// keret a byte->float konverzióhoz
__host__ void convByte2FloatGPU(float* dev_dst, unsigned char* dev_src, int count) {
	//	csak kiszámítjuk az indított rács strukturáját, és elindítuk a megfelelő kernelt.
	dim3 grid;
	dim3 block;
	grid.x = min(DIM_GRID_LIN, (count + DIM_BLOCK_LIN - 1)/DIM_BLOCK_LIN);
	block.x = DIM_BLOCK_LIN;

	byte2FloatKernel<<<grid, block>>>(dev_dst, dev_src, count);

	return;
}

// keret a byte->float konverzióhoz
__host__ void convFloat2ByteGPU(unsigned char* dev_dst, float* dev_src, int count) {
	dim3 grid;
	dim3 block;
	grid.x = min(DIM_GRID_LIN, (count + DIM_BLOCK_LIN - 1)/DIM_BLOCK_LIN);
	block.x = DIM_BLOCK_LIN;

	float2ByteKernel<<<grid, block>>>(dev_dst, dev_src, count);

	return;
}