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

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

// mnożenie macierzy
#include <windows.h>

// rozmiar tablicy
const int N = 32;

// rozmiar bloku wątków
const int BLOCK_SIZE = 16;

/////////////////////////////////////////////////////////////////////
// funkcja kodu jądra wykonująca mnożenie macierzy C = A*B
// korzysta wyłącznie z pamięci globalnej
__global__ void findNearest(int* A, int* B, int* C)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;

    int nearest_idx;
    float nearest_dist = FLT_MAX;
    int x = A[idx];
    int y = B[idx];

    for (int i = 0; i < N; i++) {
        if (i != idx) {
            float dist = 0;
            dist = sqrtf(powf(x - A[i], 2) + powf(y - B[i],2));
            if (dist < nearest_dist) {
                nearest_dist = dist;
                nearest_idx = i;
            }
        }
    }

    __syncthreads();

    C[idx] = nearest_idx;
}

/////////////////////////////////////////////////////////////////////
int main()
{
    // 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 macierzy A, B, C
    size_t size = N * sizeof(int);
    int *tabA_cpu = (int*)malloc(N * sizeof(int));
    int *tabB_cpu = (int*)malloc(N * sizeof(int));
    int *tabC_cpu = (int*)malloc(N * sizeof(int));

    if (!tabA_cpu || !tabB_cpu || !tabC_cpu)
    {
        printf("malloc failed!\n");

        return 0;
    }

    for (int i = 0; i < N; ++i) {
        tabA_cpu[i] = i;
        tabB_cpu[i] = i;
        tabC_cpu[i] = i;
    }

    int *tabA_gpu;
    int *tabB_gpu;
    int *tabC_gpu;

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

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

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

        return 0;
    }

    if (cudaMemcpy(tabA_gpu, tabA_cpu, size, cudaMemcpyHostToDevice) != cudaSuccess)
    {
        printf("cudaMemcpy failed!\n");

        return 0;
    }
    if (cudaMemcpy(tabB_gpu, tabB_cpu, size, cudaMemcpyHostToDevice) != cudaSuccess)
    {
        printf("cudaMemcpy failed!\n");

        return 0;
    }
    if (cudaMemcpy(tabC_gpu, tabC_cpu, size, cudaMemcpyHostToDevice) != cudaSuccess)
    {
        printf("cudaMemcpy failed!\n");

        return 0;
    }

    // wywołanie funkcji kodu jądra; wykonywana przez dwuwymiarowe bloki wątków;
    // bloki zgrupowane są w dwuwymiarową siatkę
    dim3 dimBlock(BLOCK_SIZE);
    dim3 dimGrid(N / BLOCK_SIZE);

    // utworzenie zdarzeń do pomiaru czasu obliczeń
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    // zarejestrowanie zdarzenia - początek obliczeń
    cudaEventRecord(start, 0);

    // wywołanie odpowiedniego kodu jądra

    findNearest << <dimGrid, dimBlock >> >(tabA_gpu, tabB_gpu, tabC_gpu);

    // zarejestrowanie zdarzenia - koniec obliczeń
    cudaEventRecord(stop, 0);

    // synchronizacja (oczekiwanie na zakończenie zdarzenia stop)
    cudaEventSynchronize(stop);

    // wyznaczenie czasu obliczeń
    float ms = 0;
    cudaEventElapsedTime(&ms, start, stop);
    printf("running time on GPU:  %20f [ms]\n", ms);

    // usunięcie zdarzeń
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    // odczytanie wynikowej macierzy C z pamięci karty
    if (cudaMemcpy(tabC_cpu, tabC_gpu, size, cudaMemcpyDeviceToHost) != cudaSuccess)
    {
        printf("cudaMemcpy failed!\n");

        return false;
    }

    // zwolnienie pamięci przydzielonej na karcie
    cudaFree(tabA_gpu);
    cudaFree(tabB_gpu);
    cudaFree(tabC_gpu);

    cudaDeviceReset();

    for (int i = 0; i < N; i++) {
        printf(i + " nearest point id: " + tabC_cpu[i]);
        printf("\n");
    }

    return 0;
}