高性能的图像算法通常使用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;
}