Notes on CS344: Introduction to Parallel Programming

 int *h_buffer = new int[num_elements];
 memset(h_buffer, 0, num_elements * sizeof(int));
 
 int *d_buffer_in;
 cudaMalloc(&d_buffer_in, num_elements * sizeof(int));
 cudaMemcpy(d_buffer_in,   h_buffer,   num_elements * sizeof(int), cudaMemcpyHostToDevice);
 
 int *d_buffer_out;
 cudaMalloc(&d_buffer_out, num_elements * sizeof(int));

 const dim3 blockSize(512, 1, 1); // specify 512 threads for a thread block 
 const dim3 gridSize(ceil(num_elements / blockSize.x), 1, 1); // will launch ceil(num_elements / blockSize.x) number of thread blocks

kernel<<<gridSize, blockSize>>>(d_buffer_in, d_buffer_out, num_elements);

kernel<<<gridSize, blockSize, number_of_elements_in_shared_mem * sizeof(int)>>>(d_buffer_in, d_buffer_out, num_elements);

cudaDeviceSynchronize();

__global__
void rgba_to_greyscale(const uchar4* const rgbaImage, unsigned char* const greyImage, int numRows, int numCols)
{
  //First create a mapping from the 2D block and grid locations
  //to an absolute 2D location in the image, then use that to
  //calculate a 1D offset
  uint tidx = blockIdx.x * blockDim.x + threadIdx.x;
  uint tidy = blockIdx.y * blockDim.y + threadIdx.y;

  uint pos_1d = numCols * tidy + tidx;

  if ( (tidx < numCols) && (tidy < numRows) )
  {
      float r = (float)rgbaImage[pos_1d].x;
      float g = (float)rgbaImage[pos_1d].y;
      float b = (float)rgbaImage[pos_1d].z;
      float channelSum = __fadd_rn(__fadd_rn(__fmul_rn(.299f , r) , __fmul_rn(.587f , g)) , __fmul_rn(.114f , b));

      greyImage[pos_1d] = static_cast<unsigned char>(channelSum);
  }
}

__global__
void separateChannels(const uchar4* const inputImageRGBA,
                      int numRows, int numCols,
                      unsigned char* const redChannel,
                      unsigned char* const greenChannel,
                      unsigned char* const blueChannel)
{
  uint idx = blockIdx.x * blockDim.x + threadIdx.x;
  uint idy = blockIdx.y * blockDim.y + threadIdx.y;
  
  if ( idx >= numCols || idy >= numRows )
  {
     return;
  }

  uchar4 pixel = inputImageRGBA[idy * numCols + idx];

  redChannel[idy * numCols + idx] = pixel.x;
  greenChannel[idy * numCols + idx] = pixel.y;
  blueChannel[idy * numCols + idx] = pixel.z;

}

__global__
void gaussian_blur(const unsigned char* const inputChannel,
                   unsigned char* const outputChannel,
                   int numRows, int numCols,
                   const float* const filter, const int filterWidth)
{  
  // NOTE: If a thread's absolute position 2D position is within the image, but some of
  // its neighbors are outside the image, then you will need to be extra careful. Instead
  // of trying to read such a neighbor value from GPU memory (which won't work because
  // the value is out of bounds), you should explicitly clamp the neighbor values you read
  // to be within the bounds of the image. If this is not clear to you, then please refer
  // to sequential reference solution for the exact clamping semantics you should follow.

  uint idx = blockIdx.x * blockDim.x + threadIdx.x;
  uint idy = blockIdx.y * blockDim.y + threadIdx.y;
  
  if ( idx >= numCols || idy >= numRows )
  {
     return;
  }

  float sum = 0.0;
  for (int filter_y = -filterWidth/2; filter_y <= filterWidth/2; ++filter_y) {
      for (int filter_x = -filterWidth/2; filter_x <= filterWidth/2; ++filter_x) {
          int image_vert_clamp = MIN(MAX(filter_y + idy, 0), static_cast<int>(numRows - 1));
          int image_horiz_clamp = MIN(MAX(filter_x + idx, 0), static_cast<int>(numCols - 1));

          float img_value = static_cast<float>(inputChannel[image_vert_clamp * numCols + image_horiz_clamp]);
          float filter_value = filter[(filterWidth/2 + filter_y)*filterWidth + (filterWidth/2 + filter_x)];
          sum +=  img_value * filter_value; 
      }
  }

  outputChannel[idy * numCols + idx] = sum;
}

__global__
void recombineChannels(const unsigned char* const redChannel,
                       const unsigned char* const greenChannel,
                       const unsigned char* const blueChannel,
                       uchar4* const outputImageRGBA,
                       int numRows,
                       int numCols)
{
  const int2 thread_2D_pos = make_int2( blockIdx.x * blockDim.x + threadIdx.x,
                                        blockIdx.y * blockDim.y + threadIdx.y);

  const int thread_1D_pos = thread_2D_pos.y * numCols + thread_2D_pos.x;

  if (thread_2D_pos.x >= numCols || thread_2D_pos.y >= numRows)
  {
     return;
  }

  unsigned char red   = redChannel[thread_1D_pos];
  unsigned char green = greenChannel[thread_1D_pos];
  unsigned char blue  = blueChannel[thread_1D_pos];

  //Alpha should be 255 for no transparency
  uchar4 outputPixel = make_uchar4(red, green, blue, 255);

  outputImageRGBA[thread_1D_pos] = outputPixel;
}