histogram_1

// Histogram Equalization

#include <wb.h>

#define HISTOGRAM_LENGTH 256

//@@ insert code here

#define wbCheck(stmt)                                                     \
  do {                                                                    \
    cudaError_t err = stmt;                                               \
    if (err != cudaSuccess) {                                             \
      wbLog(ERROR, "Failed to run stmt ", #stmt);                         \
      wbLog(ERROR, "Got CUDA error ...  ", cudaGetErrorString(err));      \
      return -1;                                                          \
    }                                                                     \
  } while (0)

__global__ void FloattoUChar(float *Pin, unsigned char *Pout, int imgsize){
  unsigned int t = blockIdx.xblockDim.x + threadIdx.x;
  if (t < imgsize){
    Pout[t] =  (unsigned char) ((HISTOGRAM_LENGTH - 1) * Pin[t]);
  }
}

__global__ void ColortoGray(unsigned char *Pin, unsigned char *Pout, int width, int height, int channels){
  unsigned int t = blockIdx.x*blockDim.x + threadIdx.x;
  unsigned int imgsize = width*height*channels;
  if (t < imgsize){
    int rgbOffset = t * channels;
    unsigned char r = Pin[rgbOffset    ];
    unsigned char g = Pin[rgbOffset + 1];
    unsigned char b = Pin[rgbOffset + 2];

    Pout[grayOffset] = (unsigned char) (0.21*r + 0.71*g + 0.07*b); ///// check this line again
  }
}

__global__ void GraytoHist(unsigned char *Pin, unsigned int *Pout, int imgsize){
  __shared__ unsigned int hist[HISTOGRAM_LENGTH];

  unsigned int t = blockIdx.x*blockDim.x + threadIdx.x;
  if (threadIdx.x < HISTOGRAM_LENGTH){ // set histogram initial values to zero
    hist[t] = 0;
  }
  __syncthreads();

  if (t < imgsize){     // feeding the histogram, Pin[t] always from 0 to 255, so hist[Pin[t]] matches
    atomicAdd(&hist[Pin[t]], 1);
  }
  __syncthreads();

  if (threadIdx.x < HISTOGRAM_LENGTH){    // adding results from different block to the output
    atomicAdd(&Pout[threadIdx.x], hist[threadIdx.x]);
  }
}


int main(int argc, char **argv) {
  wbArg_t args;
  int imageWidth;
  int imageHeight;
  int imageChannels;
  wbImage_t inputImage;
  wbImage_t outputImage;
  float *hostInputImageData;
  float *hostOutputImageData;
  const char *inputImageFile;

  float *deviceInput;
  unsigned char *deviceUChar;
  unsigned char *deviceGray;
  unsigned int *deviceHist;



  args = wbArg_read(argc, argv); /* parse the input arguments */

  inputImageFile = wbArg_getInputFile(args, 0);

  //Import data and create memory on host
  inputImage = wbImport(inputImageFile);
  imageWidth = wbImage_getWidth(inputImage);
  imageHeight = wbImage_getHeight(inputImage);
  imageChannels = wbImage_getChannels(inputImage);
  outputImage = wbImage_new(imageWidth, imageHeight, imageChannels);

  int imgsize = imageWidth*imageHeight*imageChannels;

  hostInputImageData = wbImage_getData(inputImage);
  hostOutputImageData = wbImage_getData(outputImage);

  wbCheck(cudaMalloc((void**)&deviceInput, imgsize * sizeof(float)));
  wbCheck(cudaMalloc((void**)&deviceUChar, imgsize * sizeof(unsigned char)));
  wbCheck(cudaMalloc((void**)&deviceGray,  imageWidth*imageHeight* sizeof(unsigned char)));
  wbCheck(cudaMalloc((void**)&deviceHist, imageWidth*imageHeight * sizeof(unsigned int)));

  wbCheck(cudaMemcpy(deviceInput, hostInputImageData, imgsize * sizeof(float), cudaMemcpyHostToDevice));
  wbCheck(cudaMemset((void *) deviceHist, 0, HISTOGRAM_LENGTH * sizeof(unsigned int)));

  dim3 DimGrid(ceil(1.0*imgsize/(HISTOGRAM_LENGTH)), 1, 1);
  dim3 DimBlock(HISTOGRAM_LENGTH, 1, 1);

  FloattoUChar <<<DimGrid, DimBlock>>>(deviceInput, deviceUChar, imgsize);
  ColortoGray <<<DimGrid, DimBlock>>>(deviceUChar, deviceGray, imageWidth, imageHeight, imageChannels);
  GraytoHist <<<DimGrid, DimBlock>>>(deviceGray, deviceHist, imageWidth*imageHeight);




  wbSolution(args, outputImage);

  cudaFree(deviceInput);
  cudaFree(deviceUChar);
  cudaFree(deviceGray);

  return 0;
}