daily pastebin goal
64%
SHARE
TWEET

zmenseniscuda

a guest Dec 7th, 2017 63 Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. // ***********************************************************************
  2. //
  3. // Demo program for education in subject
  4. // Computer Architectures and Paralel Systems.
  5. // Petr Olivka, dep. of Computer Science, FEI, VSB-TU Ostrava
  6. // email:petr.olivka@vsb.cz
  7. //
  8. // Example of CUDA Technology Usage.
  9. // Image transformation from RGB to BW schema.
  10. //
  11. // ***********************************************************************
  12.  
  13. #include <cuda_runtime.h>
  14. #include <device_launch_parameters.h>
  15. #include <stdio.h>
  16.  
  17.  
  18. // RESIZE + MRIZKA
  19.  
  20. __global__ void kernel_resize(uchar4 *first, uchar4 *final, int size2x, int size2y)
  21. //( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
  22.  
  23. {
  24.         // souradnice vlakna - zjistujeme vlakna podobne jako pixely
  25.         int y = blockDim.y * blockIdx.y + threadIdx.y;
  26.         int x = blockDim.x * blockIdx.x + threadIdx.x;
  27.         if (x >= size2x) return;
  28.         if (y >= size2y) return;
  29.        
  30.  
  31.         uchar4 bgr = first[(y * 4 * size2x) + x * 2] ;
  32.    
  33.  
  34.     // vykreslí mřížku
  35.     // sloupce
  36.     for (int i = 0; i<x; i++)
  37.     {
  38.         //co 20pix nebo o 1 vedle
  39.         // kdyby chtel co 40 tak to bude x%40 = 0 || x%40= 39
  40.         if (x % 20 == 0 || x % 20 ==1)
  41.         {
  42.  
  43.             bgr.x = bgr.y = bgr.z = 250;
  44.            
  45.         }
  46.     }
  47.    
  48.     // radky
  49.     for (int j = 0; j<y; j++)
  50.     {
  51.         if (y % 20 == 0 || y % 20 ==1)
  52.         {
  53.  
  54.             bgr.x = bgr.y = bgr.z = 250;
  55.            
  56.         }
  57.     }
  58.  
  59.    
  60.     // vsechny tri barevne slozky budou mit stejnou hodnotu
  61.     //bgr.x = bgr.y = bgr.z = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
  62.  
  63.     //uchar4 bgr = first[(y * 4 * size2x) + x * 2] ;
  64.    
  65.     final [ y * size2x + x ] = bgr;
  66.  
  67.     //final [y * size2x + x] = first[(y * 4 * size2x) + x * 2] = bgr;
  68. }
  69.  
  70. void joinImage(uchar4 *original, uchar4 *rotated, int width, int height, int width2, int height2)
  71. {
  72.         cudaError_t cerr;
  73.  
  74.         // alokace pameti ve videokarte
  75.         uchar4 *cudaOriginal;
  76.         uchar4 *cudaRotate;
  77.         cerr = cudaMalloc(&cudaOriginal, width * height * sizeof(uchar4));                      // alokuje, pozdeji zjistuje jestli je vse ok
  78.         if (cerr != cudaSuccess)
  79.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  80.  
  81.         cerr = cudaMalloc(&cudaRotate, width2 * height2 * sizeof(uchar4));
  82.         if (cerr != cudaSuccess)
  83.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  84.  
  85.         // prenos barevneho obrazku do videokarty
  86.         cerr = cudaMemcpy(cudaOriginal, original, width * height * sizeof(uchar4), cudaMemcpyHostToDevice);
  87.         if (cerr != cudaSuccess)
  88.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  89.  
  90.         int block = 16;                                                                         // 16 pixelu
  91.         dim3 blocks((width2 + block - 1) / block, (height2 + block - 1) / block);               // ma 3 parametry, ale resime jen 2D - proto -1
  92.         dim3 threads(block, block);
  93.  
  94.         // vytvoreni bloku s vlakny, matice vlaken muze byt vetsi, nez samotny obrazek!
  95.         kernel_resize << < blocks, threads >> >(cudaOriginal, cudaRotate, width2, height2);
  96.  
  97.         if ((cerr = cudaGetLastError()) != cudaSuccess)
  98.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  99.  
  100.         // prenos dat z videokarty
  101.         cerr = cudaMemcpy(rotated, cudaRotate, width2 * height2 * sizeof(uchar4), cudaMemcpyDeviceToHost);
  102.         if (cerr != cudaSuccess)
  103.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  104.  
  105.         // uvolneni alokovane pameti ve videokarte
  106.         cudaFree(cudaRotate);
  107.         cudaFree(cudaOriginal);
  108.  
  109. }
  110.  
  111.  
  112.  
  113. /*
  114. // BLACK AND WHITE
  115.  
  116. // Demo kernel to tranfrom RGB color schema to BW schema
  117. __global__ void kernel_grayscale( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey, char barva )
  118. {
  119.         // X,Y coordinates and check image dimensions
  120.         int y = blockDim.y * blockIdx.y + threadIdx.y;
  121.         if ( y >= sizey ) return;
  122.         int x = blockDim.x * blockIdx.x + threadIdx.x;
  123.         if ( x >= sizex ) return;
  124.  
  125.         // Get point from color picture
  126.         uchar4 bgr = color_pic[ y * sizex + x ];
  127.  
  128.         // All R-G-B channels will have the same value
  129.         switch (barva){
  130.         case'B': bgr.y = bgr.z = 0;
  131.                 break;
  132.         case 'G': bgr.x = bgr.y = bgr.z = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
  133.                 break;
  134.         }
  135.        
  136.         // Store BW point to new image
  137.         bw_pic[ y * sizex + x ] = bgr;
  138. }
  139.  
  140. void run_grayscale( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey, char barva )
  141. {
  142.         cudaError_t cerr;
  143.         // Memory allocation in GPU device
  144.         uchar4 *cudaColorPic;
  145.         uchar4 *cudaBWPic;
  146.         cerr = cudaMalloc( &cudaColorPic, sizex * sizey * sizeof( uchar4 ) );
  147.         if ( cerr != cudaSuccess )
  148.                 printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );    
  149.  
  150.         cerr = cudaMalloc( &cudaBWPic, sizex * sizey * sizeof( uchar4 ) );
  151.         if ( cerr != cudaSuccess )
  152.                 printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );    
  153.  
  154.         // Copy color image to GPU device
  155.         cerr = cudaMemcpy( cudaColorPic, color_pic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyHostToDevice );
  156.         if ( cerr != cudaSuccess )
  157.                 printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );    
  158.  
  159.         int block = 16;
  160.         dim3 blocks( ( sizex + block - 1 ) / block, ( sizey + block - 1 ) / block );
  161.         dim3 threads( block, block );
  162.  
  163.         // Grid creation, size of grid must be greater than image
  164.         kernel_grayscale<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey, barva );
  165.  
  166.         if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
  167.                 printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  168.  
  169.         // Copy new image from GPU device
  170.         cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
  171.         if ( cerr != cudaSuccess )
  172.                 printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );    
  173.  
  174.         // Free memory
  175.         cudaFree( cudaColorPic );
  176.         cudaFree( cudaBWPic );
  177. }
  178. */
  179.  
  180.  
  181.  
  182. // ROTACE
  183.  
  184. __global__ void kernel_rotate(uchar4 *color_pic, uchar4 *rotated_pic, int sizex, int sizey)
  185. {
  186.         //Souradnice vlaken, kontrola
  187.         int y = blockDim.y * blockIdx.y + threadIdx.y;
  188.         int x = blockDim.x * blockIdx.x + threadIdx.x;
  189.         if (x >= sizex) return;
  190.         if (y >= sizey) return;
  191.  
  192.         // Upside down
  193.         //rotated_pic[y * sizex + x] = color_pic[(sizey - y - 1) * sizex + x];
  194.  
  195.  
  196.         rotated_pic[(y * sizex) + x] = color_pic[(y * sizex) + (sizex - x - 1)];
  197.        
  198.  
  199. }
  200.  
  201. void rotate_image(uchar4 *color_pic, uchar4 *rotated_pic, int sizex, int sizey)
  202. {
  203.         cudaError_t cerr;
  204.         // Memory allocation in GPU device
  205.         uchar4 *cudaColorPic;
  206.         uchar4 *cudaRotatedPic;
  207.  
  208.         cerr = cudaMalloc(&cudaColorPic, sizex * sizey * sizeof(uchar4));
  209.         if (cerr != cudaSuccess)
  210.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  211.  
  212.         cerr = cudaMalloc(&cudaRotatedPic, sizex * sizey * sizeof(uchar4));
  213.         if (cerr != cudaSuccess)
  214.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  215.  
  216.         // Copy color image to GPU device
  217.         cerr = cudaMemcpy(cudaColorPic, color_pic, sizex * sizey * sizeof(uchar4), cudaMemcpyHostToDevice);
  218.         if (cerr != cudaSuccess)
  219.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  220.  
  221.         int block = 16; //VELIKOST BLOKU
  222.         dim3 blocks((sizex + block - 1) / block, (sizey + block - 1) / block); //kolik blocku je v gridu
  223.         dim3 threads(block, block); //kolik threadu je v blocku
  224.  
  225.         // Grid creation, size of grid must be greater than image
  226.         kernel_rotate<<< blocks, threads >>>(cudaColorPic, cudaRotatedPic, sizex, sizey);
  227.  
  228.         if ((cerr = cudaGetLastError()) != cudaSuccess)
  229.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  230.  
  231.         // Copy new image from GPU device
  232.         cerr = cudaMemcpy(rotated_pic, cudaRotatedPic, sizex * sizey * sizeof(uchar4), cudaMemcpyDeviceToHost);
  233.         if (cerr != cudaSuccess)
  234.                 printf("CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString(cerr));
  235.  
  236.         // Free memory
  237.         cudaFree(cudaColorPic);
  238.         cudaFree(cudaRotatedPic);
  239. }
  240.  
  241.  
  242. __global__ void kernel_flip( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
  243. {
  244.     // X,Y coordinates and check image dimensions
  245.     int y = blockDim.y * blockIdx.y + threadIdx.y;
  246.     if ( y >= sizey ) return;
  247.     int x = blockDim.x * blockIdx.x + threadIdx.x;
  248.     if ( x >= sizex ) return;
  249.  
  250.     //180
  251.     bw_pic[y * sizex + x] = color_pic[(sizey - y) * sizex + x];
  252.  
  253.     //zrcadleni
  254.     //bw_pic[ y * sizex + x ] = color_pic[ y * sizex + sizex - x ];
  255.  
  256. }
  257.  
  258.  
  259. void run_flip( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
  260. {
  261.     cudaError_t cerr;
  262.     // Memory allocation in GPU device
  263.     uchar4 *cudaColorPic;
  264.     uchar4 *cudaBWPic;
  265.     cerr = cudaMalloc( &cudaColorPic, sizex * sizey * sizeof( uchar4 ) );
  266.     if ( cerr != cudaSuccess )
  267.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  268.  
  269.     cerr = cudaMalloc( &cudaBWPic, sizex * sizey * sizeof( uchar4 ) );
  270.     if ( cerr != cudaSuccess )
  271.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  272.  
  273.     // Copy color image to GPU device
  274.     cerr = cudaMemcpy( cudaColorPic, color_pic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyHostToDevice );
  275.     if ( cerr != cudaSuccess )
  276.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  277.  
  278.     int block = 16;
  279.     dim3 blocks( ( sizex + block - 1 ) / block, ( sizey + block - 1 ) / block );
  280.     dim3 threads( block, block );
  281.  
  282.     // Grid creation, size of grid must be greater than image
  283.     kernel_flip<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey );
  284.  
  285.     if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
  286.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  287.  
  288.     // Copy new image from GPU device
  289.     cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
  290.     if ( cerr != cudaSuccess )
  291.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  292.  
  293.     // Free memory
  294.     cudaFree( cudaColorPic );
  295.     cudaFree( cudaBWPic );
  296. }
  297.  
  298.  
  299.  
  300. __global__ void kernel_mriz( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
  301. {
  302.     // X,Y coordinates and check image dimensions
  303.     int y = blockDim.y * blockIdx.y + threadIdx.y;
  304.     if ( y >= sizey ) return;
  305.     int x = blockDim.x * blockIdx.x + threadIdx.x;
  306.     if ( x >= sizex ) return;
  307.  
  308.     uchar4 bgr = color_pic[ y * sizex + x ];
  309.    
  310.     // vykresli se jen obsah krize
  311.     /*
  312.     if (y < 100)
  313.     {
  314.         if(x < 180 || x > 220)
  315.             bgr.x = bgr.y = bgr.z = 100;
  316.     }
  317.     else if (y > 140)
  318.     {
  319.         if(x < 180 || x > 220)
  320.             bgr.x = bgr.y = bgr.z = 100;
  321.     }
  322.  
  323.     bw_pic[ y * sizex + x ] = bgr;*/
  324.  
  325.  
  326.     // nakresli kriz
  327.     /*
  328.     if (x > 100 && x < 280)
  329.         {
  330.             if (y > 148 && y < 150)
  331.             {
  332.                 bgr.x = bgr.y = bgr.z =250;
  333.             }
  334.         }
  335.    
  336.         if (y > 50 && y < 450)
  337.         {
  338.             if (x > 180 && x < 200)
  339.             {
  340.              bgr.x = bgr.y = bgr.z = 250;
  341.             }
  342.         }*/
  343.  
  344.  
  345.     // vykreslí mřížku
  346.     // sloupce
  347.     for (int i = 0; i<x; i++)
  348.     {
  349.         //co 20pix nebo o 1 vedle
  350.         // kdyby chtel co 40 tak to bude x%40 = 0 || x%40= 39
  351.         if (x % 20 == 0 || x % 20 ==1)
  352.         {
  353.  
  354.             bgr.x = bgr.y = bgr.z =250;
  355.            
  356.         }
  357.     }
  358.    
  359.     // radky
  360.     for (int j = 0; j<y; j++)
  361.     {
  362.         if (y % 20 == 0 || y % 20 ==1)
  363.         {
  364.  
  365.             bgr.x = bgr.y = bgr.z =250;
  366.            
  367.         }
  368.     }
  369.  
  370.     // vsechny tri barevne slozky budou mit stejnou hodnotu
  371.     //bgr.x = bgr.y = bgr.z = bgr.x * 0.11 + bgr.y * 0.59 + bgr.z * 0.30;
  372.  
  373.     bw_pic[ y * sizex + x ] = bgr;
  374. }
  375.  
  376. void run_mriz( uchar4 *color_pic, uchar4* bw_pic, int sizex, int sizey )
  377. {
  378.     cudaError_t cerr;
  379.     // Memory allocation in GPU device
  380.     uchar4 *cudaColorPic;
  381.     uchar4 *cudaBWPic;
  382.     cerr = cudaMalloc( &cudaColorPic, sizex * sizey * sizeof( uchar4 ) );
  383.     if ( cerr != cudaSuccess )
  384.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  385.  
  386.     cerr = cudaMalloc( &cudaBWPic, sizex * sizey * sizeof( uchar4 ) );
  387.     if ( cerr != cudaSuccess )
  388.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  389.  
  390.     // Copy color image to GPU device
  391.     cerr = cudaMemcpy( cudaColorPic, color_pic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyHostToDevice );
  392.     if ( cerr != cudaSuccess )
  393.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  394.  
  395.     int block = 16;
  396.     dim3 blocks( ( sizex + block - 1 ) / block, ( sizey + block - 1 ) / block );
  397.     dim3 threads( block, block );
  398.  
  399.     // Grid creation, size of grid must be greater than image
  400.     kernel_mriz<<< blocks, threads >>>( cudaColorPic, cudaBWPic, sizex, sizey );
  401.  
  402.     if ( ( cerr = cudaGetLastError() ) != cudaSuccess )
  403.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  404.  
  405.     // Copy new image from GPU device
  406.     cerr = cudaMemcpy( bw_pic, cudaBWPic, sizex * sizey * sizeof( uchar4 ), cudaMemcpyDeviceToHost );
  407.     if ( cerr != cudaSuccess )
  408.         printf( "CUDA Error [%d] - '%s'\n", __LINE__, cudaGetErrorString( cerr ) );
  409.  
  410.     // Free memory
  411.     cudaFree( cudaColorPic );
  412.     cudaFree( cudaBWPic );
  413. }
RAW Paste Data
Top