How to pass the cv::Mat or cv::cuda::GpuMat to custom CUDA kernel?

edit

I usually pass GpuMat to kernel like this:

__global__ void Kernel_CalPhaseInfo2D(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 CalculatePhaseInfo2D(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, M_PI, output0);

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

return;

}