// odwracanie kolejności elementów w tablicy jednowymiarowej

#include <stdio.h>
#include <windows.h>
#include <limits.h>


// liczba wątków
const int N_THR = 512;

// rozmiar tablicy
const int N = (1 << 23);

//////////////////////////////////////////////////////////////////
// kod jądra - odwraca kolejność elementów w tablicy
__global__ void reverseArray(int *d_in)
{
	// fragmenty tablicy w pamięci współdzielonej
	__shared__ int s_data1[N_THR];
	__shared__ int s_data2[N_THR];

	// indeksy elementów tablicy
	int in1 = blockDim.x * blockIdx.x + threadIdx.x;
	int in2 = N - blockDim.x * (blockIdx.x + 1) + threadIdx.x;

	// indeks w tablicy w pamięci współdzielonej
	int out = blockDim.x - 1 - threadIdx.x;

	// zapisanie elementów tablicy w pamięci współdzielonej
	s_data1[out] = d_in[in1];
	s_data2[out] = d_in[in2];

	// synchronizacja wątków w bloku
	__syncthreads();

	// zamiana elementów - przepisanie z pamięci współdzielonej
	// do pamięci globalnej
	d_in[in1] = s_data2[threadIdx.x];
	d_in[in2] = s_data1[threadIdx.x];


}

__global__ void findClosest(int *d_x ,int *d_y, int *d_out)
{
	// fragmenty tablicy w pamięci współdzielonej
	__shared__ int s_data1[N_THR];
	__shared__ int s_data2[N_THR];

	// indeksy elementów tablicy
	int in1 = blockDim.x * blockIdx.x + threadIdx.x;

	// indeks w tablicy w pamięci współdzielonej
	int x = d_x[in1];
	int y = d_y[in1];


	// zapisanie elementów tablicy w pamięci współdzielonej
	s_data1[out] = d_in[in1];
	s_data2[out] = d_in[in2];

	// synchronizacja wątków w bloku
	__syncthreads();
	int min_value = INT_MAX;
	int index;
	for (int i = 0; i < N; i++) {

		if (i != in1) {
			//func
		}
	}
	// zamiana elementów - przepisanie z pamięci współdzielonej
	// do pamięci globalnej
	d_in[in1] = s_data2[threadIdx.x];
	d_in[in2] = s_data1[threadIdx.x];
}

//////////////////////////////////////////////////////////////////
int main(int argc, char* argv[])
{
	// wybór karty na której zostana wykonane obliczenia
	if (cudaSetDevice(0) != cudaSuccess)
	{
		printf("cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");

		return 0;
	}

	// utworzenie tablicy
	int *a_cpu = (int*)malloc(N * sizeof(int));
	if (!a_cpu)
	{
		printf("malloc failed!\n");

		return 0;
	}
	for (int i = 0; i < N; ++i)
		a_cpu[i] = i;

	// przydzielenie pamięci na karcie
	int *a_gpu;
	if (cudaMalloc((void**)&a_gpu, N * sizeof(int)) != cudaSuccess)
	{
		printf("cudaMalloc failed!\n");

		return 0;
	}

	// przekopiowanie tablicy do pamięci karty
	if (cudaMemcpy(a_gpu, a_cpu, N * sizeof(int), cudaMemcpyHostToDevice) != cudaSuccess)
	{
		printf("cudaMemcpy failed!\n");

		return 0;
	}

	// rozmiar bloku i siatki bloków
	dim3 dimBlock(N_THR);
	dim3 dimGrid((N / 2) / N_THR);

	printf("N: %d, dimBlock: %d, dimGrid: %d\n\n", N, dimBlock.x, dimGrid.x);

	// utworzenie zdarzeń do pomiaru czasu obliczeń
	cudaEvent_t start, stop;
	if (cudaEventCreate(&start) != cudaSuccess)
	{
		printf("cudaEventCreate failed!\n");

		return 0;
	}
	if (cudaEventCreate(&stop) != cudaSuccess)
	{
		printf("cudaEventCreate failed!\n");

		return 0;
	}

	// zarejestrowanie zdarzenia - początek obliczeń
	if (cudaEventRecord(start, 0) != cudaSuccess)
	{
		printf("cudaEventRecord failed!\n");

		return 0;
	}

	// wywołanie kodu jądra
	reverseArray << < dimGrid, dimBlock >> >(a_gpu);

	// sprawdzenie czy udało się wywołać kod jądra
	if (cudaGetLastError() != cudaSuccess)
	{
		printf("kernel invocation failed!\n");

		return 0;
	}

	// zarejestrowanie zdarzenia - koniec obliczeń
	if (cudaEventRecord(stop, 0) != cudaSuccess)
	{
		printf("cudaEventRecord failed!\n");

		return 0;
	}

	// synchronizacja (oczekiwanie na zakończenie zdarzenia stop)
	if (cudaEventSynchronize(stop) != cudaSuccess)
	{
		printf("cudaEventSynchronize failed!\n");

		return 0;
	}

	// wyznaczenie czasu obliczeń
	float ms = 0;
	if (cudaEventElapsedTime(&ms, start, stop) != cudaSuccess)
	{
		printf("cudaEventElapsedTime failed!\n");

		return 0;
	}

	printf("running time on GPU:  %20f [ms]\n", ms);

	// przekopiowanie tablicy z pamięci karty
	if (cudaMemcpy(a_cpu, a_gpu, N * sizeof(int), cudaMemcpyDeviceToHost) != cudaSuccess)
	{
		printf("cudaMemcpy failed!\n");

		return 0;
	}

	// sprawdzenie poprawności zamiany
	for (int i = 0; i < N; i++)
	{
		if (a_cpu[i] != N - 1 - i)
		{
			printf("error!\n");

			return 0;
		}
	}

	int n2 = N / 2;
	int n1 = N - 1;
	int j;

	DWORD dw1 = GetTickCount();

	// zamiana elementów tablicy przez CPU
	for (int i = 0; i < n2; i++)
	{
		j = a_cpu[i];
		a_cpu[i] = a_cpu[n1 - i];
		a_cpu[n1 - i] = j;
	}

	DWORD dw2 = GetTickCount();

	printf("running time on CPU:  %20u [ms]\n", dw2 - dw1);

	// usunięcie zdarzeń
	if (cudaEventDestroy(start) != cudaSuccess)
	{
		printf("cudaEventDestroy failed!\n");

		return 0;
	}
	if (cudaEventDestroy(stop) != cudaSuccess)
	{
		printf("cudaEventDestroy failed!\n");

		return 0;
	}

	// zwolnienie pamięci
	cudaFree(a_gpu);
	free(a_cpu);

	// usunięcie kontekstu GPU
	cudaDeviceReset();
	int testInteger;
	printf("Press a key: ");
	scanf("%d", &testInteger);

	return 0;
}