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 調用時候，多流 的使用；