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

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

#include <stdio.h>
#include <windows.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];
}

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

        return 0;
}