API tools faq
paste
Guest User

zmenseniscuda

a guest
Dec 7th, 2017
366
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 13.25 KB | None | 0 0
raw download clone embed print report
  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:[email protected]
  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
Public Pastes
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!