#include <stdio.h>
#include <assert.h>


__global__ void coalesced_float3_kernel (float* out, float* in, float value) {
	//sdilena pamět, velikost je definovana ve funkci main
	extern __shared__ float s_data[];

	//index vlakna v ramci vsech bloku
	int index = blockIdx.x * blockDim.x + threadIdx.x;

	//ulozime jednotlive slozky float3 do sdilene pameti
	//v s_data je pouzit threadIdx, protoze se jedna o pamet v ramci jednoho bloku
	//sdruzeny pristup do globalni pameti
	s_data[threadIdx.x] = in[index];
	s_data[threadIdx.x+64] = in[index+64];
	s_data[threadIdx.x+128] = in[index+128];

	//nez budeme pokracovat dale, je nutne vsechny zkopirovat z globalni pameti do sdilene
	__syncthreads();

	//pretypujeme pozadovane cislo na float3
	float3 number = ((float3*)s_data)[threadIdx.x];

	//prace s cislem typu float3
	number.x += value;
	number.y += value;
	number.z += value;

	//pretypovani cisla typu float3 zpet na float* pole
	((float3*)s_data)[threadIdx.x] = number;

	//pockame dokud neni prace s cislem dokoncena u vsech vlaken
	__syncthreads();

	//vse prekopirujeme ze sdilene pameti na vystup do globalni pameti
	//sdruzeny pristup do globalni pameti
	out[index] = s_data[threadIdx.x];
	out[index+64] = s_data[threadIdx.x+64];
	out[index+128] = s_data[threadIdx.x+128];
}

int main (int argc, char** argv) {
	//pole pro praci s float3
	float3 *array_in_h, *array_in_d, *array_out_h, *array_out_d;

	//pomocna promenna
	float value;

	//velikost pole (pocet float3 elementu v poli)
	size_t array_size = 8*64;

	//velikost pole v bytech
	size_t array_size_bytes = array_size * sizeof(float3);

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

	//pocet bloku
	int num_blocks = array_size/num_threads_per_block;

	//velikost sdilene pameti v bytech (sdilena pamet je pouze v ramci jednoho bloku!)
	int shared_mem_size = num_threads_per_block * sizeof (float3);

	//alokace pozadovane pameti pro pole typu float3
	array_in_h = (float3*)malloc (array_size_bytes);

	//naplneni pole cisly
	for (unsigned int idx = 0; idx < array_size; idx++) {
		array_in_h[idx].x = (float)idx;
		array_in_h[idx].y = (float)idx;
		array_in_h[idx].z = (float)idx;
	}

	//alokace vystupniho a vstupniho pole v globalni pameti
	cudaMalloc ((void**) &array_out_d, array_size_bytes);
	cudaMalloc ((void**) &array_in_d, array_size_bytes);

	//kopirovani naplneneho pole z hosta do globalni pameti zarizeni
	cudaMemcpy (array_in_d, array_in_h, array_size_bytes, cudaMemcpyHostToDevice);

	//hodnota, ktera se bude pricitat k jednotlivym slozkam (x, y, z) elementu v poli typu float3
	value = 3.0f;

	//volani kernelu
	coalesced_float3_kernel<<<num_blocks, num_threads_per_block, shared_mem_size>>> ((float*)array_out_d, (float*)array_in_d, value);

	//pockame dokud nezkonci vsechny vlakna
	cudaThreadSynchronize ();

	//alokace pameti pro vystupni pole (v hostu)
	array_out_h = (float3*)malloc (array_size_bytes);

	//kopirovani vystupniho pole v globalni pameti do vystupniho pole v pameti hosta
	cudaMemcpy (array_out_h, array_out_d, array_size_bytes, cudaMemcpyDeviceToHost);

	//kontrola, zda soucet je v poradku
	for (unsigned int idx = 0; idx < array_size; idx++) {
		assert (array_out_h[idx].x == idx+value);
		assert (array_out_h[idx].y == idx+value);
		assert (array_out_h[idx].z == idx+value);
	}

	//pokud ano, vypise se zprava nize:
	printf ("Ok!\n");

	//uvolneni alokovane pameti vstupniho a vystupniho pole v pameti hosta
	free (array_in_h);
	free (array_out_h);

	//uvolneni alokovane pameti vstupniho a vystupniho pole v globalni pameti zarizeni
	cudaFree (array_in_d);
	cudaFree (array_out_d);

	//ukonceni programu
	return 0;
}