opencv cpp 中使用 cuda 核

前提条件

• 电脑安装英伟达显卡；
• 安装英伟达驱动，版本建议10以上；
• 安装 cuda 和 cudnn；
• github 下载 opencv，opencv_contribute;本地 cmake 编译 opencv；

目标

• 在 cpp 中，使用本身写的 cuda 核函数；

写这篇文章的缘由

最近在写 cuda 加速，而后发现，把核函数卸载 .cu 文件中，调用写在.cpp文件中，总会报出错，好比：

C2039    “atomicAdd”: 不是“`global namespace'”的成员
C2039    “atomicMin”: 不是“`global namespace'”的成员
C2664    “uchar1 cv::cudev::max(const uchar1 &,const uchar1 &)”: 没法将参数 1 从“const unsigned char”转换为“const uchar1 &”

花费了一些时间，在网上查找，没有找到；因而想，写出来，记录本身的解决过程，帮助他人尽快找到答案；

//.cpp
//cuda 相关头文件
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda_device_runtime_api.h"

//标准库
#include <iostream>
#include <string>

//opencv 使用的头文件
#include "opencv2/core/core.hpp"  
#include "opencv2/highgui/highgui.hpp"  
#include "opencv2/imgproc/imgproc.hpp"

// 引用一些opencv已经有的头文件；
#include "opencv2/cudaarithm.hpp"

//cpp万万不能引用
//#include "opencv2/cudev.hpp"

////.cu
#include <cuda.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda_device_runtime_api.h"

//使用 opencv 的形参类型 
#include "opencv2/cudev.hpp"

​ 起初，网上查到 “opencv2/cudev.hpp” 是opencv cuda 的基础头文件；因而天然引用到了 cpp 和 cu 文件中，一直报错，一个一个找，最终是由于 cpp 中引用“opencv2/cudev.hpp” 头文件，产生编译错误；

​ 若是要使用 opencv 的 形参类型，好比 cv::cuda::PtrStep，cv::cuda::PtrStepSz 能够在 .cu 中引入 “opencv2/cudev.hpp” 头文件；

​ 看 cuda 的 example，通常把 核函数A 和 核函数的调用AA 都放在 .cu文件中，.cpp 中再调用AA 就好。

代码 以下：

.cpp 文件

//.cpp
#include<iostream>

#include <cuda.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda_runtime_api.h"

#include "opencv2/cudacodec.hpp" //
//#include "opencv2/cudaimgproc.hpp" // 
//#include "opencv2/cudaobjdetect.hpp"
//#include "opencv2/cudaarithm.hpp"

using namespace cv;
using namespace std;

//初始化
cv::cuda::Stream sdr2hdrStream;

cv::cuda::GpuMat dev_src_yuv8_GPU(height * 3/2, width, CV_8UC1);
cv::cuda::GpuMat dev_src_RGB_32F_GPU(height, width, CV_32FC3);

Mat src = imread("lena.jpg");
Mat sc_yuv;
cvtColor(src, sc_yuv, COLOR_BGR2YUV_I420);

dev_src_yuv8_GPU.upload(sc_yuv, sdr2hdrStream);
cv::cuda::PtrStep<uchar> psrc(dev_src_yuv8_GPU.data, dev_src_yuv8_GPU.step);
cv::cuda::PtrStep<float3> pdst(dev_src_RGB_32F_GPU.ptr<float3>(0), dev_src_RGB_32F_GPU.step);

yuv2rgb420pCudaGPUMat(psrc, pdst, width, height, sdr2hdrStream);

//等待 stream 执行完
sdr2hdrStream.waitForCompletion();
Mat ttttt;
dev_src_RGB_32F_GPU.download(ttttt);
cv::imwrite("tttt_dTmp_dst.png", ttttt * 255);

.cu 文件

#include <cuda.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda_device_runtime_api.h"

#include "cuda_runtime_api.h"
#include "cuda_texture_types.h"

#include "opencv2/cudev.hpp" //cuda 

// cuda threads
int hlg_threads = 16;

//核函数
__global__ void yuv2rgb420p_private_GPUMat(cv::cuda::PtrStep<uchar> src, cv::cuda::PtrStep<float3> dst, int w, int h) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    if (x < w && y < h) {
        //未优化前
        //float y_ = (src(y,x) - 16) / 219.0;
        //优化除法
        float y_ = __fdividef((src(y, x) - 16), 219.0F);
        float u_ = __fdividef((src(h + (y >> 2), x >> 1) - 128), 224.0F);
        float v_ = __fdividef((src(((h * 5) >> 2) + (y >> 2), x >> 1) - 128), 224.0F);
        //R
        dst(y, x).x = __saturatef(__fmul_rn(1.4746F, v_) + y_);
        //B
        dst(y, x).z = __saturatef(__fmul_rn(1.8814F, u_) + y_);
        //G
        dst(y, x).y = __saturatef(__fmul_rn(1.4749F, y_) - __fmul_rn(0.3875F, dst(y, x).x) - __fmul_rn(0.0875F, dst(y, x).z));
    }
}

//核函数 调用
void yuv2rgb420pCudaGPUMat(cv::cuda::PtrStep<uchar> src, cv::cuda::PtrStep<float3> dst,    int w, int h, cv::cuda::Stream& sdr2hdrStream_) {

    int bx = (w + hlg_threads - 1) / hlg_threads;
    int by = (h + hlg_threads - 1) / hlg_threads;
    dim3 blocks(bx, by);
    dim3 threads(hlg_threads, hlg_threads);
    cudaStream_t s = cv::cuda::StreamAccessor::getStream(sdr2hdrStream_);
    yuv2rgb420p_private_GPUMat << <blocks, threads, 0, s >> > (src, dst, w, h);

}

核函数，能够优化的点有以下几个：

• 定义变量，使用时候，浮点数 写为 0.9F，而不是 0.9；避免 double 转换 float；
• 除法 使用 __fdividef，大概 20个时钟周期，比表达式 “9 / 5” ， 36 个时钟周期快一些；
• __saturatef，0--1 截断函数，小于0 为0，大于1 为 1 ；其余为输入值；

计划

• 计划写 cuda 纹理读取，texture 的使用；对比 纹理对象 cudaTextureObject_t 的使用；
• cuda 调用时候，多流 的使用；