CoE Cluster november 2011/CUDA

// 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>