***********************************************************************
//
// Demo program for education in subject
// Computer Architectures and Paralel Systems.
// Petr Olivka, dep. of Computer Science, FEI, VSB-TU Ostrava
// email:petr.olivka@vsb.cz
//
// Example of CUDA Technology Usage without unified memory.
//
// Manipulation with prepared image.
//
// ***********************************************************************

#include <stdio.h>
#include <cuda_device_runtime_api.h>
#include <cuda_runtime.h>
#include \"pic_type.h\"
#include <cmath>

// Every threads identifies its position in grid and in block and modify image
__global__ void kernel_animation( CUDA_Pic cuda_pic )
{
    // X,Y coordinates
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    if ( x >= cuda_pic.Size.x ) return;
    if ( y >= cuda_pic.Size.y ) return;

    // Point [x,y] selection from image
    uchar3 bgr = cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ];

    if ((((x - cuda_pic.Size.x / 2) * (x - cuda_pic.Size.x / 2) + (y - cuda_pic.Size.y / 2) * (y - cuda_pic.Size.y / 2)) < 10000) && (((x - cuda_pic.Size.x / 2) * (x - cuda_pic.Size.x / 2) + (y - cuda_pic.Size.y / 2) * (y - cuda_pic.Size.y / 2)) > 9000))
    {
        bgr.x = 0;
        bgr.y = 0;
        bgr.z = 255;
        cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ] = bgr;
    }

    if (x > 125 && x < 130 && y > 90 && y < 210)
    {
        bgr.x = 0;
        bgr.y = 0;
        bgr.z = 0;
        cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ] = bgr;
    }

    if (x > 170 && x < 175 && y > 90 && y < 210)
    {
        bgr.x = 0;
        bgr.y = 0;
        bgr.z = 255;
        cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ] = bgr;
    }

    if (x + 79 <= y && x + 84 >= y  && x > 99 && x < 129)
    {
        bgr.x = 0;
        bgr.y = 0;
        bgr.z = 0;
        cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ] = bgr;
    }

    if (x - 32 >= abs(300 - y) && x - 37 <= abs(300 - y) && x > 129 && x < 159)
    {
        bgr.x = 0;
        bgr.y = 0;
        bgr.z = 0;
        cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ] = bgr;
    }

    if (x - 80 >= y && x - 85 <= y && x > 174 && x < 199)
    {
        bgr.x = 0;
        bgr.y = 0;
        bgr.z = 255;
        cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ] = bgr;
    }

    if (x + 33 <= abs(300 - y) && x + 38 >= abs(300 - y) && x > 146 && x < 174)
    {
        bgr.x = 0;
        bgr.y = 0;
        bgr.z = 255;
        cuda_pic.P_uchar3[ y * cuda_pic.Size.x + x ] = bgr;
    }
}

void cu_run_animation( CUDA_Pic pic, uint2 block_size )
{
    cudaError_t cerr;

    CUDA_Pic cudaPic;
    cudaPic.Size = pic.Size;

    // Memory allocation in GPU device
    cerr = cudaMalloc( &cudaPic.P_void, cudaPic.Size.x * cudaPic.Size.y * sizeof( uchar3 ) );
    if ( cerr != cudaSuccess )
        printf( \"CUDA Error [%d] - \'%s\'\\n\", __LINE__, cudaGetErrorString( cerr ) );

    // Copy data to GPU device
    cerr = cudaMemcpy( cudaPic.P_void, pic.P_void, cudaPic.Size.x * cudaPic.Size.y * sizeof( uchar3 ), cudaMemcpyHostToDevice );
    if ( cerr != cudaSuccess )