CUDA总结：opencv图像金字塔函数分析

基于 opencv 3.1.0

相关源码
..\sources\modules\cudalegacy\src\image_pyramid.cpp
..\sources\modules\cudalegacy\src\cuda\NCVPyramid.cu

类
cv::cuda::ImagePyramidImpl

金字塔分解，调用层次：
cv::cuda::ImagePyramidImpl::ImagePyramidImpl ->
cv::cuda::device::pyramid::downsampleX2 ->
kernelDownsampleX2_gpu ->
kernelDownsampleX2 (cuda kernel)

template<typename T>
__global__ void kernelDownsampleX2(T *d_src,    //d for device
                                   Ncv32u srcPitch,
                                   T *d_dst,
                                   Ncv32u dstPitch,
                                   NcvSize32u dstRoi)
{
    Ncv32u i = blockIdx.y * blockDim.y + threadIdx.y;
    Ncv32u j = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < dstRoi.height && j < dstRoi.width)
    {
    //srcPitch是每行的字节长度
    //两行合并为一行
        T *d_src_line1 = (T *)((Ncv8u *)d_src + (2 * i + 0) * srcPitch);
        T *d_src_line2 = (T *)((Ncv8u *)d_src + (2 * i + 1) * srcPitch);
        T *d_dst_line = (T *)((Ncv8u *)d_dst + i * dstPitch);

        T p00 = d_src_line1[2*j+0];
        T p01 = d_src_line1[2*j+1];
        T p10 = d_src_line2[2*j+0];
        T p11 = d_src_line2[2*j+1];
     //2倍抽取：两行相邻四个点取平均
     //也可以采用高斯滤波
        d_dst_line[j] = _average4(p00, p01, p10, p11);
    }
}

获得金字塔第n层图像，调用层次：

cv::cuda::ImagePyramidImpl::getLayer->
cv::cuda::device::pyramid::interpolateFrom1->
kernelInterpolateFrom1_gpu->
kernelInterpolateFrom1 (cuda kernel)

template<typename T>
__global__ void kernelInterpolateFrom1(T *d_srcTop,
                                       Ncv32u srcTopPitch,
                                       NcvSize32u szTopRoi,
                                       T *d_dst,
                                       Ncv32u dstPitch,
                                       NcvSize32u dstRoi)
{
    Ncv32u i = blockIdx.y * blockDim.y + threadIdx.y;
    Ncv32u j = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < dstRoi.height && j < dstRoi.width)
    {
        Ncv32f ptTopX = 1.0f * (szTopRoi.width - 1) * j / (dstRoi.width - 1);
        Ncv32f ptTopY = 1.0f * (szTopRoi.height - 1) * i / (dstRoi.height - 1);
        Ncv32u xl = (Ncv32u)ptTopX;   
        Ncv32u xh = xl+1;
        Ncv32f dx = ptTopX - xl;
        Ncv32u yl = (Ncv32u)ptTopY;
        Ncv32u yh = yl+1;
        Ncv32f dy = ptTopY - yl;
//相邻两行之间插值，获得新一行
        T *d_src_line1 = (T *)((Ncv8u *)d_srcTop + yl * srcTopPitch);
        T *d_src_line2 = (T *)((Ncv8u *)d_srcTop + yh * srcTopPitch);
        T *d_dst_line = (T *)((Ncv8u *)d_dst + i * dstPitch);

        T p00, p01, p10, p11;
        p00 = d_src_line1[xl];
        p01 = xh < szTopRoi.width ? d_src_line1[xh] : p00;
        p10 = yh < szTopRoi.height ? d_src_line2[xl] : p00;
        p11 = (xh < szTopRoi.width && yh < szTopRoi.height) ? d_src_line2[xh] : p00;
        typedef typename TConvBase2Vec<Ncv32f, NC(T)>::TVec TVFlt;
        TVFlt m_00_01 = _lerp<T, TVFlt>(p00, p01, dx);   //线性插值
        TVFlt m_10_11 = _lerp<T, TVFlt>(p10, p11, dx);
        TVFlt mixture = _lerp<TVFlt, TVFlt>(m_00_01, m_10_11, dy);
        T outPix = _pixDemoteClampZ<TVFlt, T>(mixture);   //像素值饱和处理

        d_dst_line[j] = outPix;
    }
}

线性插值
_lerp函数，支持多通道数据，下面是单通道的定义：

template<typename Tin, typename Tout> struct __lerp_CN<Tin, Tout, 1> {
static __host__ __device__ Tout _lerp_CN(const Tin &a, const Tin &b, Ncv32f d)
{
    typedef typename TConvVec2Base<Tout>::TBase TB;  //定义输出数据类型
    return _pixMake(TB(b.x * d + a.x * (1 - d)));   
    //!b*d+a*（1-d） 等价于 (b-a)*d+a
    //! d取值[0，1]，代表线性插值系数，d越大，插值越接近于b，d越小，插值越接近于a
}};

---

能否利用纹理内存的硬件插值特性加速下采样和上采样处理？
金字塔的cuda实现采用线性插值法，而cpu实现则采用高斯滤波，两者的计算结果是不一样的，前者的实时性更好，精度更差。

相比cpu版本的代码，cuda版本的代码并没有实现拉普拉斯金字塔，也没有金字塔重构。