opencv cpp 中使用 cuda 核

c/c++

浏览数:47

2020-6-15

AD:资源代下载服务

前提条件

  • 电脑安装英伟达显卡;
  • 安装英伟达驱动,版本建议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 调用时候,多流 的使用;

作者:mastxu