#include <stdio.h>

__global__ void numbers (int* out, int* in) {
	//definujeme promennou s_data jako sdilenou pamet
	extern __shared__ int s_data[];

	// spocitame index do pole in a out pro konkretni vlakno
	// pripominam, ze pracujeme pouze s jednorozmernym polem,
	// tudiz jsme definovali x-slozku velikosti mrizky a bloku v execution configuration.
	int idx = blockDim.x * blockIdx.x + threadIdx.x;

	//delime cislo in[idx] dvema a ukladameho do sdilene pameti
	s_data[threadIdx.x] = in[idx]/2;

	//blokuj dokud tohoto bodu nedosahnou vsechny vlakna
	__syncthreads ();

	// ukladame do vystupniho pole
	// ma smysl, pokud idx nepresahne velikost vstupniho (a vystupniho) pole
	out[idx] = s_data[threadIdx.x];
}

int main (int argc, char** argv) {
	//pomocna promena
	bool result = true;

	//ukazatel na pole integeru pro praci s cisly v kernelu
	int* num_h;

	//ukazatel na pole int v globalni pameti
	int* num_d;

	//ukazatel na vystupni pole v globalni pameti
	int* num_out_d;

	//pozadovana velikost pole
	size_t num_size = 128*512;

	//pocet vlaken na jeden blok
	int num_threads_per_block = 128;

	//velikost mrizky
	int num_blocks = num_size/num_threads_per_block;

	//pozadovana velikost pole v bytech
	size_t num_size_bytes = sizeof (int)*num_size;

	//alokace pameti pole
	num_h = (int*)malloc (num_size_bytes);

	//alokujeme pole num_d v globalni pameti
	cudaMalloc ((void**) &num_d, num_size_bytes);

	//alokujeme vystupni pole num_out_d v globalni pameti
	cudaMalloc ((void**) &num_out_d, num_size_bytes);

	//naplnime pole cisly
	for (unsigned int i = 0; i < num_size; i++) {
		num_h[i] = i*2;
	}

	//kopirovani pole int z hosta do globalni pameti
	cudaMemcpy (num_d, num_h, num_size_bytes, cudaMemcpyHostToDevice);

	//spocitame velikost sdilene pameti (pocet vlaken v jednom bloku * velikost integeru)
	size_t shared_mem_size = num_threads_per_block*sizeof(int);

	//volame kernel
	numbers<<<num_blocks, num_threads_per_block, shared_mem_size>>> (num_out_d, num_d);

	//cekej, dokud vsechny vlakna nezkonci
	cudaThreadSynchronize();

	//nyni muzeme zkopirovat vystupni data num_out_d do num_h
	cudaMemcpy (num_h, num_out_d, num_size_bytes, cudaMemcpyDeviceToHost);

	//zkontrolujeme, zda vysledek je spravny
	for (unsigned int i = 0; i < num_size; i++) {
		if (i != num_h[i]) {
			result = false;
			printf ("Vysledek je spatny! %d\n", i);
			break;
		}
	}

	//pokud je vysledek spravny, vytiskneme oznameni
	if (result) {
		printf ("Vysledek je spravny!\n");
	}

	//uvolneni prostredku globalni pameti
	cudaFree (num_d);
	cudaFree (num_out_d);

	//uvolneni prostredku v pameti hosta
	free (num_h);
	return 0;
}