CUDA仅处理OpenCV 16位灰度级Mat中总列数的一半


甲基蓝
2025-02-15 12:00:34 (5小时前)
  1. 我正在制作一个初学CUDA程序,它基本上使用OpenCV执行输入灰度图像的下采样。经过测试,它与8位灰度图像配合得很好,但噪点很大......

2 条回复
  1. 0# 那年 | 2019-08-31 10-32



    首先,为了处理16位图像,像素数据必须被解释为16位宽的数据类型,其可能是

    unsigned short

    要么

    short

    。请记住,我们只需要


    </强>
    图像数据为

    unsigned short

    类型;不要输入它。为此,我们将把图像数据指针转换为所需的类型,如下例所示:




    1. unsigned short ptr16 = reinterpret_cast<unsigned short>(im1.ptr());

    2. </code>


    作为上述步骤的结果,我们还必须为16位数据类型创建单独的内核。我们可以通过将内核定义为C ++模板来巧妙地做到这一点。
    所以内核可能如下所示:




    1. template
      global void binCUDAKernel(T input, T output, int binDim, int outputWidth, int outputHeight, int inputWstep, int outputWstep, int nChannels)
      {
      int outXind = blockIdx.x blockDim.x + threadIdx.x;
      int outYind = blockIdx.y
      blockDim.y + threadIdx.y;

    2. if ((outXind < outputWidth) && (outXind > outputWidth/2) && (outYind < outputHeight)) // Only run threads in output image coordinate range
    3. {
    4.     if (nChannels == 1) // Test only for greyscale images
    5.     {
    6.         // Calculate x & y index of input binned pixels corresponding to current output pixel
    7.         int inXstart = outXind * binDim;
    8.         int inYstart = outYind * binDim;
    9.         // Perform binning on identified input pixels
    10.         float sum = 0;
    11.         for (int binY = inYstart; binY < (inYstart + binDim); binY++) {
    12.             for (int binX = inXstart; binX < (inXstart + binDim); binX++) {
    13.                 int input_tid = binY * inputWstep + binX;
    14.                 sum += float(input[input_tid]);
    15.             }
    16.         }
    17.         // Establish output thread index in current output pixel index
    18.         int output_tid = outYind * outputWstep + outXind;
    19.         // Assign binned pixel value to output pixel
    20.         output[output_tid] = static_cast<T>(sum / (binDim*binDim));
    21.     }
    22. }
    23. }

    24. </code>


    使用自定义CUDA内核处理OpenCV Mat时的另一个重要问题是图像步骤必须除以数据类型的大小(以字节为单位)。对于16位图像,单个像素的大小为16位(2个字节),因此内核中使用的步骤必须除以2.请记住,不应修改原始步骤。只是作为内核参数传递的步长值应该被划分。



    结合上述修复,最终的CPU代码可能如下所示:




    1. void binFilter(const cv::Mat input, std::string output_file)
      {
      // 2X2 binning
      int binDim = 2;

    2. // Create blank output image & calculate size of input and output
    3. cv::Size outsize(input.size().width / binDim, input.size().height / binDim);
    4. cv::Mat output(outsize, input.type());
    5. const int inputBytes = input.step * input.rows;
    6. const int outputBytes = output.step * output.rows;
    7. // Allocate memory in device
    8. unsigned char *d_input, *d_output;
    9. gpuErrchk(cudaMalloc<unsigned char>(&d_input, inputBytes));
    10. gpuErrchk(cudaMalloc<unsigned char>(&d_output, outputBytes));
    11. // Copy input image to device
    12. gpuErrchk(cudaMemcpy(d_input, input.ptr(), inputBytes, cudaMemcpyHostToDevice));
    13. // Configure size of block and grid
    14. const dim3 block(16, 16);
    15. const dim3 grid((output.cols + block.x - 1) / block.x, (output.rows + block.y - 1) / block.y); // Additional block for rounding up
    16. int depth = input.depth();
    17. // Execute kernel
    18. if (input.depth() == CV_16U)
    19. {
    20.     typedef unsigned short t16;
    21.     t16* input16 = reinterpret_cast<t16*>(d_input);
    22.     t16* output16 = reinterpret_cast<t16*>(d_output);
    23.     int inputStep16 = input.step / sizeof(t16);
    24.     int outputStep16 = output.step / sizeof(t16);
    25.     binCUDAKernel <t16> <<<grid, block>>> (input16, output16, binDim, output.cols, output.rows, inputStep16, outputStep16, input.channels());
    26. }
    27. else
    28. {
    29.     binCUDAKernel <unsigned char> <<<grid, block>>> (d_input, d_output, binDim, output.cols, output.rows, input.step, output.step, input.channels());   
    30. }
    31. gpuErrchk(cudaPeekAtLastError());
    32. // Wait for all threads to finish
    33. //gpuErrchk(cudaDeviceSynchronize());
    34. // Copy output image from device back to host (cudaMemcpy is a blocking instruction)
    35. gpuErrchk(cudaMemcpy(output.ptr(), d_output, outputBytes, cudaMemcpyDeviceToHost));
    36. // Free device memory
    37. gpuErrchk(cudaFree(d_input));
    38. gpuErrchk(cudaFree(d_output));
    39. // Write image to specified output_file path
    40. cv::imwrite(output_file, output);
    41. }

    42. </code>


    由于分箱算法的逻辑,输出图像中的噪声似乎是引入的混叠。例如,它与使用最近邻方法重新采样图像非常相似。



    更新:



    上面提到的计算像素的存储器地址的方法没有记录,只是直觉的结果,所以它看起来有点不同寻常。
    OpenCV和其他库使用的另一种方法避免了图像步骤划分的混淆。给定像素的x和y索引,它如下进行:




    1. 将图像数据指针重新解释为字节表示(

      unsigned char

      )。


    2. 使用y索引和图像步骤计算图像行的起始地址。


    3. 将行开始地址重新解释为所需类型(

      unsigned short


      )。


    4. 访问行开始指针的x索引。

    5. </醇>


      使用此方法,我们可以计算灰度图像的像素存储器地址,如下所示:




      1. template
        T getPixelAddress(unsigned char data, int x, int y, int step)
        {
        T row = (T)((unsigned char)(data) + y step);
        return row + x;
        }

      2. </code>


      在上述方法中,步长值是原始值,没有任何除法。


登录 后才能参与评论