Advertisement
Guest User

zmenseniscuda

a guest
Dec 7th, 2017
346
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 13.25 KB | None | 0 0
  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. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement