之前給大家分享過(guò)opencv在jetson nano 2gb和ubuntu設(shè)備中使用并且展示了一些人臉識(shí)別等的小demo。但是對(duì)于圖像處理，使用gpu加速是很常見 .(以下概念介紹內(nèi)容來(lái)自百科和網(wǎng)絡(luò)其他博主文章)

GPU介紹(從GPU誕生之日起，GPU的設(shè)計(jì)邏輯與CPU的設(shè)計(jì)邏輯相差很多。GPU從誕生之日起，它的定位是3D圖形渲染設(shè)備。在設(shè)計(jì)GPU時(shí)從其功能出發(fā)，把更多的晶體管用于數(shù)據(jù)處理。這使得GPU相比CPU有更強(qiáng)的單精度浮點(diǎn)運(yùn)算能力。人們?yōu)榱顺浞掷肎PU的性能，使用了很多方法。這)加速處理是比較常見的。

在而GPU加速的軟件實(shí)現(xiàn)中我們可能會(huì)聽到 opengl opencl cuda這些名詞。下面再給大家分別介紹一下這幾種區(qū)別：

OpenGL（英語(yǔ)：Open Graphics Library，譯名：開放圖形庫(kù)或者“開放式圖形庫(kù)”）是用于渲染2D、3D矢量圖形的跨語(yǔ)言、跨平臺(tái)的應(yīng)用程序編程接口（API）。這個(gè)接口由近350個(gè)不同的函數(shù)調(diào)用組成，用來(lái)繪制從簡(jiǎn)單的圖形比特到復(fù)雜的三維景象。而另一種程序接口系統(tǒng)是僅用于Microsoft Windows上的Direct3D。OpenGL常用于CAD、虛擬現(xiàn)實(shí)、科學(xué)可視化程序和電子游戲開發(fā)。

OpenCl(是由蘋果（Apple）公司發(fā)起，業(yè)界眾多著名廠商共同制作的面向異構(gòu)系統(tǒng)通用目的并行編程的開放式、免費(fèi)標(biāo)準(zhǔn)，也是一個(gè)統(tǒng)一的編程環(huán)境。便于軟件開發(fā)人員為高性能計(jì)算服務(wù)器、桌面計(jì)算系統(tǒng)、手持設(shè)備編寫高效輕便的代碼，而且廣泛適用于多核心處理器(CPU)、圖形處理器(GPU)、Cell類型架構(gòu)以及數(shù)字信號(hào)處理器(DSP)等其他并行處理器，在游戲、娛樂(lè)、科研、醫(yī)療等各種領(lǐng)域都有廣闊的發(fā)展前景。)

從2007年以后，基于CUDA和OpenCL這些被設(shè)計(jì)成具有近似于高階語(yǔ)言的語(yǔ)法特性的新GPGPU語(yǔ)言，降低了人們使用GPGPU的難度，平緩了開始時(shí)的學(xué)習(xí)曲線。使得在GPGPU領(lǐng)域，OpenGL中的GLSL逐漸退出了人們的視線。來(lái)源：

簡(jiǎn)單說(shuō)OpenCL與OpenGL一樣，都是基于硬件API的編程。OpenGL是針對(duì)圖形的，而OpenCL則是針對(duì)并行計(jì)算的API，而針對(duì)并行計(jì)算下面還有一個(gè)cuda。

CUDA（Compute Unified Device Architecture，統(tǒng)一計(jì)算架構(gòu)）是由英偉達(dá)NVIDIA所推出的一種整合技術(shù)，是該公司對(duì)于GPGPU的正式名稱。透過(guò)這個(gè)技術(shù)，使用者可利用NVIDIA的GeForce 8以后的GPU和較新的Quadro GPU進(jìn)行計(jì)算。亦是首次可以利用GPU作為C-編譯器的開發(fā)環(huán)境。NVIDIA行銷的時(shí)候，往往將編譯器與架構(gòu)混合推廣，造成混亂。實(shí)際上，CUDA可以相容OpenCL或者自家的C-編譯器。無(wú)論是CUDA C-語(yǔ)言或是OpenCL，指令最終都會(huì)被驅(qū)動(dòng)程式轉(zhuǎn)換成PTX代碼，交由顯示核心計(jì)算。

通俗的介紹，cuda是nvidia公司的生態(tài)，它前者是配備完整工具包、針對(duì)單一供應(yīng)商(NVIDIA)的成熟的開發(fā)平臺(tái)，opencl是一個(gè)開源的標(biāo)準(zhǔn)。

CUDA是NVIDIA GPU編程語(yǔ)言， OpenCL是異構(gòu)計(jì)算庫(kù)。CUDA和C++雖然都可以用nvcc編譯，但C++只能在CPU上跑，CUDA只能在GPU上跑；而OpenCL并不局限于某個(gè)計(jì)算設(shè)備，旨在將同樣的任務(wù)通過(guò)其提供的抽象接口在多種硬件上運(yùn)行（CPU，GPU，FPGA，etc）

跨平臺(tái)性和通用性上OpenCL占有很大優(yōu)勢(shì)（這也是很多National Laboratory使用OpenCL進(jìn)行科學(xué)計(jì)算的最主要原因）。OpenCL支持包括ATI,NVIDIA,Intel,ARM在內(nèi)的多類處理器，并能支持運(yùn)行在CPU的并行代碼，同時(shí)還獨(dú)有Task-Parallel Execution Mode，能夠更好的支持Heterogeneous Computing。這一點(diǎn)是僅僅支持?jǐn)?shù)據(jù)級(jí)并行并僅能在NVIDIA眾核處理器上運(yùn)行的CUDA無(wú)法做到的。

在開發(fā)者友好程度CUDA在這方面顯然受更多開發(fā)者青睞。原因在于其統(tǒng)一的開發(fā)套件(CUDA Toolkit, NVIDIA GPU Computing SDK以及NSight等等)、非常豐富的庫(kù)(cuFFT, cuBLAS, cuSPARSE, cuRAND, NPP, Thrust)以及NVCC(NVIDIA的CUDA編譯器)所具備的PTX(一種SSA中間表示，為不同的NVIDIA GPU設(shè)備提供一套統(tǒng)一的靜態(tài)ISA)代碼生成、離線編譯等更成熟的編譯器特性。相比之下，使用OpenCL進(jìn)行開發(fā)，只有AMD對(duì)OpenCL的驅(qū)動(dòng)相對(duì)成熟。?來(lái)源

今天的主題來(lái)自介紹cuda的使用，前面這部分概念性的介紹（來(lái)源都是網(wǎng)絡(luò)，只有一小部分是自己寫的）只是幫助大家好理解今天要使用的cuda，那么開始進(jìn)入正題。（自己本身的電腦帶NVIDIA的顯卡，以及手里面有jetson nano的NVIDIA板卡，所以才使用cuda，其他朋友使用請(qǐng)注意自己手中的硬件是否有NVIDIA的顯卡）。

接下來(lái)，我大致分為三個(gè)部分介紹，一、cuda并行計(jì)算原理介紹 ，二、cuda環(huán)境安裝和直接使用，三、opencv中cuda安裝和使用

歡迎關(guān)注微信公眾號(hào)：羽林君，或者添加作者個(gè)人微信：become_me

cuda并行計(jì)算原理介紹

GPU背景介紹

GPU里有很多Compute Unit(計(jì)算單元), 這些單元是由專門的處理邏輯, 很多register, 和L1 cache 組成. Memory access subsystem 把GPU 和RAM 連起來(lái)(通常是個(gè)L2 cache). Threads 以SIMT的形式運(yùn)行: 多個(gè)thread共享一個(gè)instruction unit. 對(duì)NV GPU來(lái)說(shuō), 32 thread 組成一個(gè)warps, 對(duì)AMD GPU 來(lái)說(shuō), 64 threads 叫做一個(gè)wave fronts. 本文只使用warp的定義. 想達(dá)到最高速度, 一定要討論SIMT, 因?yàn)樗绊懼鴐emory access 也會(huì)造成code的序列化(serialization, parallelism的反面)

kernel function里會(huì)指出哪些code由一條thread處理, host program會(huì)決定多少條thread來(lái)處理一個(gè)kernel. 一個(gè)work group 里的thread可以通過(guò)barrier synchronization, 共享L1 cache 來(lái)互相協(xié)力. 再由compute unit處理這些work group. 能同時(shí)被conpute unit 運(yùn)行的work group數(shù)量有限, 所以很多得等到其他完成之后再運(yùn)行. 需要處理的work group的size和數(shù)量還有最大并行數(shù)量被稱為kernel的execution configuration(運(yùn)行配置).

kernel 的占用率就是指同時(shí)運(yùn)行的thread數(shù)量除以最大數(shù)量. 傳統(tǒng)建議就是提升這個(gè)占用率來(lái)獲得更好的性能, 但也有一些其他因素比如ILP, MLP, instruction latencies, 也很關(guān)鍵.內(nèi)容來(lái)源

GPU架構(gòu)特點(diǎn)

首先我們先談一談串行計(jì)算和并行計(jì)算。我們知道，高性能計(jì)算的關(guān)鍵利用多核處理器進(jìn)行并行計(jì)算。

當(dāng)我們求解一個(gè)計(jì)算機(jī)程序任務(wù)時(shí)，我們很自然的想法就是將該任務(wù)分解成一系列小任務(wù)，把這些小任務(wù)一一完成。在串行計(jì)算時(shí)，我們的想法就是讓我們的處理器每次處理一個(gè)計(jì)算任務(wù)，處理完一個(gè)計(jì)算任務(wù)后再計(jì)算下一個(gè)任務(wù)，直到所有小任務(wù)都完成了，那么這個(gè)大的程序任務(wù)也就完成了。

串行計(jì)算的缺點(diǎn)非常明顯，如果我們擁有多核處理器，我們可以利用多核處理器同時(shí)處理多個(gè)任務(wù)時(shí)，而且這些小任務(wù)并沒(méi)有關(guān)聯(lián)關(guān)系（不需要相互依賴，比如我的計(jì)算任務(wù)不需要用到你的計(jì)算結(jié)果），那我們?yōu)槭裁催€要使用串行編程呢？為了進(jìn)一步加快大任務(wù)的計(jì)算速度，我們可以把一些獨(dú)立的模塊分配到不同的處理器上進(jìn)行同時(shí)計(jì)算（這就是并行），最后再將這些結(jié)果進(jìn)行整合，完成一次任務(wù)計(jì)算。下圖就是將一個(gè)大的計(jì)算任務(wù)分解為小任務(wù)，然后將獨(dú)立的小任務(wù)分配到不同處理器進(jìn)行并行計(jì)算，最后再通過(guò)串行程序把結(jié)果匯總完成這次的總的計(jì)算任務(wù)。

所以，一個(gè)程序可不可以進(jìn)行并行計(jì)算，關(guān)鍵就在于我們要分析出該程序可以拆分出哪幾個(gè)執(zhí)行模塊，這些執(zhí)行模塊哪些是獨(dú)立的，哪些又是強(qiáng)依賴強(qiáng)耦合的，獨(dú)立的模塊我們可以試著設(shè)計(jì)并行計(jì)算，充分利用多核處理器的優(yōu)勢(shì)進(jìn)一步加速我們的計(jì)算任務(wù)，強(qiáng)耦合模塊我們就使用串行編程，利用串行+并行的編程思路完成一次高性能計(jì)算。

接下來(lái)我們談?wù)凜PU和GPU有什么區(qū)別，他們倆各自有什么特點(diǎn)，我們?cè)谡劜⑿?、串行?jì)算時(shí)多次談到“多核”的概念，現(xiàn)在我們先從“核”的角度開始這個(gè)話題。首先CPU是專為順序串行處理而優(yōu)化的幾個(gè)核心組成。而GPU則由數(shù)以千計(jì)的更小、更高效的核心組成，這些核心專門為同時(shí)處理多任務(wù)而設(shè)計(jì)，可高效地處理并行任務(wù)。也就是，CPU雖然每個(gè)核心自身能力極強(qiáng)，處理任務(wù)上非常強(qiáng)悍，無(wú)奈他核心少，在并行計(jì)算上表現(xiàn)不佳；反觀GPU，雖然他的每個(gè)核心的計(jì)算能力不算強(qiáng)，但他勝在核心非常多，可以同時(shí)處理多個(gè)計(jì)算任務(wù)，在并行計(jì)算的支持上做得很好。

GPU和CPU的不同硬件特點(diǎn)決定了他們的應(yīng)用場(chǎng)景，CPU是計(jì)算機(jī)的運(yùn)算和控制的核心，GPU主要用作圖形圖像處理。圖像在計(jì)算機(jī)呈現(xiàn)的形式就是矩陣，我們對(duì)圖像的處理其實(shí)就是操作各種矩陣進(jìn)行計(jì)算，而很多矩陣的運(yùn)算其實(shí)可以做并行化，這使得圖像處理可以做得很快，因此GPU在圖形圖像領(lǐng)域也有了大展拳腳的機(jī)會(huì)。下圖表示的就是一個(gè)多GPU計(jì)算機(jī)硬件系統(tǒng)，可以看出，一個(gè)GPU內(nèi)存就有很多個(gè)SP和各類內(nèi)存，這些硬件都是GPU進(jìn)行高效并行計(jì)算的基礎(chǔ)。

現(xiàn)在再?gòu)臄?shù)據(jù)處理的角度來(lái)對(duì)比CPU和GPU的特點(diǎn)。CPU需要很強(qiáng)的通用性來(lái)處理各種不同的數(shù)據(jù)類型，比如整型、浮點(diǎn)數(shù)等，同時(shí)它又必須擅長(zhǎng)處理邏輯判斷所導(dǎo)致的大量分支跳轉(zhuǎn)和中斷處理，所以CPU其實(shí)就是一個(gè)能力很強(qiáng)的伙計(jì)，他能把很多事處理得妥妥當(dāng)當(dāng)，當(dāng)然啦我們需要給他很多資源供他使用（各種硬件），這也導(dǎo)致了CPU不可能有太多核心（核心總數(shù)不超過(guò)16）。而GPU面對(duì)的則是類型高度統(tǒng)一的、相互無(wú)依賴的大規(guī)模數(shù)據(jù)和不需要被打斷的純凈的計(jì)算環(huán)境，GPU有非常多核心（費(fèi)米架構(gòu)就有512核），雖然其核心的能力遠(yuǎn)沒(méi)有CPU的核心強(qiáng)，但是勝在多， 在處理簡(jiǎn)單計(jì)算任務(wù)時(shí)呈現(xiàn)出“人多力量大”的優(yōu)勢(shì)，這就是并行計(jì)算的魅力。內(nèi)容來(lái)源

整理一下兩者特點(diǎn)就是：

• 1.CPU：擅長(zhǎng)流程控制和邏輯處理，不規(guī)則數(shù)據(jù)結(jié)構(gòu)，不可預(yù)測(cè)存儲(chǔ)結(jié)構(gòu)，單線程程序，分支密集型算法2.GPU：擅長(zhǎng)數(shù)據(jù)并行計(jì)算，規(guī)則數(shù)據(jù)結(jié)構(gòu)，可預(yù)測(cè)存儲(chǔ)模式

CUDA存儲(chǔ)器類型：
每個(gè)線程擁有自己的?register寄存器?and?loacal?memory?局部?jī)?nèi)存
每個(gè)線程塊擁有一塊?shared?memory?共享內(nèi)存
所有線程都可以訪問(wèn)?global?memory?全局內(nèi)存

還有，可以被所有線程訪問(wèn)的
?????只讀存儲(chǔ)器：
?????constant?memory?(常量?jī)?nèi)容)?and?texture?memory
????
a.?寄存器Register
???寄存器是GPU上的高速緩存器，其基本單元是寄存器文件，每個(gè)寄存器文件大小為32bit.
?? Kernel中的局部(簡(jiǎn)單類型)變量第一選擇是被分配到Register中。
???????特點(diǎn)：每個(gè)線程私有，速度快。

b.?局部存儲(chǔ)器?local?memory

???當(dāng)register耗盡時(shí)，數(shù)據(jù)將被存儲(chǔ)到local?memory。
???如果每個(gè)線程中使用了過(guò)多的寄存器，或聲明了大型結(jié)構(gòu)體或數(shù)組，
???或編譯器無(wú)法確定數(shù)組大小，線程的私有數(shù)據(jù)就會(huì)被分配到local?memory中。
???
???????特點(diǎn)：每個(gè)線程私有；沒(méi)有緩存，慢。
???????注：在聲明局部變量時(shí)，盡量使變量可以分配到register。如：
???????unsigned?int?mt[3];
???????改為：unsigned int mt0, mt1, mt2;

c.?共享存儲(chǔ)器?shared?memory
???????可以被同一block中的所有線程讀寫
???????特點(diǎn)：block中的線程共有；訪問(wèn)共享存儲(chǔ)器幾乎與register一樣快.

d.?全局存儲(chǔ)器?global?memory
??特點(diǎn)：所有線程都可以訪問(wèn)；沒(méi)有緩存

e.?常數(shù)存儲(chǔ)器constant?memory
???用于存儲(chǔ)訪問(wèn)頻繁的只讀參數(shù)
???特點(diǎn)：只讀；有緩存；空間小(64KB)
???注：定義常數(shù)存儲(chǔ)器時(shí)，需要將其定義在所有函數(shù)之外，作用于整個(gè)文件

f.?紋理存儲(chǔ)器?texture?memory
???????是一種只讀存儲(chǔ)器，其中的數(shù)據(jù)以一維、二維或者三維數(shù)組的形式存儲(chǔ)在顯存中。
???在通用計(jì)算中，其適合實(shí)現(xiàn)圖像處理和查找，對(duì)大量數(shù)據(jù)的隨機(jī)訪問(wèn)和非對(duì)齊訪問(wèn)也有良好的加速效果。
???????特點(diǎn)：具有紋理緩存，只讀。
threadIdx，blockIdx,?blockDim,?gridDim之間的區(qū)別與聯(lián)系
?在啟動(dòng)kernel的時(shí)候，要通過(guò)指定gridsize和blocksize才行
?dim3?gridsize(2,2);???//?2行*2列*1頁(yè)?形狀的線程格，也就是說(shuō)?4個(gè)線程塊
??? gridDim.x，gridDim.y，gridDim.z相當(dāng)于這個(gè)dim3的x，y，z方向的維度，這里是2*2*1。
序號(hào)從0到3，且是從上到下的順序，就是說(shuō)是下面的情況:
?具體到?線程格?中每一個(gè)?線程塊的 id索引為：
?
??? grid 中的 blockidx 序號(hào)標(biāo)注情況為：??????0???? 2 
??????????????????????????????????????????1?????3
???????????
???dim3?blocksize(4,4);??//?線程塊的形狀，4行*4列*1頁(yè)，一個(gè)線程塊內(nèi)部共有?16個(gè)線程
?
???blockDim.x，blockDim.y，blockDim.z相當(dāng)于這個(gè)dim3的x，y，z方向的維度，
???這里是4*4*1.序號(hào)是0-15，也是從上到下的標(biāo)注：

block中的?threadidx?序號(hào)標(biāo)注情況???0???????4???????8??????12?
??????1???????5???????9??????13
??????2???????6???????10?????14
??????3???????7???????11?????15
1.??1維格子，1維線程塊，N個(gè)線程======

?實(shí)際的線程id?tid?=??blockidx.x?*?blockDim.x?+?threadidx.x
?塊id???0?1?2?3?
?線程id?0?1?2?3?4
2.?1維格子，2D維線程塊
?塊id???1?2?3?
?線程id??0??2
????????1??3
?????????塊id???????????塊線程總數(shù)
?實(shí)際的線程id?tid?=??blockidx.x?*?blockDim.x?*?blockDim.y?+?
?????????當(dāng)前線程行數(shù)????每行線程數(shù)
??????threadidx.y?*?blockDim.x???+?
?????????當(dāng)前線程列數(shù)
??????threadidx.x

更加詳細(xì)信息，大家可以參考這兩篇文章：

https://www.cnblogs.com/skyfsm/p/9673960.html

https://github.com/Ewenwan/ShiYanLou/tree/master/CUDA

安裝cuda使用環(huán)境

cuda使用中 首先我們直接安裝CUDA toolkit （cuDNN是用于配置深度學(xué)習(xí)使用，本次我沒(méi)有使用到，就是下載cuDNN的庫(kù)放置到對(duì)應(yīng)的鏈接庫(kù)目錄，大家可以自行去搜索安裝），然后直接調(diào)用cuda的api進(jìn)行一些計(jì)算的測(cè)試，還有是可以安裝opencv中的cuda接口，通過(guò)opencv中的cuda使用。

直接使用cuda

直接安裝CUDA ，首先需要下載安裝包

CUDA toolkit（toolkit就是指工具包）cuDNN 注：cuDNN 是用于配置深度學(xué)習(xí)使用（這次我沒(méi)有使用就沒(méi)有下載）

官方教程

CUDA：Installation Guide Windows :: CUDA Toolkit Documentation（鏈接）

注：按照自己版本去查看資料，打開鏈接后，右上角處有一個(gè)older選項(xiàng)，大家可以選擇自己對(duì)應(yīng)的版本文檔。

cuDNN：Installation Guide :: NVIDIA Deep Learning cuDNN Documentation（鏈接）

安裝cuda要注意一個(gè)版本匹配問(wèn)題，就是我們不同的顯卡是對(duì)應(yīng)不同的驅(qū)動(dòng)。所以第一步就是查看顯卡驅(qū)動(dòng)版本，我使用的是Ubuntu，和window安裝有些區(qū)別。第一如果我們安裝過(guò)自己的顯卡驅(qū)動(dòng)，使用 nvidia-smi命令查看支持的cuda版本

如果沒(méi)有安裝可以選擇ubuntu的bash界面搜索附加驅(qū)動(dòng)，然后進(jìn)行安裝。參考

當(dāng)然由于ubuntu本身也有不同的版本支持，大家也可以直接到驅(qū)動(dòng)的官網(wǎng)進(jìn)行查找，里面有專門的ubuntu20適配版本選項(xiàng)，其中由于我使用的是ubuntu20，而CUDA11.0以上版本才支持Ubuntu20.04，所以直接就選擇了最新的。

官方地址：https://developer.nvidia.com/cuda-toolkit-archive

選擇CUDA Toolkit 11.6.0之后，里面有幾種選項(xiàng)，分別是本地deb安裝、網(wǎng)絡(luò)下載安裝和可執(zhí)行文件安裝，我因?yàn)榫W(wǎng)絡(luò)不太順暢原因選擇了第一個(gè)選項(xiàng)，最后下載下載下來(lái)的deb文件有2.7G，大家做好下載大文件的準(zhǔn)備。

wget?https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2004/x86_64/cuda-ubuntu2004.pin
sudo?mv?cuda-ubuntu2004.pin?/etc/apt/preferences.d/cuda-repository-pin-600
wget?https://developer.download.nvidia.com/compute/cuda/11.6.0/local_installers/cuda-repo-ubuntu2004-11-6-local_11.6.0-510.39.01-1_amd64.deb
sudo?dpkg?-i?cuda-repo-ubuntu2004-11-6-local_11.6.0-510.39.01-1_amd64.deb
sudo?apt-key?add?/var/cuda-repo-ubuntu2004-11-6-local/7fa2af80.pub
sudo?apt-get?update
sudo?apt-get?-y?install?cuda

這時(shí)候cuda就安裝好了，接下來(lái)就是使用了。

大家可以使用cuda命令，tab就可以出現(xiàn)對(duì)應(yīng)的命令，或者使用nvcc編譯對(duì)應(yīng)的cuda的.cu文件進(jìn)行驗(yàn)證。

大家可以通過(guò)以下鏈接里面的文檔指導(dǎo)寫符合自己需要的功能：https://docs.nvidia.com/cuda/archive/11.6.0/cuda-c-programming-guide/index.html

下面官網(wǎng)API文件目錄：

Versioned Online Documentation

官方示例：

接下來(lái)開始進(jìn)行cuda的使用：輸出hello world

#include?"cuda_runtime.h"
#include?"device_launch_parameters.h"
#include?<stdio.h>
#include?<iostream>
using?namespace?std;

__global__?void?kernel(void)?{?//帶有了__global__這個(gè)標(biāo)簽，表示這個(gè)函數(shù)是在GPU上運(yùn)行
??printf("hello?world?gpu?n");
}
int?main()?{
? kernel<<<1, 1>>>();//調(diào)用除了常規(guī)的參數(shù)之外，還增加了<<<>>>修飾,其中第一個(gè)1，代表線程格里只有一個(gè)線程塊；第二個(gè)1，代表一個(gè)線程塊里只有一個(gè)線程。
????cudaError_t?cudaStatus;

????cudaStatus?=?cudaDeviceSynchronize();
????if?(cudaStatus?!=?cudaSuccess)?{
????????fprintf(stderr,?"cudaDeviceSynchronize?returned?error?code?%d?after?launching?addKernel!n",?cudaStatus);
????}

??return?0;
}

這里面記得使用cudaDeviceSynchronize函數(shù)，等待GPU完成運(yùn)算。如果不加的話，你將看不到printf的輸出

編譯選項(xiàng) nvcc test_cuda_hello.cu

這個(gè)程序中GPU調(diào)用的函數(shù)和普通的C程序函數(shù)的區(qū)別

函數(shù)的調(diào)用除了常規(guī)的參數(shù)之外，還增加了<<<>>>修飾。調(diào)用通過(guò)<<<參數(shù)1,參數(shù)2>>>，用于說(shuō)明內(nèi)核函數(shù)中的線程數(shù)量，以及線程是如何組織的。以線程格（Grid）的形式組織，每個(gè) 線程格 由若干個(gè) 線程塊（block）組成，而每個(gè)線程塊又由若干個(gè)線程（thread）組成。是以block為單位執(zhí)行的。

接下來(lái)，我們?cè)賮?lái)使用一個(gè)demo，使用一個(gè)簡(jiǎn)單的計(jì)算，將其放置于GPU進(jìn)行計(jì)算

//?輸入變量?以指針?lè)绞絺鬟f===================================
#include?<iostream>
#include?<cuda.h>
#include?<cuda_runtime.h>
#include?<stdio.h>

//?輸入變量?全部為指針??類似數(shù)據(jù)的引用
__global__?void?gpuAdd(int?*d_a,?int?*d_b,?int?*d_c)?
{
?*d_c?=?*d_a?+?*d_b;
}

int?main(void)?
{
?//?CPU變量
?int?h_a,h_b,?h_c;
?//?CPU??指針?變量?指向?GPU數(shù)據(jù)地址
?int?*d_a,*d_b,*d_c;
?//?初始化CPU變量
?h_a?=?1;
?h_b?=?4;
??
?//?分配GPU?變量?jī)?nèi)存
?cudaMalloc((void**)&d_a,?sizeof(int));
?cudaMalloc((void**)&d_b,?sizeof(int));
?cudaMalloc((void**)&d_c,?sizeof(int));
??
?//?輸入變量?CPU?拷貝到?GPU???右?到?左
?cudaMemcpy(d_a,?&h_a,?sizeof(int),?cudaMemcpyHostToDevice);
?cudaMemcpy(d_b,?&h_b,?sizeof(int),?cudaMemcpyHostToDevice);
??
?//?調(diào)用核函數(shù)
?gpuAdd?<<?<1,?1?>>?>?(d_a,?d_b,?d_c);
??
?//?拷貝GPU數(shù)據(jù)結(jié)果?d_c?到?CPU變量
?cudaMemcpy(&h_c,?d_c,?sizeof(int),?cudaMemcpyDeviceToHost);
?printf("Passing?Parameter?by?Reference?Output:?%d?+?%d?=?%dn",?h_a,?h_b,?h_c);
??
?//?清理GPU內(nèi)存?Free?up?memory?
?cudaFree(d_a);
?cudaFree(d_b);
?cudaFree(d_c);
?return?0;
}

CUDA代碼中比較重要的函數(shù):

cudaMalloc與C語(yǔ)言中的malloc函數(shù)一樣，只是此函數(shù)在GPU的內(nèi)存你分配內(nèi)存。

addKernel<<<1, size>>>

這里就涉及了GPU和主機(jī)之間的內(nèi)存交換了，cudaMalloc是在GPU的內(nèi)存里開辟一片空間，
然后通過(guò)操作之后，這個(gè)內(nèi)存里有了計(jì)算出來(lái)內(nèi)容，再通過(guò)cudaMemcpy這個(gè)函數(shù)把內(nèi)容從GPU復(fù)制出來(lái)。可以在設(shè)備代碼中使用cudaMalloc()分配的指針進(jìn)行設(shè)備內(nèi)存讀寫操作；cudaMalloc()分配的指針傳遞給在設(shè)備上執(zhí)行的函數(shù)；但是不可以在主機(jī)代碼中使用cudaMalloc()分配的指針進(jìn)行主機(jī)內(nèi)存讀寫操作（即不能進(jìn)行解引用）。cudaFree與c語(yǔ)言中的free()函數(shù)一樣，只是此函數(shù)釋放的是cudaMalloc()分配的內(nèi)存。

cudaMemcpy與c語(yǔ)言中的memcpy函數(shù)一樣，只是此函數(shù)可以在主機(jī)內(nèi)存和GPU內(nèi)存之間互相拷貝數(shù)據(jù)。此外與C中的memcpy()一樣，以同步方式執(zhí)行，即當(dāng)函數(shù)返回時(shí)，復(fù)制操作就已經(jīng)完成了，并且在輸出緩沖區(qū)中包含了復(fù)制進(jìn)去的內(nèi)容。相應(yīng)的有個(gè)異步方式執(zhí)行的函數(shù)cudaMemcpyAsync().

該函數(shù)第一個(gè)參數(shù)是目的指針，第二個(gè)參數(shù)是源指針，第三個(gè)參數(shù)是復(fù)制內(nèi)存的大小，第四個(gè)參數(shù)告訴運(yùn)行時(shí)源指針，這是一個(gè)什么類型的指針，即把內(nèi)存從哪里復(fù)制到哪里。第四個(gè)參數(shù)可以選用以下的形式：參考

cudaMemcpyHostToDevice：從主機(jī)復(fù)制到設(shè)備；cudaMemcpyDeviceToHost：從設(shè)備復(fù)制到主機(jī)；cudaMemcpyDeviceToDevice：從設(shè)備復(fù)制到設(shè)備；cudaMemcpyHostToHost：從主機(jī)復(fù)制到主機(jī)；

調(diào)用的核函數(shù)前面已經(jīng)介紹過(guò)，此處不再贅述。還有其他函數(shù)大家可以自行參考官網(wǎng)的API介紹。

大家也可以看《GPU高性能編程CUDA實(shí)戰(zhàn)》這本書

opencv中cuda庫(kù)安裝與使用

除了直接使用cuda，opencv也有相應(yīng)的cuda的庫(kù)，調(diào)用cuda從而在opencv實(shí)現(xiàn)gpu加速功能。這就需要安裝CUDA，并從OpenCV編譯安裝時(shí)候安裝好對(duì)應(yīng)cv::cuda庫(kù)。

首先要使用cv::cuda里面對(duì)應(yīng)的cv::cuda::add 、cv::cuda::multiply，我們需要安裝opencv要配合OpenCV Contrib庫(kù)。

OpenCV Contrib庫(kù)是非官方的第三方開發(fā)擴(kuò)充庫(kù)。通過(guò)這個(gè)庫(kù)，我們能使用如dnn、相機(jī)標(biāo)注、3D成像、ArUco、物體追蹤，甚至是需付費(fèi)的SURF、SIFT特征點(diǎn)提取算法

安裝OpenCV Contrib也是需要和opencv進(jìn)行版本匹配的。大家可以在下面位置進(jìn)行對(duì)應(yīng)版本下載：

contrib庫(kù)：https://github.com/opencv/opencv_contrib/tagsopencv版本：https://opencv.org/releases.html

對(duì)與我個(gè)人而言，opencv版本是 4.4.5

$?pkg-config?opencv4?--modversion?
4.4.5

下載好對(duì)應(yīng)的opencv和opencv_contrib庫(kù)，開始進(jìn)行編譯，和之前單獨(dú)opencv編譯類似，解壓下載好的文件，然后在opencv建立一個(gè)build目錄，在里面進(jìn)行cmake 編譯。

在這里給大家在注解一些相關(guān)選項(xiàng)的含義：

?cmake?-D?CMAKE_BUILD_TYPE=Release
????-D?ENABLE_CXX11=ON
??????????-D?CMAKE_INSTALL_PREFIX=/usr/local
??????????-D?WITH_CUDA=ON
???????????-D?CUDA_ARCH_BIN=${cuda_compute}
???????????-D?CUDA_ARCH_PTX=""
???????????-D?ENABLE_FAST_MATH=ON
???????????-D?CUDA_FAST_MATH=ON
???????????-D?WITH_CUBLAS=ON
????????????-D?WITH_LIBV4L=ON
????????????-D?WITH_GSTREAMER=ON
?????????????-D?WITH_GSTREAMER_0_10=OFF
?????????????-D?WITH_QT=ON
?????????????-D?WITH_OPENGL=ON
?????????????-D?CUDA_NVCC_FLAGS="--expt-relaxed-constexpr"?
?????????????-D?CUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-11.6
?????????????-D?WITH_TBB=ON
?????????????-D?OPENCV_EXTRA_MODULES_PATH=/home/lyn/Documents/application/opencv_contrib-4.5.5/modules??????????../

CUDA_ARCH_BIN=${cuda_compute} 顯卡算力，大家可以對(duì)照我最后給算力查詢自行查詢，也可以進(jìn)行系統(tǒng)自己查找匹配版本。

WITH_QT=ON 這是是QT的編譯選項(xiàng) 大家可以按照自己需要選擇是否用QT進(jìn)行編譯開與關(guān)。

CUDA_TOOLKIT_ROOT_DIR 這是自己電腦安裝的cuda 版本

OPENCV_EXTRA_MODULES_PATH=/home/lyn/Documents/application/opencv_contrib-4.5.5/modules 這是解壓好后的opencv_contrib目錄

安裝時(shí)候可能會(huì)遇到如下問(wèn)題：

In?file?included?from?/home/lyn/Documents/application/opencv-4.5.5/build/modules/python_bindings_generator/pyopencv_custom_headers.h:7,
?????????????????from?/home/lyn/Documents/application/opencv-4.5.5/modules/python/src2/cv2.cpp:88:
/home/lyn/Documents/application/opencv_contrib-4.5.5/modules/phase_unwrapping/misc/python/pyopencv_phase_unwrapping.hpp:2:13:?error:?‘phase_unwrapping’?in?namespace?‘cv’?does?not?name?a?type
????2?|?typedef?cv::phase_unwrapping::HistogramPhaseUnwrapping::Params?HistogramPhaseUnwrapping_Params;

打開-D ENABLE_CXX11=ON 打開C++11的選項(xiàng) （這就是上面命令使用C++11的原因）

最后編譯成功：

sudo make install這個(gè)時(shí)候opencv中的cuda庫(kù)就安裝好了

介紹一個(gè)簡(jiǎn)單例子：

CMakeLists.txt文件

#?聲明要求的?cmake?最低版本
cmake_minimum_required(?VERSION?2.8?)
#?聲明一個(gè)?cmake?工程
project(opencv_cuda)
#?設(shè)置編譯模式
set(?CMAKE_BUILD_TYPE?"Debug"?)
set(?CMAKE_CXX_FLAGS?"-std=c++11")
#添加OPENCV庫(kù)
find_package(OpenCV?REQUIRED)
#添加OpenCV頭文件
include_directories(${OpenCV_INCLUDE_DIRS})
#顯示OpenCV_INCLUDE_DIRS的值
message(${OpenCV_INCLUDE_DIRS})
add_executable(target?test_cuda_cv.cpp)
#?將庫(kù)文件鏈接到可執(zhí)行程序上
target_link_libraries(target??${OpenCV_LIBS}?)?

test_cuda_cv.cpp文件

#include?<iostream>
#include?"opencv2/opencv.hpp"
#include?<opencv2/core/cuda.hpp>?
#include?<opencv2/cudaarithm.hpp>
#include?<opencv2/core/version.hpp>

int?main?(int?argc,?char*?argv[])
{
????//Read?Two?Images?
????cv::Mat?h_img1?=?cv::imread(?"/home/lyn/Documents/work-data/test_code/opencv/learn_code/"
??????"opencv_tutorial_data-master/images/sp_noise.png");
????cv::Mat?h_img2?=?cv::imread(?"/home/lyn/Documents/work-data/test_code/opencv/learn_code/"
??????"opencv_tutorial_data-master/images/sp_noise.png");??????
????//?cv::Mat?h_img2?=?cv::imread("/home/lyn/Documents/work-data/test_code/opencv/learn_code/c++/build/filter_line.png");
??
??//?cv::namedWindow(?"1",?cv::WINDOW_FREERATIO);?
??//?cv::imshow("1",h_img1);??

??//?cv::namedWindow(?"2",?cv::WINDOW_FREERATIO);?
??//?cv::imshow("2",h_img2);??

????cv::Mat?h_result1;


????#if?0
????//?cv::add(h_img1,?h_img2,h_result1);//加法操作?api
????cv::multiply(h_img1,cv::Scalar(2,?2,?2),h_result1);//乘法操作?api

????#else
????//?定義GPU數(shù)據(jù)
????cv::cuda::GpuMat?d_result1,d_img1,?d_img2;
????
????//?CPU?到?GPU???
????d_img1.upload(h_img1);
????d_img2.upload(h_img2);
????
????//?調(diào)用GPU接執(zhí)行?mat?substract
????//?cv::cuda::add(d_img1,?d_img2,d_result1);
????cv::cuda::multiply(d_img1,cv::Scalar(2,?2,?2),d_result1);//乘法操作?api

????//?gpu?結(jié)果?到?cpu
????d_result1.download(h_result1);

????#endif?
?

????//?顯示，保存
??//?cv::namedWindow(?"Image1",?cv::WINDOW_FREERATIO);?
????cv::imshow("Image1?",?h_img1);

??//?cv::namedWindow(?"Image2",?cv::WINDOW_FREERATIO);?
????cv::imshow("Image2?",?h_img2);

??//?cv::namedWindow(?"Result_Subtraction",?cv::WINDOW_FREERATIO);?
????cv::imshow("Result_Subtraction?",?h_result1);
????cv::imwrite("/home/lyn/Documents/work-data/test_code/opencv/learn_code/c++/result_add.png",?h_result1);
????cv::waitKey();
????return?0;
}

大家也可以使用opencv中的其他cuda庫(kù)，具體函數(shù)使用在下面

官方鏈接：https://docs.opencv.org/4.5.5/d1/d1e/group__cuda.html

大家可以在當(dāng)前頁(yè)面的版本選擇中，選擇和自己一致的opencv版本，進(jìn)行參考：

例如我查詢cv::cuda::add函數(shù)，就在 Operations on Matrices -> Per-element Operations

也可以看具體的數(shù)據(jù)結(jié)構(gòu)含義：例如 cv::cuda::GpuMat

官方的解釋是：具有引用計(jì)數(shù)的 GPU 內(nèi)存的基本存儲(chǔ)類。GpuMat類似于Mat，不過(guò)管理的是GPU內(nèi)存，也就是顯存。GpuMat和Mat相比，存在以下限制：

• 不支持任意尺寸（只支持2D，也就是二維數(shù)據(jù)）沒(méi)有返回對(duì)其數(shù)據(jù)的引用的函數(shù)（因?yàn)?GPU 上的引用對(duì) CPU 無(wú)效）不支持c++的模板技術(shù)。

總之是很方便的，大家可以按照自己需要進(jìn)行查詢信息。

附錄：這是官網(wǎng)提供算力表，大家可以參考選擇自己的要求算力的顯卡。

https://developer.nvidia.com/zh-cn/cuda-gpus

結(jié)語(yǔ)

這就是我自己的一些cuda的使用分享，因?yàn)槠蓿乱黄?，我們聊聊cuda實(shí)際使用gpu加速和cpu使用的對(duì)比，以及它兩使用場(chǎng)景分析。如果大家有更好的想法和需求，也歡迎大家加我好友交流分享哈。

作者：良知猶存，白天努力工作，晚上原創(chuàng)公號(hào)號(hào)主。公眾號(hào)內(nèi)容除了技術(shù)還有些人生感悟，一個(gè)認(rèn)真輸出內(nèi)容的職場(chǎng)老司機(jī)，也是一個(gè)技術(shù)之外豐富生活的人，攝影、音樂(lè) and 籃球。關(guān)注我，與我一起同行。