使用 cuda 从 RGBA 图像中分离通道（无法显示完整图像）

Question

我正在尝试使用 cuda 将通道与图像分开。程序输出三个通道对应的图像。我得到了正确的输出，但它只显示了部分图像通道。

这是我的代码：

  // main.cpp
  void separateHelper(const uchar4 *d_rgbaImage, uchar4 *d_channel, const int numRows, const int numCols,int channel);

  std::string file_name = "test.jpg";
  cv::Mat image, rgbaImage;
  int numRows(){ return rgbaImage.rows; };
  int numCols(){ return rgbaImage.cols; };

  int main(){

  uchar4 *h_rgbaImage, *h_red, *h_green, *h_blue;
  uchar4 *d_rgbaImage, *d_red, *d_green, *d_blue;
  cv::Mat red, green, blue;
  cv::Mat redChannel, greenChannel, blueChannel;

  image = cv::imread(file_name.c_str(),CV_LOAD_IMAGE_COLOR);
  if (image.empty()){
      std::cerr << "error loading image";
      system("pause");
      exit(1);
  }

  cv::cvtColor(image,rgbaImage, CV_BGR2RGBA);
  //create space for output
  red.create(numRows(), numCols(), CV_8UC3);
  cv::cvtColor(red, redChannel, CV_BGRA2RGBA);
  green.create(numRows(), numCols(), CV_8UC3);
  cv::cvtColor(green, greenChannel, CV_BGRA2RGBA);
  blue.create(numRows(), numCols(), CV_8UC3);
  cv::cvtColor(blue, blueChannel, CV_BGRA2RGBA);

  h_rgbaImage = (uchar4*)rgbaImage.ptr<unsigned char>(0);
  h_red = (uchar4*)redChannel.ptr<unsigned char>(0);
  h_green = (uchar4*)greenChannel.ptr<unsigned char>(0);
  h_blue = (uchar4*)blueChannel.ptr<unsigned char>(0);

  //allocate memory on device
  const int numPixels = numCols()*numRows();
  checkCudaErrors(cudaMalloc((void**)&d_rgbaImage,sizeof(uchar4) * (numPixels + 500)));
  checkCudaErrors(cudaMalloc((void**)&d_red, sizeof(uchar4) * (numPixels + 500)));
  checkCudaErrors(cudaMalloc((void**)&d_green, sizeof(uchar4) * (numPixels + 500)));
  checkCudaErrors(cudaMalloc((void**)&d_blue, sizeof(uchar4) * (numPixels + 500)));

  //copy image from host to device
  checkCudaErrors(cudaMemcpy(d_rgbaImage, h_rgbaImage, sizeof(uchar4) * numPixels, cudaMemcpyHostToDevice));

  //call helper function of kernel
  separateHelper(d_rgbaImage, d_red, numRows(), numCols(),1);
  separateHelper(d_rgbaImage, d_green, numRows(), numCols(),2);
  separateHelper(d_rgbaImage, d_blue, numRows(), numCols(),3);

  //copy results back to host
  checkCudaErrors(cudaMemcpy(h_red, d_red, sizeof(uchar4) * numPixels, cudaMemcpyDeviceToHost));
  checkCudaErrors(cudaMemcpy(h_green, d_green, sizeof(uchar4) * numPixels, cudaMemcpyDeviceToHost));
  checkCudaErrors(cudaMemcpy(h_blue, d_blue, sizeof(uchar4) * numPixels, cudaMemcpyDeviceToHost));

  //change RGBA to BGR
  cv::cvtColor(redChannel,red,CV_RGBA2BGR);
  cv::cvtColor(greenChannel,green,CV_RGBA2BGR);
  cv::cvtColor(blueChannel,blue,CV_RGBA2BGR);

  cv::namedWindow("RED", CV_WINDOW_AUTOSIZE);
  cv::imshow("RED", red);
  cv::namedWindow("GREEN", CV_WINDOW_AUTOSIZE);
  cv::imshow("GREEN", green);
  cv::namedWindow("BLUE", CV_WINDOW_AUTOSIZE);
  cv::imshow("BLUE", blue);
  cv::waitKey(0);

  cudaFree(d_rgbaImage);
  cudaFree(d_red);
  cudaFree(d_green);
  cudaFree(d_blue);
  return 0;
}

这是我的 GPU 代码：

// kernel.cu
__global__ void separateChannels(const uchar4* d_rgbaImage,uchar4* d_channel, int numRows, int numCols, int channel){
  int x = threadIdx.x + blockIdx.x * blockDim.x;
  int y = threadIdx.y + blockIdx.y * blockDim.y;
  if (x >= numCols || y >= numRows)
      return;
  int index = numRows * y + x;
  if (channel == 1){
      d_channel[index].x = d_rgbaImage[index].x;
      d_channel[index].y = 0;
      d_channel[index].z = 0;
  }
  else if (channel == 2){
      d_channel[index].x = 0;
      d_channel[index].y = d_rgbaImage[index].y;
      d_channel[index].z = 0;
  }
  else if (channel == 3){
      d_channel[index].x = 0;
      d_channel[index].y = 0;
      d_channel[index].z = d_rgbaImage[index].z;
  }
  d_channel[index].w = 255;
}

void separateHelper(const uchar4 *d_rgbaImage, uchar4 *d_channel,
    const int numRows, const int numCols, int channel){


  //set grid and block size
  int blockWidth = 32;
  const dim3 blockSize(blockWidth, blockWidth, 1);
  const dim3 gridSize(((numCols)/32 + 1 ), ((numRows)/32 + 1), 1);
  //call kernel
  separateChannels <<<gridSize, blockSize >>>(d_rgbaImage, d_channel, numRows, numCols, channel);

  cudaDeviceSynchronize();
  checkCudaErrors(cudaGetLastError());
}        

错误：只有一部分图像（红色、绿色和蓝色通道图像）显示为输出。

Answer 1

我认为要么没有足够的线程分配来执行任务，要么你混淆了 x 和 y 坐标。通常，y 方向的条带分配有列，x 方向的条带分配有行。每行包含numColumns 元素，每列包含numRows 元素。当你分配线程时，你遵循这个逻辑：

int blockWidth = 32;
const dim3 blockSize(blockWidth, blockWidth, 1);
const dim3 gridSize(((numCols)/32 + 1 ), ((numRows)/32 + 1), 1);

但是当你计算指数时你没有。不应该

int index = numRows * y + x;

是：

int index = numColumns * y + x;

？

Answer 2

我太兴奋了，我刚刚解决了我的问题！

我的情况是，C++ 代码结果是正确的，但 GPU 代码结果只显示完整图像的四分之一。 那是因为当 cudaMemcpy 从设备到主机时，我设置了错误的参数 'size'。

//cudaMemcpy(h_result, d_result, imagesize, cudaMemcpyDeviceToHost);

//cudaMemcpy(h_result, d_result, imagesize*sizeof(float), cudaMemcpyDeviceToHost);

sizeof(float) 正好是 4 个字节！所以我只得到了完整图像的四分之一。

记住乘以 sizeof（数据类型）。

希望我的回答有用:)

Answer 3

由于您的网格尺寸错误，您只能获得部分图像通道。您需要在此处替换 numCols 和 numRows：

 const dim3 gridSize(((numCols)/32 + 1 ), ((numRows)/32 + 1), 1);

像这样：

const dim3 gridSize(((numRows)/32 + 1 ), ((numCols)/32 + 1), 1);

而且这里不需要多加500：

checkCudaErrors(cudaMalloc((void**)&d_blue, sizeof(uchar4) * (numPixels + 500)));