Advertisement
Not a member of Pastebin yet?
Sign Up,
it unlocks many cool features!
- // ***********************************************************************
- //
- // Demo program pro vyuku predmetu APPS (10/2021)
- // Petr Olivka, katedra informatiky, FEI, VSB-TU Ostrava
- // email:petr.olivka@vsb.cz
- //
- // Priklad pouziti CUDA technologie.
- // Prevod barevneho obrazku na odstiny sede barvy
- //
- // ***********************************************************************
- #include <cuda_runtime.h>
- #include <stdio.h>
- // Demo kernel pro prevod barevneho bodu na odstin sede.
- __global__ void kernel_grayscale( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
- {
- // souradnice vlakna, kontrola rozmeru obrazku
- 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 ];
- // 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;
- // ulozeni bodu do obrazku
- bw_pic[ y * sizex + x ] = bgr;
- }
- __global__ void kernel_otoc( uchar4 *color_pic, uchar4* newimg, int sizex, int sizey )
- {
- // souradnice vlakna, kontrola rozmeru obrazku
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- if ( y >= sizey ) return;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( x > (sizex/2) ) return;
- int xDruhaPulka = ((sizex/2)-x)+(sizex/2);
- newimg[ y * sizex + x ] =color_pic[ y * sizex +xDruhaPulka ];
- newimg[ y * sizex + xDruhaPulka ] =color_pic[ y * sizex +x ];
- /*if(sizex%2 != 0)
- {
- if(x = sizex/2)
- newimg[ y * x] =color_pic[ y * +x ];
- }*/
- }
- __global__ void kernel_vyrez(uchar4 *in, uchar4* out, int2 size, int4 vyrez)
- {
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( y >= size.y) return;
- if ( x >= size.x) return;
- out[size.x * y +x] = in[0];
- if ( y >= vyrez.z || y <= vyrez.y) return;
- if ( x >= vyrez.w || x <= vyrez.x) return;
- /* if ( y >= vyrez.z || y <= vyrez.y) return;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( x >= vyrez.w || x <= vyrez.x) return;
- if ( y >= size.y) return;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( x >= size.x) return;
- */
- int souradniceVyrezuX = x - vyrez.x;
- int souradniceVyrezuY = y - vyrez.y;
- int vyrezSizex = vyrez.w - vyrez.x;
- out[vyrezSizex * souradniceVyrezuY + souradniceVyrezuX] = in[size.x * y +x];
- //out[size.x * y +x] = in[size.x * y +x];
- }
- __global__ void kernel_roztrhni(uchar4 *orig, uchar4* newimg1, uchar4* newimg2, int sizex, int sizey)
- {
- int y = blockDim.y * blockIdx.y + threadIdx.y;
- if ( y >= sizey ) return;
- int x = blockDim.x * blockIdx.x + threadIdx.x;
- if ( x >= sizex) return;
- if(x > sizex/2)
- {
- newimg1[ y * (sizex/2) + x - (sizex/2) ] = orig[ y * sizex + x ];
- }
- else
- {
- newimg2[ y * (sizex/2) + x ] = orig[ y * sizex + x ];
- }
- }
- void run_grayscale( uchar4 *original, uchar4* output, int2 size, int4 vyrez)
- {
- cudaError_t cerr;
- uchar4 *cudaIN;
- uchar4 *cudaOUT;
- cerr = cudaMalloc( &cudaIN, size.x * size.y * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- cerr = cudaMalloc( &cudaOUT, (vyrez.w-vyrez.x) * (vyrez.z-vyrez.y) * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- /* cerr = cudaMalloc( &cudaOUT, size.x *size.y * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) ); */
- // prenos barevneho obrazku do videokarty
- cerr = cudaMemcpy( cudaIN, original, size.x * size.y * sizeof( uchar4 ), cudaMemcpyHostToDevice );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- int block = 16;
- dim3 blocks(size.x,size.y);
- dim3 threads( block, block );
- printf( "predpustenim %d %d\n", size.x, size.y);
- printf( "predpustenim %d %d, %d, %d\n", vyrez.x, vyrez.y, vyrez.w, vyrez.z);
- // vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
- kernel_vyrez<<< blocks, threads >>>( cudaIN, cudaOUT, size, vyrez);
- /* cerr = cudaMemcpy(output, cudaOUT, ( size.x )* (size.y) * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) ); */
- cerr = cudaMemcpy(output, cudaOUT, (vyrez.w-vyrez.x) * (vyrez.z-vyrez.y) * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- cudaFree( cudaIN );
- cudaFree( cudaOUT );
- // alokace pameti ve videokarte
- /*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 ) );
- // prenos barevneho obrazku do videokarty
- 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)/2)+1, ( sizey + block - 1 ) / block );
- dim3 threads( block, block );
- // vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
- kernel_otoc<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey );
- if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- // prenos dat z videokarty
- cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- */
- printf( "OTOC OK" );
- // ROZTRZENI
- /* uchar4 *cudaImg1;
- cerr = cudaMalloc( &cudaImg1, sizex* sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- uchar4 *cudaImg2;
- cerr = cudaMalloc( &cudaImg2, sizex* sizey * sizeof( uchar4 ) );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- printf( "alok OK\n" );
- block = 16;
- dim3 blocks2( (( sizex + block - 1 ) / block), ( sizey + block - 1 ) / block );
- dim3 threads2( block, block );
- printf( "predpustenim\n" );
- // vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
- kernel_roztrhni<<< blocks2, threads2 >>>( cudaColorPic, cudaImg1, cudaImg2, sizex, sizey );
- printf( "pospusteni\n" );
- cerr = cudaMemcpy( img1, cudaImg1, (sizex/2) * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- printf( "ffffffff\n" );
- cerr = cudaMemcpy( img2, cudaImg2, (sizex/2) * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
- if ( cerr != cudaSuccess )
- printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
- printf( "gggggggggg\n" );
- // uvolneni alokovane pameti ve videokarte
- cudaFree( cudaImg1 );
- cudaFree( cudaImg2 );
- printf( "hhhhhh\n" );
- cudaFree( cudaColorPic );
- cudaFree( cudaBWPic );*/
- }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement