c++ - 生成三 channel LUT 掩码的有效方法

Question

我有一些代码可以创建一个看起来大致像这样的 HSV 掩码(取自 this Japanese-language page ):

void colorExtraction(const cv::gpu::GpuMat &src,
                     cv::gpu::GpuMat       *dst)
{
    cv::Mat lut(256, 1, CV_8UC3);
    for (int i = 0; i < 256; i++)
    {
        cv::Vec3b thisHSV;
        thisHSV[0] = inHRange(i) ? 255 : 0;
        thisHSV[1] = inSRange(i) ? 255 : 0;
        thisHSV[2] = inVRange(i) ? 255 : 0;

        lut.at<cv::Vec3b>(i) = thisHSV;
    }

    /* apply LUT to input image */
    cv::gpu::GpuMat extracted(src.size(), CV_8UC3);
    cv::gpu::LUT(src, lut, extracted);

    /* divide image into each channel */
    std::vector<cv::gpu::GpuMat> channels;
    cv::gpu::split(extracted, channels);

    /* create mask */
    cv::gpu::bitwise_and(channels[0], channels[1], *dst);
    cv::gpu::bitwise_and(*dst, channels[2], *dst);
}

这可行，但尽管操作主要在 GPU 中进行，但它比我希望的要慢，这可能是由于许多中间 GpuMat 的缘故。我怀疑可能有一种很好的方法可以将这一切折叠成一个或两个调用，但我不知道它可能是什么。编写我自己的 CUDA 内核当然是一种选择，但我想看看我是否不需要重新发明轮子。

Answer 1

为了 self 回答，我最终编写了自己的内核来执行 LUT()、split() 和两个 bitwise_and()在一次调用中:

__global__ void colorExtractionKernel(cv::gpu::PtrStepSz<uchar3> const  src,
                                      cv::gpu::PtrStepSz<unsigned char> dst,
                                      cv::gpu::PtrStepSz<uchar3> const  lut)
{
    unsigned int row = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int col = blockIdx.x * blockDim.x + threadIdx.x;

    // Extract post-LUT hsv flags
    uchar3 srcHSV = src(row, col);
    unsigned char h = lut(srcHSV.x, 0).x;
    unsigned char s = lut(srcHSV.y, 0).y;
    unsigned char v = lut(srcHSV.z, 0).z;

    // Result pixel is the AND of the pixels
    dst(row, col) = (h & s & v);
}

void colorExtraction_cuda(const cv::gpu::GpuMat &src,   // input HSV image
                          cv::gpu::GpuMat       &dst,   // specified color extracted binarized image
                          cv::Mat const         &lut)   // Look-up thresholds
{
    cudaStream_t thisStream;
    gpuErrChk(cudaStreamCreate(&thisStream));

    dim3 Threads(32, 16);
    dim3 Blocks((src.cols + Threads.x - 1)/Threads.x, (src.rows + Threads.y - 1)/Threads.y);

    cv::gpu::GpuMat gpuLut(lut);

    colorExtractionKernel<<<Blocks, Threads, 0, thisStream>>>(src, dst, gpuLut);

    gpuErrChk(cudaStreamSynchronize(thisStream));
    gpuErrChk(cudaStreamDestroy(thisStream));
    gpuErrChk(cudaGetLastError());
}