Advertisement
Not a member of Pastebin yet?
Sign Up,
it unlocks many cool features!
- // ***********************************************************************
- //
- // 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.
- // Image transformation from RGB to BW schema.
- //
- // ***********************************************************************
- #include <cuda_runtime.h>
- #include <device_launch_parameters.h>
- #include <stdio.h>
- // RESIZE + MRIZKA
- __global__ void kernel_resize(uchar4 *first, uchar4 *final, int size2x, int size2y)
- //( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
- {
- // souradnice vlakna - zjistujeme vlakna podobne jako pixely
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if (x >= size2x) return;
- if (y >= size2y) return;
- uchar4 bgr = first[(y * 4 * size2x) + x * 2] ;
- // vykreslí mřížku
- // sloupce
- for (int i = 0; i<x; i++)
- {
- //co 20pix nebo o 1 vedle
- // kdyby chtel co 40 tak to bude x%40 = 0 || x%40= 39
- if (x % 20 == 0 || x % 20 ==1)
- {
- bgr.x = bgr.y = bgr.z = 250;
- }
- }
- // radky
- for (int j = 0; j<y; j++)
- {
- if (y % 20 == 0 || y % 20 ==1)
- {
- bgr.x = bgr.y = bgr.z = 250;
- }
- }
- // vsechny tri barevne slozky budou mit stejnou hodnotu
- //bgr.x = bgr.y = bgr.z = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
- //uchar4 bgr = first[(y * 4 * size2x) + x * 2] ;
- final [ y * size2x + x ] = bgr;
- //final [y * size2x + x] = first[(y * 4 * size2x) + x * 2] = bgr;
- }
- void joinImage(uchar4 *original, uchar4 *rotated, int width, int height, int width2, int height2)
- {
- cudaError_t cerr;
- // alokace pameti ve videokarte
- uchar4 *cudaOriginal;
- uchar4 *cudaRotate;
- cerr = cudaMalloc(&cudaOriginal, width * height * sizeof(uchar4)); // alokuje, pozdeji zjistuje jestli je vse ok
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- cerr = cudaMalloc(&cudaRotate, width2 * height2 * sizeof(uchar4));
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- // prenos barevneho obrazku do videokarty
- cerr = cudaMemcpy(cudaOriginal, original, width * height * sizeof(uchar4), cudaMemcpyHostToDevice);
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- int block = 16; // 16 pixelu
- dim3 blocks((width2 + block - 1) / block, (height2 + block - 1) / block); // ma 3 parametry, ale resime jen 2D - proto -1
- dim3 threads(block, block);
- // vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
- kernel_resize << < blocks, threads >> >(cudaOriginal, cudaRotate, width2, height2);
- if ((cerr = cudaGetLastError()) != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- // prenos dat z videokarty
- cerr = cudaMemcpy(rotated, cudaRotate, width2 * height2 * sizeof(uchar4), cudaMemcpyDeviceToHost);
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- // uvolneni alokovane pameti ve videokarte
- cudaFree(cudaRotate);
- cudaFree(cudaOriginal);
- }
- /*
- // BLACK AND WHITE
- // Demo kernel to tranfrom RGB color schema to BW schema
- __global__ void kernel_grayscale( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey, char barva )
- {
- // X,Y coordinates and check image dimensions
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- if ( y >= sizey ) return;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( x >= sizex ) return;
- // Get point from color picture
- uchar4 bgr = color_pic[ y * sizex + x ];
- // All R-G-B channels will have the same value
- switch (barva){
- case'B': bgr.y = bgr.z = 0;
- break;
- case 'G': bgr.x = bgr.y = bgr.z = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
- break;
- }
- // Store BW point to new image
- bw_pic[ y * sizex + x ] = bgr;
- }
- void run_grayscale( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey, char barva )
- {
- cudaError_t cerr;
- // Memory allocation in GPU device
- uchar4 *cudaColorPic;
- uchar4 *cudaBWPic;
- cerr = cudaMalloc( &cudaColorPic, sizex * sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- cerr = cudaMalloc( &cudaBWPic, sizex * sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Copy color image to GPU device
- cerr = cudaMemcpy( cudaColorPic, color_pic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyHostToDevice );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- int block = 16;
- dim3 blocks( ( sizex + block - 1 ) / block, ( sizey + block - 1 ) / block );
- dim3 threads( block, block );
- // Grid creation, size of grid must be greater than image
- kernel_grayscale<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey, barva );
- if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Copy new image from GPU device
- cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Free memory
- cudaFree( cudaColorPic );
- cudaFree( cudaBWPic );
- }
- */
- // ROTACE
- __global__ void kernel_rotate(uchar4 *color_pic, uchar4 *rotated_pic, int sizex, int sizey)
- {
- //Souradnice vlaken, kontrola
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if (x >= sizex) return;
- if (y >= sizey) return;
- // Upside down
- //rotated_pic[y * sizex + x] = color_pic[(sizey - y - 1) * sizex + x];
- rotated_pic[(y * sizex) + x] = color_pic[(y * sizex) + (sizex - x - 1)];
- }
- void rotate_image(uchar4 *color_pic, uchar4 *rotated_pic, int sizex, int sizey)
- {
- cudaError_t cerr;
- // Memory allocation in GPU device
- uchar4 *cudaColorPic;
- uchar4 *cudaRotatedPic;
- cerr = cudaMalloc(&cudaColorPic, sizex * sizey * sizeof(uchar4));
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- cerr = cudaMalloc(&cudaRotatedPic, sizex * sizey * sizeof(uchar4));
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- // Copy color image to GPU device
- cerr = cudaMemcpy(cudaColorPic, color_pic, sizex * sizey * sizeof(uchar4), cudaMemcpyHostToDevice);
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- int block = 16; //VELIKOST BLOKU
- dim3 blocks((sizex + block - 1) / block, (sizey + block - 1) / block); //kolik blocku je v gridu
- dim3 threads(block, block); //kolik threadu je v blocku
- // Grid creation, size of grid must be greater than image
- kernel_rotate<<< blocks, threads >>>(cudaColorPic, cudaRotatedPic, sizex, sizey);
- if ((cerr = cudaGetLastError()) != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- // Copy new image from GPU device
- cerr = cudaMemcpy(rotated_pic, cudaRotatedPic, sizex * sizey * sizeof(uchar4), cudaMemcpyDeviceToHost);
- if (cerr != cudaSuccess)
- printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
- // Free memory
- cudaFree(cudaColorPic);
- cudaFree(cudaRotatedPic);
- }
- __global__ void kernel_flip( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
- {
- // X,Y coordinates and check image dimensions
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- if ( y >= sizey ) return;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( x >= sizex ) return;
- //180
- bw_pic[y * sizex + x] = color_pic[(sizey - y) * sizex + x];
- //zrcadleni
- //bw_pic[ y * sizex + x ] = color_pic[ y * sizex + sizex - x ];
- }
- void run_flip( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
- {
- cudaError_t cerr;
- // Memory allocation in GPU device
- uchar4 *cudaColorPic;
- uchar4 *cudaBWPic;
- cerr = cudaMalloc( &cudaColorPic, sizex * sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- cerr = cudaMalloc( &cudaBWPic, sizex * sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Copy color image to GPU device
- cerr = cudaMemcpy( cudaColorPic, color_pic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyHostToDevice );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- int block = 16;
- dim3 blocks( ( sizex + block - 1 ) / block, ( sizey + block - 1 ) / block );
- dim3 threads( block, block );
- // Grid creation, size of grid must be greater than image
- kernel_flip<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey );
- if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Copy new image from GPU device
- cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Free memory
- cudaFree( cudaColorPic );
- cudaFree( cudaBWPic );
- }
- __global__ void kernel_mriz( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
- {
- // X,Y coordinates and check image dimensions
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- if ( y >= sizey ) return;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( x >= sizex ) return;
- uchar4 bgr = color_pic[ y * sizex + x ];
- // vykresli se jen obsah krize
- /*
- if (y < 100)
- {
- if(x < 180 || x > 220)
- bgr.x = bgr.y = bgr.z = 100;
- }
- else if (y > 140)
- {
- if(x < 180 || x > 220)
- bgr.x = bgr.y = bgr.z = 100;
- }
- bw_pic[ y * sizex + x ] = bgr;*/
- // nakresli kriz
- /*
- if (x > 100 && x < 280)
- {
- if (y > 148 && y < 150)
- {
- bgr.x = bgr.y = bgr.z =250;
- }
- }
- if (y > 50 && y < 450)
- {
- if (x > 180 && x < 200)
- {
- bgr.x = bgr.y = bgr.z = 250;
- }
- }*/
- // vykreslí mřížku
- // sloupce
- for (int i = 0; i<x; i++)
- {
- //co 20pix nebo o 1 vedle
- // kdyby chtel co 40 tak to bude x%40 = 0 || x%40= 39
- if (x % 20 == 0 || x % 20 ==1)
- {
- bgr.x = bgr.y = bgr.z =250;
- }
- }
- // radky
- for (int j = 0; j<y; j++)
- {
- if (y % 20 == 0 || y % 20 ==1)
- {
- bgr.x = bgr.y = bgr.z =250;
- }
- }
- // vsechny tri barevne slozky budou mit stejnou hodnotu
- //bgr.x = bgr.y = bgr.z = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
- bw_pic[ y * sizex + x ] = bgr;
- }
- void run_mriz( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
- {
- cudaError_t cerr;
- // Memory allocation in GPU device
- uchar4 *cudaColorPic;
- uchar4 *cudaBWPic;
- cerr = cudaMalloc( &cudaColorPic, sizex * sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- cerr = cudaMalloc( &cudaBWPic, sizex * sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Copy color image to GPU device
- cerr = cudaMemcpy( cudaColorPic, color_pic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyHostToDevice );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- int block = 16;
- dim3 blocks( ( sizex + block - 1 ) / block, ( sizey + block - 1 ) / block );
- dim3 threads( block, block );
- // Grid creation, size of grid must be greater than image
- kernel_mriz<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey );
- if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Copy new image from GPU device
- cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // Free memory
- cudaFree( cudaColorPic );
- cudaFree( cudaBWPic );
- }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement