cuda_4

// ***********************************************************************
//
// Demo program for education in subject
// Computer Architectures and Parallel Systems.
// Petr Olivka, dep. of Computer Science, FEI, VSB-TU Ostrava
// email:petr.olivka@vsb.cz
//
// Example of CUDA Technology Usage with unified memory.
//
// Image transformation from RGB to BW schema.
//
// ***********************************************************************

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

// Demo kernel to transform RGB color schema to BW schema
__global__ void kernel_grayscale( CUDA_Pic colorPic, CUDA_Pic bwPic ) {
    // X,Y coordinates and check image dimensions
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    if ( y >= colorPic.Size.y ) return;
    if ( x >= colorPic.Size.x ) return;



    // Get point from color picture
    uchar3 bgr = colorPic.P_uchar3[ y * colorPic.Size.x + x ];

    // Store BW point to new image
    bwPic.P_uchar1[ y * bwPic.Size.x + x ].x = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
}

__global__ void kernel_resize( CUDA_Pic pic_orig, CUDA_Pic pic_small ) {
    // X,Y coordinates and check image dimensions
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    if ( y >= pic_orig.Size.y ) return;
    if ( x >= pic_orig.Size.x ) return;

    if(x % 2 == 0 && y % 2 == 0){
        //pic_orig.P_uchar3[ y * pic_orig.Size.x + x ] = ( uchar3 ) { 0, 255, 0 };

        uchar3 lh = pic_orig.P_uchar3[ y * pic_orig.Size.x + x ];
        uchar3 ph = pic_orig.P_uchar3[ y * pic_orig.Size.x + x + 1 ];
        uchar3 ld = pic_orig.P_uchar3[ (y + 1) * pic_orig.Size.x + x ];
        uchar3 pd = pic_orig.P_uchar3[ (y + 1) * pic_orig.Size.x + x + 1 ];

        uchar3 novy_pixel;
        novy_pixel.x = (lh.x + ph.x + ld.x + pd.x) / 4;
        novy_pixel.y = (lh.y + ph.y + ld.y + pd.y) / 4;
        novy_pixel.z = (lh.z + ph.z + ld.z + pd.z) / 4;


        int newPosition = (y / 2) * pic_small.Size.x + (x / 2);

        pic_small.P_uchar3[ newPosition ] = novy_pixel;
    }

    // Get point from color picture
    //uchar3 bgr = colorPic.P_uchar3[ y * colorPic.Size.x + x ];

    // Store BW point to new image
    //bwPic.P_uchar1[ y * bwPic.Size.x + x ].x = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
    //pic_small.P_uchar3[ newPosition ] = ( uchar3 ) { 0, 255, 0 };
    //pic_small.P_uchar3[ newPosition ] = pic_orig.P_uchar3[ y * pic_orig.Size.x + x ];
}

__global__ void kernel_blur( CUDA_Pic pic_orig, CUDA_Pic pic_blurred ) {
    // X,Y coordinates and check image dimensions
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    if ( y >= pic_orig.Size.y ) return;
    if ( x >= pic_orig.Size.x ) return;

    if(true){
        //pic_orig.P_uchar3[ y * pic_orig.Size.x + x ] = ( uchar3 ) { 0, 255, 0 };

        uchar3 lh = pic_orig.P_uchar3[ y * pic_orig.Size.x + x ];
        uchar3 ph = pic_orig.P_uchar3[ y * pic_orig.Size.x + x + 1 ];
        uchar3 ld = pic_orig.P_uchar3[ (y + 1) * pic_orig.Size.x + x ];
        uchar3 pd = pic_orig.P_uchar3[ (y + 1) * pic_orig.Size.x + x + 1 ];

        uchar3 novy_pixel;
        novy_pixel.x = (lh.x + ph.x + ld.x + pd.x) / 4;
        novy_pixel.y = (lh.y + ph.y + ld.y + pd.y) / 4;
        novy_pixel.z = (lh.z + ph.z + ld.z + pd.z) / 4;


        //int newPosition = (y / 2) * pic_small.Size.x + (x / 2);
        //pic_blurred.P_uchar3[ newPosition ] = novy_pixel;
        pic_blurred.P_uchar3[ y * pic_orig.Size.x + x ] = novy_pixel;
        //pic_blurred.P_uchar3[ y * pic_orig.Size.x + x + 1 ] = novy_pixel;
        //pic_blurred.P_uchar3[ (y + 1) * pic_orig.Size.x + x ] = novy_pixel;
        //pic_blurred.P_uchar3[ (y + 1) * pic_orig.Size.x + x + 1 ] = novy_pixel;


    }

    // Get point from color picture
    //uchar3 bgr = colorPic.P_uchar3[ y * colorPic.Size.x + x ];

    // Store BW point to new image
    //bwPic.P_uchar1[ y * bwPic.Size.x + x ].x = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
    //pic_small.P_uchar3[ newPosition ] = ( uchar3 ) { 0, 255, 0 };
    //pic_small.P_uchar3[ newPosition ] = pic_orig.P_uchar3[ y * pic_orig.Size.x + x ];
}

__global__ void kernel_blur2(CUDA_Pic org_pic, CUDA_Pic new_pic)
{
    // X,Y coordinates and check image dimensions
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    if (y == 0) return;
    if (x == 0) return;
    if (y+1 >= org_pic.Size.y) return;
    if (x+1 >= org_pic.Size.x) return;

    uchar3 bgr[9];
    bgr[0] = org_pic.P_uchar3[(y - 1) * org_pic.Size.x + x - 1];
    bgr[1] = org_pic.P_uchar3[(y - 1) * org_pic.Size.x + x];
    bgr[2] = org_pic.P_uchar3[(y - 1) * org_pic.Size.x + x + 1];

    bgr[3] = org_pic.P_uchar3[y * org_pic.Size.x + x - 1];
    bgr[4] = org_pic.P_uchar3[y * org_pic.Size.x + x];
    bgr[5] = org_pic.P_uchar3[y * org_pic.Size.x + x + 1];

    bgr[6] = org_pic.P_uchar3[(y + 1) * org_pic.Size.x + x - 1];
    bgr[7] = org_pic.P_uchar3[(y + 1) * org_pic.Size.x + x];
    bgr[8] = org_pic.P_uchar3[(y + 1) * org_pic.Size.x + x + 1];

    uint3 avg;
    avg.x = 0;
    avg.y = 0;
    avg.z = 0;

    for (int i = 0; i < 9; i++) {
        avg.x += bgr[i].x;
        avg.y += bgr[i].y;
        avg.z += bgr[i].z;
    }

    avg.x = avg.x / 9;
    avg.y = avg.y / 9;
    avg.z = avg.z / 9;

    uchar3 avgChar;
    avgChar.x = avg.x;
    avgChar.y = avg.y;
    avgChar.z = avg.z;

    new_pic.P_uchar3[(y)* new_pic.Size.x + (x)] = avgChar;
}



__global__ void kernel_rotate( CUDA_Pic pic_orig, CUDA_Pic pic_rotated, float angle ) {
    // X,Y coordinates and check image dimensions
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    if ( y >= pic_orig.Size.y ) return;
    if ( x >= pic_orig.Size.x ) return;

    int newx = x * cos(angle) + y * sin(angle);
    int newy = y * cos(angle) - x * sin(angle);
    newy = pic_rotated.Size.y + newy;

    int newPosition = (newy) * pic_rotated.Size.x + (newx);

    if(0 <= newPosition && newPosition < pic_rotated.Size.x * pic_rotated.Size.y )
        pic_rotated.P_uchar3[ newPosition ] =  pic_orig.P_uchar3[ y * pic_orig.Size.x + x ];
    //pic_rotated.P_uchar3[ newPosition ] = pic_orig.P_uchar3[ y * pic_orig.Size.x + x ];

    // Get point from color picture
    //uchar3 bgr = colorPic.P_uchar3[ y * colorPic.Size.x + x ];

    // Store BW point to new image
    //bwPic.P_uchar1[ y * bwPic.Size.x + x ].x = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
    //pic_small.P_uchar3[ newPosition ] = ( uchar3 ) { 0, 255, 0 };
    //pic_small.P_uchar3[ newPosition ] = pic_orig.P_uchar3[ y * pic_orig.Size.x + x ];
}






void cu_run_grayscale( CUDA_Pic colorPic, CUDA_Pic bwPic ) {
    cudaError_t cerr;

    // Grid creation, size of grid must be equal or greater than images
    int block_size = 16;
    dim3 blocks( ( colorPic.Size.x + block_size - 1 ) / block_size, ( colorPic.Size.y + block_size - 1 ) / block_size );
    dim3 threads( block_size, block_size );


    kernel_grayscale<<< blocks, threads >>>( colorPic, bwPic );

    if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
        printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

    cudaDeviceSynchronize();
}

void cu_run_resize( CUDA_Pic pic_orig, CUDA_Pic pic_small ) {
    cudaError_t cerr;

    // Grid creation, size of grid must be equal or greater than images
    int block_size = 16;
    dim3 blocks( ( pic_orig.Size.x + block_size - 1 ) / block_size, ( pic_orig.Size.y + block_size - 1 ) / block_size );
    dim3 threads( block_size, block_size );


    kernel_resize<<< blocks, threads >>>( pic_orig, pic_small );

    if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
        printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

    cudaDeviceSynchronize();
}

void cu_run_blur( CUDA_Pic pic_orig, CUDA_Pic pic_blurred ) {
    cudaError_t cerr;

    // Grid creation, size of grid must be equal or greater than images
    int block_size = 16;
    dim3 blocks( ( pic_orig.Size.x + block_size - 1 ) / block_size, ( pic_orig.Size.y + block_size - 1 ) / block_size );
    dim3 threads( block_size, block_size );


    kernel_blur<<< blocks, threads >>>( pic_orig, pic_blurred );

    if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
        printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

    cudaDeviceSynchronize();
}

void cu_run_rotate( CUDA_Pic pic_orig, CUDA_Pic pic_rotated, float angle ) {
    cudaError_t cerr;

    // Grid creation, size of grid must be equal or greater than images
    int block_size = 16;
    dim3 blocks( ( pic_orig.Size.x + block_size - 1 ) / block_size, ( pic_orig.Size.y + block_size - 1 ) / block_size );
    dim3 threads( block_size, block_size );


    kernel_rotate<<< blocks, threads >>>( pic_orig, pic_rotated, angle );

    if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
        printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );

    cudaDeviceSynchronize();

}