#include <stdio.h>

/////////////////////
//kernel_hello_cuda//
/////////////////////
__global__ void kernel_hello_cuda (size_t num_chars, char* in, char* out) {
//klicove slovo __global__ definuje, ze se jedna o kernel

	//pokud bude id vlakna vetsi nez je pocet znaku v retezci, zkonci
	//nepotrebujeme vic vlaken, nez je nezbytne nutne (kopirujeme vcetne 12.znaku, tedy prazdneho znaku \0)
	if (threadIdx.x > num_chars) {
		return;
	}

	//kopirujeme obracene retezec !ADUC olleH tak, aby vznikl retezec Hello CUDA!
	//kazde vlakno vezme jeden znak (z retezce in) a ulozi ho na spravne misto (v retezci out)
	out[threadIdx.x] = in[num_chars-threadIdx.x-1];
}

//////////////////////
//hlavni funkce main//
//////////////////////
int main (int argc, char** argv) {
	//alokace pameti pro retezec znaku !ADUC olleH  (nesmime zapomenout na 12. znak \0)
	char* hello_cuda_h = (char*)malloc (sizeof(char)*12);
	
	//kopirovani retezce znaku !ADUC olleH
	strcpy (hello_cuda_h, "!ADUC olleH");

	//retezec znaku pro !ADUC olleH (alokace retezce bude na GPU)
	char* hello_cuda_reverse_d;

	//retezec znaku pro Hello CUDA! (alokace retezce bude na GPU)
	char* hello_cuda_d;

	//velikost retezce znaku v bytech
	int hello_cuda_size = sizeof(char)*(strlen(hello_cuda_h));

	//alokace retezce na GPU (konkretne v global memory) a predani ukazatele
	//prvni parametr: ukazatel na alokovanou pamet
	//druhy parametr: pozadovana velikost pro alokaci pameti
	//vystupni parametr: je enumerator cudaError_t (o nem v dalsich dilech)
	cudaMalloc ((void**) &hello_cuda_reverse_d, hello_cuda_size);

	//alokace retezce na GPU (konkretne v global memory) a predani ukazatele
	cudaMalloc ((void**) &hello_cuda_d, hello_cuda_size);

	//nastaveni prazdneho retezce hello_cuda_d
	cudaMemset (hello_cuda_d, 0, hello_cuda_size);

	//kopirovani retezce z pameti hosta do pameti zarizeni (GPU), konkretne do globalni pameti GPU
	//prvni parametr: ukazatel cilove pameti
	//druhy paremetr: ukazatel zdrojove pameti
	//treti parametr: velikost kopirovane pameti v bytech
	//ctvrty parametr: enumerator definujici z jakeho druhu pameti do jakeho druhu pameti bude kopirovano
	//vystupni parametr: je enumerator cudaError_t
	cudaMemcpy (hello_cuda_reverse_d, hello_cuda_h, hello_cuda_size, cudaMemcpyHostToDevice);

	//volame kernel
	//exekuce instrukci probiha jiz na zarizeni, nikoliv na hostovi!
	kernel_hello_cuda<<<1,17>>> (strlen(hello_cuda_h), hello_cuda_reverse_d, hello_cuda_d);

	//synchronizace vlaken
	//jinymi slovy: cekej dokud nezkonci vsechny vlakna
	cudaThreadSynchronize();

	//kopirovani retezce z pameti zarizeni zpet do pameti hosta
	cudaMemcpy (hello_cuda_h, hello_cuda_d, hello_cuda_size, cudaMemcpyDeviceToHost);

	//nyni je v retezci hello_cuda_h obracena veta Hello CUDA!, kterou vytiskneme
	printf ("%s\n", hello_cuda_h);

	//uvolneni prostredku v pameti zarizeni (GPU)
	//prvni parametr: ukazatel alokovane pameti pomoci funkce cudaMalloc (neda se pouzit pro pamet alokovanou funkci malloc!)
	//vystupni parametr: je enumerator cudaError_t
	cudaFree (hello_cuda_reverse_d);

	//uvolneni prostredku v pameti zarizeni (GPU)
	cudaFree (hello_cuda_d);

	//uvolneni prostredku v pameti hosta
	free (hello_cuda_h);

	//ukonceni programu
	return 0;
}