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

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <windows.h>
#include <cmath>
#include <limits>
#include <float.h>
#include <cstdlib>
#include <ctime>

// 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];
}

/////////////////////////////////////////////////////////////////////
// szuka najblizszego punktu
__global__ int GetClosest(int *x, int *y/*, int index, int n*/)
{
	__shared__ int s_data1[N_THR];

	int a = x[index];
	int b = y[index];
	double minDistance = DBL_MAX;
	double tmp;
	int minIndex;

	for (int i = 0; i < n; i++) {
		if (i != index) {
			tmp = sqrt(pow(x[i] - a, 2) + pow(y[i] - b, 2));
			if (tmp < minDistance) {
				minDistance = tmp;
				minIndex = i;
			}
		}
	}
	return minIndex;
}

//////////////////////////////////////////////////////////////////
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;
	}


	//**********************************

	int *X = (int*)malloc(N * sizeof(int));
	int *Y = (int*)malloc(N * sizeof(int));

	for (int i = 0; i < N; i++) {
		X[i] = rand() % 100;
		Y[i] = rand() % 100;
	}

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

		return 0;
	}
	if (cudaMalloc((void**)&Y_gpu, N * sizeof(int)) != cudaSuccess)
	{
		printf("cudaMalloc failed!\n");

		return 0;
	}

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

		return 0;
	}
	if (cudaMemcpy(Y_gpu, Y, 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
	GetClosest <<< dimGrid, dimBlock >> >(X_gpu,Y_gpu,N);

	// 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();

	return 0;
}