2. 程式人生 > >cuda與openCV結合程式設計（一）

cuda與openCV結合程式設計（一）

阿新 • • 發佈：2019-02-14
    學習計算機影象處理演算法的童鞋，就不得不學習cuda，為啥呢？因為影象處理一般都是矩陣運算，動不動就是百萬的計算量這個時候優化計算時間是必不可少的。openCV本身提供了很多cuda函式，能夠滿足大多數使用者的需求。但是也不絕對，有時候我們需要自己定義一個核心函式進行優化，當然你也可以用openGL或者多執行緒，openCV也提供較好的支援，掌握一種或多種加速演算法，對程式設計師特別是演算法工程師來講很重要。
    閒話不多說，再學習了cuda的基礎以後【cuda並行程式設計基礎（一）、cuda並行程式設計基礎（二）】，我們其實就具備與opencv聯合程式設計的能力，雖然不是最佳優化，但已經可以滿足大多數需求了。
一、cuda與openCV結合方法
    （下面僅涉及windows環境）
    1.我們知道，cuda程式碼一般以.cu結尾（windows，其他系統除外，下同），它的編譯器是nvcc，編譯時它會將CPU程式碼和GPU程式碼分開，CPU部分其實與gcc編譯差不多，GPU部分就按照nvcc的規則編譯，這玩意其實並不複雜;
    2.openCV的程式碼一般都是以.cpp結尾，它的編譯器一般是gcc、g++(或者其他相似編譯器),那麼可不可以將openCV程式碼用nvcc編譯呢?答案是肯定的,但在windows系統,你得把它改為.cu結尾.
    3.所以，在windows系統裡面，你有兩種辦法讓openCV結合cuda程式設計：
        a.openCV正常編譯，cuda程式碼編譯好後，作為靜態庫引入openCV呼叫；
        b.openCV與cuda程式碼混在一起，統一用nvcc編譯。
二、如何編寫程式碼

    (下面openCV基於3.2.0版本)
    openCV是一個非常強大的視覺演算法庫，當然也支援cuda咯。
    cv::cuda是一個專門處理cuda的名稱空間，你在這個名稱空間裡面可以看到很多已經整合好的函式。
        如：cuda::remap()、cuda::add()等
    我們要用到的是cuda::PtrStepSz<T>的模板，以及cuda::GpuMat
    比如：如果我們有一個cuda::GpuMat型別的img，我們怎麼傳入cuda裡面呢？答案就是，直接將img傳到cuda::PtrStepSz裡面，他們是不是等同，但是可以互傳資料，具體見樣例。至於傳到cuda::PtrStepSz裡面如何操作，那就跟cuda差不多了。
    除了cuda::PtrStepSz，openCV還有其他介面可以提供互傳，自己去摸索啦，這裡就不囉嗦了。     至於cuda與openCV的結合程式設計效率問題？哈，誰用誰知道，你不用也無需知道，有興趣自己去測一下咯，反正筆者是牆裂推薦的，後面有空再講效率問題。
三、常見錯誤

    1.cudaErrorMemoryAllocation，主要是申請空間太大，超出了GPU限制；
    2.cudaErrorLaunchFailure，訪問了非法地址，比如index超過了陣列大小；
    3.cuda與vs2015結合程式設計，偶爾會出現抽筋的問題，比如你這次編譯出錯，改正了以後再編譯還出錯，建議要重新編譯時，把以前的編譯生成的東西全刪掉，這樣就保險多了，筆者遇見多次這種情況；
    4.<<<>>>核心符號報錯，要確定它出現在cu檔案裡而不是cpp檔案裡，cu檔案會顯示紅色，不用管它；
    5.靜態庫的編寫規範，額，自己上網研究吧，其實我寫得也不太規範，吐槽一下，網上好多技術文章抄來抄去很沒意思，很多大牛又寫得太過高深，研究不出個所以然來，也希望能夠在各個層次都有合適的文章介紹吧，這樣入門和進階也不會太困難。
第一個程式，直接在cu檔案實現cuda與opencv結合程式設計，非常重要哦

一般我們不這麼使用，因為cuda作為獨立的程式設計方式，放在一起容易混亂，而且為支援高速運算，一般都使用c運算，而不是c++
//opencv_cuda.cu:使用自定義函式，實現cuda版本圖片翻轉
//authored by alpc40
//version：visual studio 2015\cuda toolkit 8.0\openCV 3.2.0
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include<opencv2/opencv.hpp>
#include<iostream>
using namespace std;
using namespace cv;
#ifdef _DEBUG
#pragma comment ( lib,"opencv_core320d.lib")
#pragma comment ( lib,"opencv_highgui320d.lib")
#pragma comment ( lib,"opencv_calib3d320d.lib")
#pragma comment ( lib,"opencv_imgcodecs320d.lib")
#pragma comment ( lib,"opencv_imgproc320d.lib")
#pragma comment ( lib,"opencv_cudaimgproc320d.lib")
#pragma comment ( lib,"opencv_cudaarithm320d.lib")
#pragma comment ( lib,"cudart.lib")
#else
#pragma comment ( lib,"opencv_core320.lib")
#pragma comment ( lib,"opencv_highgui320.lib")
#pragma comment ( lib,"opencv_calib3d320.lib")
#pragma comment ( lib,"opencv_imgcodecs320.lib")
#pragma comment ( lib,"opencv_imgproc320.lib")
#pragma comment ( lib,"opencv_cudaimgproc320.lib")
#pragma comment ( lib,"opencv_cudaarithm320.lib")
#pragma comment ( lib,"cudart.lib")
#endif
//出錯處理函式
#define CHECK_ERROR(call){\
    const cudaError_t err = call;\
    if (err != cudaSuccess)\
    {\
        printf("Error:%s,%d,",__FILE__,__LINE__);\
        printf("code:%d,reason:%s\n",err,cudaGetErrorString(err));\
        exit(1);\
    }\
}
//核心函式：實現上下翻轉
__global__ void swap_image_kernel(cuda::PtrStepSz<uchar3> cu_src, cuda::PtrStepSz<uchar3> cu_dst, int h, int w)
{
    //計算的方法：參看前面兩文
    unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
    //為啥要這樣限制：參看前面兩文
    if (x < cu_src.cols && y < cu_src.rows)
    {
        //為何不是h-y-1,而不是h-y，自己思考哦
        cu_dst(y, x) = cu_src(h - y - 1, x);
    }
}
//呼叫函式，主要處理block和grid的關係
void swap_image(cuda::GpuMat src,cuda::GpuMat dst,int h, int w)
{
    assert(src.cols == w && src.rows ==h);
    int uint = 32;
    //參考前面兩文的block和grid的計算方法，注意不要超過GPU限制
    dim3 block(uint, uint);
    dim3 grid((w + block.x - 1) / block.x, (h + block.y - 1) / block.y);
    printf("grid = %4d %4d %4d\n",grid.x,grid.y,grid.z);
    printf("block= %4d %4d %4d\n",block.x,block.y,block.z);
    swap_image_kernel << <grid, block >> > (src,dst,h,w);
    //同步一下，因為計算量可能很大
    CHECK_ERROR(cudaDeviceSynchronize());
}
int main(int argc,char **argv)
{
    Mat src, dst;
    cuda::GpuMat cu_src, cu_dst;
    int h, w;
    //根據argv[1]讀入圖片資料，BGR格式讀進來
    src = imread(argv[1]);
    //檢測是否正確讀入
    if (src.data == NULL)
    {
        cout << "Read image error" << endl;
        return -1;
    }
    h = src.rows; w = src.cols;
    cout <<"圖片高：" << h << ",圖片寬：" << w << endl;
    //上傳CPU影象資料到GPU，跟cudaMalloc和cudaMemcpy很像哦，其實upload裡面就是這麼寫的
    cu_src.upload(src);
    //申請GPU空間，也可以到函式裡申請，不管怎樣總要申請，要不然核心函式會爆掉哦
    cu_dst = cuda::GpuMat(h, w, CV_8UC3, Scalar(0, 0, 0));
    //申請CPU空間
    dst = Mat(h, w, CV_8UC3, Scalar(0, 0, 0));
    //呼叫函式swap_image,由該函式呼叫核心函式，這樣層次分明，不容易出錯
    //當然你也可以直接在這裡呼叫核心函式，東西太多程式碼容易亂
    swap_image(cu_src,cu_dst,h, w);
    //下載GPU資料到CPU，與upload()對應
    cu_dst.download(dst);
    //顯示cpu影象，如果安裝了openCV集成了openGL,那可以直接顯示GpuMat
    imshow("dst",dst);
    //等待按鍵
    waitKey();
    //寫圖片到檔案
    if(argc==3)
        imwrite(argv[2],dst);
    return 0;
}
第二個程式，使用靜態庫的方式實現cuda與openCV的結合，非常重要哦
這種方式兩相分離，更好實現了這種功能 //swap_image.cu:生成swap_image.lib，供主函式呼叫 
//authored by alpc40
//version：visual studio 2015\cuda toolkit 8.0\openCV 3.2.0
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include<opencv2/opencv.hpp>
using namespace cv;
//出錯處理函式
#define CHECK_ERROR(call){\
    const cudaError_t err = call;\
    if (err != cudaSuccess)\
    {\
        printf("Error:%s,%d,",__FILE__,__LINE__);\
        printf("code:%d,reason:%s\n",err,cudaGetErrorString(err));\
        exit(1);\
    }\
}
//核心函式：實現上下翻轉
__global__ void swap_image_kernel(cuda::PtrStepSz<uchar3> cu_src, cuda::PtrStepSz<uchar3> cu_dst, int h, int w)
{
    //計算的方法：參看前面兩文
    unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
    //為啥要這樣限制：參看前面兩文
    if (x < cu_src.cols && y < cu_src.rows)
    {
        //為何不是h-y-1,而不是h-y，自己思考哦
        cu_dst(y, x) = cu_src(h - y - 1, x);
    }
}
//呼叫函式，主要處理block和grid的關係，注意extern哦，它是庫檔案編寫規範
extern "C" void swap_image(cuda::GpuMat src, cuda::GpuMat dst, int h, int w)
{
    assert(src.cols == w && src.rows == h);
    int uint = 32;
    //參考前面兩文的block和grid的計算方法，注意不要超過GPU限制
    dim3 block(uint, uint);
    dim3 grid((w + block.x - 1) / block.x, (h + block.y - 1) / block.y);
    printf("grid = %4d %4d %4d\n", grid.x, grid.y, grid.z);
    printf("block= %4d %4d %4d\n", block.x, block.y, block.z);
    swap_image_kernel << <grid, block >> > (src, dst, h, w);
    //同步一下，因為計算量可能很大
    CHECK_ERROR(cudaDeviceSynchronize());
}

//opencv_cuda.cpp:第二個程式主函式,使用自編靜態庫，實現cuda版本圖片翻轉
//authored by alpc40
//version：visual studio 2015\cuda toolkit 8.0\openCV 3.2.0
#include <stdio.h>
#include<opencv2/opencv.hpp>
#include<iostream>
using namespace std;
using namespace cv;
#ifdef _DEBUG
#pragma comment ( lib,"opencv_core320d.lib")
#pragma comment ( lib,"opencv_highgui320d.lib")
#pragma comment ( lib,"opencv_calib3d320d.lib")
#pragma comment ( lib,"opencv_imgcodecs320d.lib")
#pragma comment ( lib,"opencv_imgproc320d.lib")
#pragma comment ( lib,"opencv_cudaimgproc320d.lib")
#pragma comment ( lib,"opencv_cudaarithm320d.lib")
#pragma comment ( lib,"cudart.lib")
#pragma comment ( lib,"swap_image.lib")//別忘了加庫
#else
#pragma comment ( lib,"opencv_core320.lib")
#pragma comment ( lib,"opencv_highgui320.lib")
#pragma comment ( lib,"opencv_calib3d320.lib")
#pragma comment ( lib,"opencv_imgcodecs320.lib")
#pragma comment ( lib,"opencv_imgproc320.lib")
#pragma comment ( lib,"opencv_cudaimgproc320.lib")
#pragma comment ( lib,"opencv_cudaarithm320.lib")
#pragma comment ( lib,"cudart.lib")
#pragma comment ( lib,"swap_image.lib")//別忘了加庫
#endif
//這個宣告很重要,呼叫靜態庫
extern "C" void swap_image(cuda::GpuMat src,cuda::GpuMat dst,int w,int h);
int main(int argc, char **argv)
{
    Mat src, dst;
    cuda::GpuMat cu_src, cu_dst;
    int h, w;
    //根據argv[1]讀入圖片資料，BGR格式讀進來
    src = imread(argv[1]);
    //檢測是否正確讀入
    if (src.data == NULL)
    {
        cout << "Read image error" << endl;
        return -1;
    }
    h = src.rows; w = src.cols;
    cout << "圖片高：" << h << ",圖片寬：" << w << endl;
    //上傳CPU影象資料到GPU，跟cudaMalloc和cudaMemcpy很像哦，其實upload裡面就是這麼寫的
    cu_src.upload(src);
    //申請GPU空間，也可以到函式裡申請，不管怎樣總要申請，要不然核心函式會爆掉哦
    cu_dst = cuda::GpuMat(h, w, CV_8UC3, Scalar(0, 0, 0));
    //申請CPU空間
    dst = Mat(h, w, CV_8UC3, Scalar(0, 0, 0));
    //呼叫函式swap_image,由該函式呼叫核心函式，這樣層次分明，不容易出錯
    //當然你也可以直接在這裡呼叫核心函式，東西太多程式碼容易亂
    swap_image(cu_src, cu_dst, h, w);
    //下載GPU資料到CPU，與upload()對應
    cu_dst.download(dst);
    //顯示cpu影象，如果安裝了openCV集成了openGL,那可以直接顯示GpuMat
    imshow("dst", dst);
    //等待按鍵
    waitKey();
    //寫圖片到檔案
    if (argc == 3)
        imwrite(argv[2], dst);
    return 0;
}
PS：以下為圖片效果，順便為我家贛南臍橙打個廣告   

相關推薦

cuda與openCV結合程式設計（一）

    學習計算機影象處理演算法的童鞋，就不得不學習cuda，為啥呢？因為影象處理一般都是矩陣運算，動不動就是百萬的計算量這個時候優化計算時間是必不可少的。openCV本身提供了很多cuda函式，能夠滿足大多數使用者的需求。但是也不絕對，有時候我們需要自己定義一個核心函式進

effective C++筆記--模板與泛型程式設計（一）

文章目錄 瞭解隱式介面和編譯器多型 瞭解typename的雙重意義 學習處理模板化基類內的名稱 將與引數無關的程式碼抽離template 瞭解隱式介面和編譯器多型 ．　面向物件程式設計世界總是以顯式介面和執行期多型解決問題。比

C++ STL標準庫與泛型程式設計（一）

泛型程式設計，就是使用模板為主要工具來編寫程式。其中沒有太多的面向物件的觀念，不涉及虛擬函式的使用。 使用C++標準庫 C++標準庫：以程式碼形式給出，放於各種標頭檔案( header files )內，經過編譯後才能使用。 所有新式的 headers 內的元件封裝於 namespace

Python多執行緒與多程序程式設計（一） 就這麼簡單

""" <axiner>宣告：（錯了另刂扌丁我） （如若有誤，請記得指出喲，謝謝了！！！） """ 先來了解一個概念，GIL？ GIL的全稱為Global Interpreter Lock, 全域性直譯器鎖。 Python程式碼的執行由Python 虛擬機器(也叫直譯器主

Python下opencv使用筆記（一）（圖像簡單讀取、顯示與儲存）

操作 灰度圖 清晰 ren from tty ims 圖像 type 寫在之前 從去年開始關註python這個軟件，途中間間斷斷看與學過一些關於python的東西。感覺python確實是一個簡單優美、easy上手的腳本編程語言，眾多的第三方庫使得py

利用CUDA進行GPU程式設計（一）

安裝CUDA工具 進行GPU程式設計的第一步，是對程式設計環境進行搭建，小編選擇的是NVIDIA提供的CUDA toolkit, 使用該工具的硬體基礎是電腦顯示卡需要是N卡（即NVIDIA系列顯示卡），通過在電腦中工作管理員的顯示介面卡中檢視自己電腦的顯示卡資訊，也可以在魯大師等軟體中直

併發程式設計（一）：執行緒基礎、執行緒之間的共享與協作

一、基礎概念 1.1 CPU核心數、執行緒數 **兩者的關係：**cpu的核心數與執行緒數是1:1的關係，例如一個8核的cpu，支援8個執行緒同時執行。但在intel引入超執行緒技術以後，cpu與執行緒數的關係就變成了1:2。此外在開發過程中並沒感覺到執行緒的限制，那是因為cpu

Python + OpenCV 學習筆記（一）>>> 影象的載入與儲存

影象屬性： 通道數目 高與寬 畫素資料 影象型別 顯示一張圖片： import cv2 as cv src = cv.imread("/home/pi/Desktop/apple.jpg") cv.namedWindow("input im

C語言面向物件程式設計（一）：封裝與繼承

最近在用 C 做專案，之前用慣了 C++ ，轉回頭來用C 還真有點不適應。 C++ 語言中自帶面向物件支援，如封裝、繼承、多型等面向物件的基本特徵。 C 原本是面向過程的語言，自身沒有內建這些特性，但我們還是可以利用 C 語言本身已有的特性來實現面向物件的一些基本特徵。接下來

Netty深入分析與Dubbo實戰解析（一）——網路程式設計模型介紹

Linux網路程式設計模型介紹 Linux核心將所有外部裝置都看作一個檔案來操作，對一個檔案的讀寫操作會呼叫核心提供的系統命令，返回一個file descriptor（fd，檔案描述符）。而對一個socket的讀寫也會有相應的描述符。描述符就是一個數字，它指向核心中的一個結構體（檔案路徑

linux系統程式設計（一）--時間與錯誤處理

#inlcude<error.h> errno perror(); //列印錯誤 strerror //列印錯誤 #inlcude<time.h> //使用CLOCKS_PER_SEC //常量可以得到秒數 clock_t \\ typedef

linux學習筆記之shell程式設計（一）正則表示式與字元處理

shell程式設計 基礎正則表示式 正則和萬用字元的區別：正則是包含匹配，匹配檔案內容，grep，awk等支援正則表示式。萬用字元是完全匹配，匹配檔名，例如find，ls不認識正則表示式 ####正則表示式常用的字元（注意區別於萬用字元裡面的符號）#### -*

併發程式設計（十一）—— Java 執行緒池 實現原理與原始碼深度解析（一）

史上最清晰的執行緒池原始碼分析 鼎鼎大名的執行緒池。不需要多說!!!!! 這篇部落格深入分析 Java 中執行緒池的實現。 總覽 下圖是 java 執行緒池幾個相關類的繼承結構：    先簡單說說這個繼承結構，Executor 位於最頂層，也是最簡單的，就一個 execute(

《GPU高效能程式設計CUDA實戰》學習筆記（一）

第一天讀這本書，先將封面讓大家看看吧 這兩張圖好大啊，不過讀者會更加清晰的看到作者資訊，這樣也不錯。 近年來英偉達在CUDA上float運算效能基本上已經超過了cpu，並且gpu（Graphics Procdss Unit）程式設計難度也接近cpu程式設計。NVID

網路程式設計（一）——淺析web伺服器與瀏覽器的實現原理

我們基本每天都在通過WEB瀏覽器，去瀏覽一些新聞，看看視訊之類的。眾所周知，這就是所謂的B/S結構（Browser/Server，瀏覽器/伺服器模式），是WEB興起後的一種網路結構模式，WEB瀏覽器是客戶端最主要的應用軟體。那順道就來簡單的看一下，所謂的Web伺服器(例如知名

OpenGL基礎圖形程式設計（一）OpenGL與3D圖形世界

一、OpenGL與3D圖形世界1.1、OpenGL使人們進入三維圖形世界　　我們生活在一個充滿三維物體的三維世界中，為了使計算機能精確地再現這些物體，我們必須能在三維空間描繪這些物體。我們又生活在一個充滿資訊的世界中，能否儘快地理解並運用這些資訊將直接影響事業的成敗，所以我

socket程式設計（一），實現伺服器與客戶端簡單通訊

本節主講客戶端向服務傳送資訊，伺服器轉發給客戶端，當然也可以稍微改一下程式碼可以實現互傳訊息，不再贅述。 難點在於伺服器端的程式碼思路： （1）主程式Socket socket=server.acc

[1.2]Spark core程式設計（一）之RDD總論與建立RDD的三種方式

參考 場景 RDD的理解 一、RDD是基於工作集的應用抽象;是分散式、函數語言程式設計的抽象。 MapReduce:基於資料集的處理。兩者的共同特徵：位置感知（具體資料在哪裡）、容錯、負載均衡。 基於資料集的處理：從物理儲存裝置上載入資料，然

Python與Qt UI的圖形介面程式設計（一）

本例項實現一個"Hello Kitty!"例子，簡單介紹Qt程式設計的基本流程，以及Qt程式的編繹執行方式,例項效果圖如圖所示。   這是一個簡單的例子，整個對話方塊只有一個按鈕，單擊該按鈕，對話方塊關閉，退出程式。 實現程式碼如下： from PyQt4.QtGu

Python核心原始碼解析與C/CPP-API拓展程式設計（一）PyObject

開發十年，就只剩下這套架構體系了！ >>>