Difference between revisions of "CoE Cluster november 2011/CUDA"

From Teknologisk videncenter
Jump to: navigation, search
(Created page with "// Read image from file, remove blue channel and write onto another file #include <iostream> #include <string> #include <sstream> #include <cv.h> #include <highgui.h> using name...")
 
Line 1: Line 1:
// Read image from file, remove blue channel and write onto another file
+
<pre>// Read image from file, remove blue channel and write onto another file
  
#include <iostream>
+
include &lt;iostream&gt;
#include <string>
 
#include <sstream>
 
#include <cv.h>
 
#include <highgui.h>
 
using namespace cv;
 
using namespace std;
 
#include "PixelMat.hpp"
 
  
void startCUDA(int blocks, int threads, unsigned char* img, int rows, int cols, int step);
+
include &lt;string&gt;
 +
 
 +
include &lt;sstream&gt;
 +
 
 +
include &lt;cv.h&gt;
 +
 
 +
include &lt;highgui.h&gt;
  
int main(int argc, char* argv[])
+
using namespace cv; using namespace std;
{
 
  int blocks = 0;
 
  int threads = 0;
 
  
  if (argc > 4)
+
include "PixelMat.hpp"
    {
 
      istringstream buf1(argv[1]);
 
      buf1 >> blocks;
 
      istringstream buf2(argv[2]);
 
      buf2 >> threads;
 
      PixelMat img = (PixelMat) imread(argv[3],-1);
 
      int rows = img.rows;
 
      int cols = img.cols;
 
      int step = img.step;
 
  
      fprintf(stderr, "%d\n", step);
+
void startCUDA(int blocks, int threads, unsigned char* img, int rows, int cols, int step);
  
      startCUDA(blocks, threads, img.data, rows, cols, step);
+
int main(int argc, char* argv[]) {
  
      imwrite(argv[4], img);
+
int blocks = 0;
    }
+
int threads = 0;
  else
 
    cout << "Usage: " << argv[0] << " <blocks> <threads> <inputfile> <outputfile>" << endl;
 
  
  return 0;
+
if (argc &gt; 4)
}
+
  {
 +
    istringstream buf1(argv[1]);
 +
    buf1 &gt;&gt; blocks;
 +
    istringstream buf2(argv[2]);
 +
    buf2 &gt;&gt; threads;
 +
    PixelMat img = (PixelMat) imread(argv[3],-1);
 +
    int rows = img.rows;
 +
    int cols = img.cols;
 +
    int step = img.step;
  
 +
    fprintf(stderr, "%d\n", step);
  
 +
    startCUDA(blocks, threads, img.data, rows, cols, step);
  
// removeblue.cu
+
    imwrite(argv[4], img);
 +
  }
 +
else
 +
  cout &lt;&lt; "Usage: " &lt;&lt; argv[0] &lt;&lt; " &lt;blocks&gt; &lt;threads&gt; &lt;inputfile&gt; &lt;outputfile&gt;" &lt;&lt; endl;
  
// Read image from file, remove blue channel and write onto another file
+
return 0;
  
 +
} </pre>
 +
<br>
 +
<pre><span style="display: none;" id="1323422515452S"> </span>// removeblue.cu // Read image from file, remove blue channel and write onto another file //#include &lt;stdio.h&gt; //#include &lt;stdlib.h&gt; //#include &lt;unistd.h&gt; __device__ void RGB(int x, int y, unsigned char* m, int step, int r, int g, int b) { unsigned char *p; p = ((unsigned char *) (m + step*x))+3*y; *p = (unsigned char) b; *(p+1) = (unsigned char) g; *(p+2) = (unsigned char) r; } __device__ void RGB(int x, int y, unsigned char* m, int step, int rgb) { unsigned char *p; p = ((unsigned char *) (m + step*x)+3*y); *p = (unsigned char) (rgb&amp;0xff); *(p+1) = (unsigned char) ((rgb&gt;&gt;8)&amp;0xff); *(p+2) = (unsigned char) ((rgb&gt;&gt;16)&amp;0xff); } __device__ int RGB(int x, int y, unsigned char* m, int step) { unsigned char *p; unsigned int b; p = ((unsigned char *) (m + step*x)+3*y); b = *p+((*(p+1))&lt;&lt;8)+((*(p+2))&lt;&lt;16); }
 +
__global__ void Count(unsigned char *img, int rows, int cols, int step)<br>{<br> float rowfac = ((float) rows)/gridDim.x;<br> int rowstart = blockIdx.x*rowfac;<br> int rowend = (blockIdx.x+1)*rowfac;<br> float colfac = ((float) cols)/blockDim.x;<br> int colstart = threadIdx.x*colfac;<br> int colend = (threadIdx.x+1)*colfac;<br> for (int i=rowstart; i&lt;rowend; i++)<br> {<br> for (int j=colstart; j&lt;colend; j++)<br> RGB(i, j, img, step, RGB(i, j, img, step)&amp;0x0000ff);<br> }<br>}
  
//#include <stdio.h>
+
void startCUDA(int blocks, int threads, unsigned char* img, int rows, int cols, int step)<br>{<br> if (img!=NULL)<br> {<br> dim3 dimBlock(threads, threads, 1);<br> dim3 dimGrid(blocks, 1, 1);
//#include <stdlib.h>
 
//#include <unistd.h>
 
  
__device__ void RGB(int x, int y, unsigned char* m, int step, int r, int g, int b)
+
unsigned char *CUDAimg;<br> cudaMalloc((void**) &amp;CUDAimg, rows*cols*3);<br> cudaMemcpy(CUDAimg, img, rows*cols*3, cudaMemcpyHostToDevice);<br> Count&lt;&lt;&lt;dimGrid, dimBlock&gt;&gt;&gt;(CUDAimg, rows, cols, step);<br> cudaMemcpy(img, CUDAimg, rows*cols*3, cudaMemcpyDeviceToHost);<br> cudaFree(CUDAimg);<br> }<br>}<span style="display: none;" id="1323422515704E"> </span><br>
{
+
</pre>
  unsigned char *p;
 
  p = ((unsigned char *) (m + step*x))+3*y;
 
  *p = (unsigned char) b;
 
  *(p+1) = (unsigned char) g;
 
  *(p+2) = (unsigned char) r;
 
}
 
__device__ void RGB(int x, int y, unsigned char* m, int step, int rgb)
 
{
 
  unsigned char *p;
 
  p = ((unsigned char *) (m + step*x)+3*y);
 
  *p = (unsigned char) (rgb&0xff);
 
  *(p+1) = (unsigned char) ((rgb>>8)&0xff);
 
  *(p+2) = (unsigned char) ((rgb>>16)&0xff);
 
}
 
__device__ int RGB(int x, int y, unsigned char* m, int step)
 
{
 
  unsigned char *p;
 
  unsigned int b;
 
  p = ((unsigned char *) (m + step*x)+3*y);
 
  b = *p+((*(p+1))<<8)+((*(p+2))<<16);
 

Revision as of 10:25, 9 December 2011

// Read image from file, remove blue channel and write onto another file

include <iostream>

include <string>

include <sstream>

include <cv.h>

include <highgui.h>

using namespace cv; using namespace std;

include "PixelMat.hpp"

void startCUDA(int blocks, int threads, unsigned char* img, int rows, int cols, int step);

int main(int argc, char* argv[]) { 

 int blocks = 0;
int threads = 0;

 if (argc > 4)
   {
     istringstream buf1(argv[1]);
     buf1 >> blocks;
     istringstream buf2(argv[2]);
     buf2 >> threads;
     PixelMat img = (PixelMat) imread(argv[3],-1);
     int rows = img.rows;
     int cols = img.cols;
     int step = img.step;

     fprintf(stderr, "%d\n", step);

     startCUDA(blocks, threads, img.data, rows, cols, step);

     imwrite(argv[4], img);
   }
else
   cout << "Usage: " << argv[0] << " <blocks> <threads> <inputfile> <outputfile>" << endl;

 return 0;

} 


<span style="display: none;" id="1323422515452S"> </span>// removeblue.cu // Read image from file, remove blue channel and write onto another file //#include <stdio.h> //#include <stdlib.h> //#include <unistd.h> __device__ void RGB(int x, int y, unsigned char* m, int step, int r, int g, int b) { unsigned char *p; p = ((unsigned char *) (m + step*x))+3*y; *p = (unsigned char) b; *(p+1) = (unsigned char) g; *(p+2) = (unsigned char) r; } __device__ void RGB(int x, int y, unsigned char* m, int step, int rgb) { unsigned char *p; p = ((unsigned char *) (m + step*x)+3*y); *p = (unsigned char) (rgb&0xff); *(p+1) = (unsigned char) ((rgb>>8)&0xff); *(p+2) = (unsigned char) ((rgb>>16)&0xff); } __device__ int RGB(int x, int y, unsigned char* m, int step) { unsigned char *p; unsigned int b; p = ((unsigned char *) (m + step*x)+3*y); b = *p+((*(p+1))<<8)+((*(p+2))<<16); }
__global__ void Count(unsigned char *img, int rows, int cols, int step)<br>{<br> float rowfac = ((float) rows)/gridDim.x;<br> int rowstart = blockIdx.x*rowfac;<br> int rowend = (blockIdx.x+1)*rowfac;<br> float colfac = ((float) cols)/blockDim.x;<br> int colstart = threadIdx.x*colfac;<br> int colend = (threadIdx.x+1)*colfac;<br> for (int i=rowstart; i<rowend; i++)<br> {<br> for (int j=colstart; j<colend; j++)<br> RGB(i, j, img, step, RGB(i, j, img, step)&0x0000ff);<br> }<br>}

void startCUDA(int blocks, int threads, unsigned char* img, int rows, int cols, int step)<br>{<br> if (img!=NULL)<br> {<br> dim3 dimBlock(threads, threads, 1);<br> dim3 dimGrid(blocks, 1, 1);

unsigned char *CUDAimg;<br> cudaMalloc((void**) &CUDAimg, rows*cols*3);<br> cudaMemcpy(CUDAimg, img, rows*cols*3, cudaMemcpyHostToDevice);<br> Count<<<dimGrid, dimBlock>>>(CUDAimg, rows, cols, step);<br> cudaMemcpy(img, CUDAimg, rows*cols*3, cudaMemcpyDeviceToHost);<br> cudaFree(CUDAimg);<br> }<br>}<span style="display: none;" id="1323422515704E"> </span><br>