cuda kernel中访问cv::cuda::Gpumat的方法

高性能的图像算法通常使用GPU加速，OpenCV中的cuda模块提供了常用的算法函数，可直接在GPU中运行。对于复杂的应用，cuda模块中的函数无法满足要求，这时需要自己写cuda kernel。

以下代码为cuda kernel中访问OpenCV的数据结构cv::cuda::Gpumat的示例：

#include <cuda_runtime.h>
#include <opencv2/opencv.hpp>
#include <opencv2/core.hpp>
#include <opencv2/cudaarithm.hpp>


__global__ void Kernel_CalPhase2D(const cv::cuda::PtrStepSzf dReal, const cv::cuda::PtrStepSzf dImag,
                                      const float d_Pi, cv::cuda::PtrStepSzf dOutput)
{
    int iCol = blockIdx.x * blockDim.x + threadIdx.x; int iRow = blockIdx.y * blockDim.y + threadIdx.y;

    if (iCol < dReal.cols && iRow < dReal.rows && iRow >= 0 && iCol >= 0)
    {
        dOutput(iRow, iCol) = (atan2(dImag(iRow, iCol), dReal(iRow, iCol)) * (-1) + d_Pi) / (2 * d_Pi);
    }
}

void CalculatePhase2D(const cv::InputArray _input0, const cv::InputArray _input1, cv::OutputArray _output)
{
    const cv::cuda::GpuMat input0 = _input0.getGpuMat();
    const cv::cuda::GpuMat input1 = _input1.getGpuMat();

    _output.create(input0.size(), input0.type());

    cv::cuda::GpuMat output0 = _output.getGpuMat();

    dim3 cthreads(32, 32);
    dim3 cblocks(
                static_cast<int>(std::ceil(input0.size().width /
                                           static_cast<double>(cthreads.x))),
                static_cast<int>(std::ceil(input0.size().height /
                                           static_cast<double>(cthreads.y))));

    Kernel_CalPhaseInfo2D << <cblocks, cthreads >> > (input0, input1, 3.14159f, output0);

    if (cudaSuccess != cudaGetLastError())
        std::cout << "CalculatePhaseInfo2D(): gave an error" << std::endl;

    return;
}