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


__global__ void noncoalesced_float3_kernel (float3* out, float3* in, float value) {
	//index vlakna v ramci vsech bloku
	int idx = blockIdx.x * blockDim.x + threadIdx.x;

	//prace s cislem typu float3
	in[idx].x += value;
	in[idx].y += value;
	in[idx].z += value;

	//ulozime element do globalni pameti do vystupniho pole
	out[idx] = in[idx];
}

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;

	//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
	noncoalesced_float3_kernel<<<num_blocks, num_threads_per_block>>> (array_out_d, 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;
}