cuda(compute unified device architecture),是顯示卡廠商nvidia推出的運算平台。是一種通用平行計算架構,該架構使gpu能夠解決複雜的計算問題。說白了就是我們可以使用gpu來並行完成像神經網路、影象處理演算法這些在cpu上跑起來比較吃力的程式。通過gpu和高並行,我們可以大大提高這些演算法的執行速度。
有的同學可能知道,在cpu和gpu上跑同乙個神經網路,由於其大量的浮點數權重計算以及可高並行化,其速度的差距往往在10倍左右,原本需要睡一覺才能看到的訓練結果也許看兩集動漫就ok了。
gpu並行在影象處理方面更是應用廣泛,大家知道影象處理實際上是對影象的二維矩陣進行處理,影象的尺寸都是幾百乘幾百的,很容易就是上萬個畫素的操作,隨便搞個什麼平滑演算法,匹配演算法等等的影象演算法在cpu上跑個幾十秒都是很正常的,對於影象處理,神經網路這種大矩陣計算,往往是可以並行化的,通過gpu並行化處理往往能夠成倍的加速。
綜上所述,去學習一下怎麼在gpu上開個幾千個執行緒過把優化癮還是一件很愜意的事情,更何況cuda為我們提供了這麼優秀的計算平台,可以直接使用c/c++寫出在顯示晶元上執行的程式,還是一件很讚的事情。
不過cuda程式設計需要注意的點是很多的,有很多因素如果忽略了會大大降低速度,寫的不好的cuda程式可能會比cpu程式還慢。所以優化和並行是一門很大的學問,需要我們去不斷學習與了解。
cuda發展到現在說實話已經比較成熟了,當然在使用的時候偶爾會出現各種各樣的問題(充滿血與淚),但就談安裝來說已經很簡單了,這裡以vs2013和cuda 7.0為例(現在已經到cuda7.5了,我需要使用zed攝像頭,而它只支援7.0,所以電腦上裝的7.0)。
現在的cuda安裝還是很簡單的,等安裝結束之後就會發現cuda for visual studio已經安裝成功了,我們也不需要去新增什麼環境變數,這些工作安裝程式都幫我們做好了~之後我們開啟vs,也不需要繁瑣的各種引庫的過程了,我們只需要新建乙個cuda工程就可以了~
建立好工程之後,會發現已經自帶了乙個矩陣相乘的示例**kernel.cu,二話不說直接ctrl+f5編譯執行,如果沒報什麼編譯錯誤執行成功那就恭喜同學你跑了你的第乙個我cuda程式~kernel.cu
注意:這裡我再多說幾句,我關於各種錯誤的解決經驗。cuda還是會經常出現各式各樣的問題的,我自己就遇到過好幾個。
(1)首先最簡單的乙個,你的工程路徑不能有中文。。。好多個版本了都沒解決這個問題。
(3)有時候還會出現下面這個錯誤,這個也很奇葩,我隔了一周沒寫cuda程式,然後再寫的時候原來沒問題的程式都編譯不過了,周天就給我來了這麼個開門黑,重灌了各種版本的cuda仍然不行,弄了兩天才莫名其妙的弄好,這個貌似是因為.net的問題,我在控制面板-解除安裝程式-啟用或關閉windows功能 裡把.net4.5關了,開啟了.net3.5 , 重啟,然後,還是不行,我已經準備要重灌電腦了,去吃了個晚飯回來,莫名其妙行了。
(4)我還遇到過核函式進不去的情況,也是莫名其妙出現的,就是下面會講到的__global__
函式,最後被迫重灌了遍cuda,然後還是不行,重啟,結果行了。
總之大家看到我遇到的奇葩問題就知道了,這玩意有時候還是很脆弱的,什麼防毒軟體,系統更新啥的都可能隨時乾掉你的cuda,所以防患於未然還是把這些玩意都關了吧。
我知道cuda安裝的還是比較慢的,安裝的時候還是來看一下關於gpu和cuda架構的一些基礎知識吧~
上圖是cpu與gpu的對比圖,對於浮點數操作能力,cpu與gpu的能力相差在gpu更適用於計算強度高,多並行的計算中。因此,gpu擁有更多電晶體,而不是像cpu一樣的資料cache和流程控制器。這樣的設計是因為多平行計算的時候每個資料單元執行相同程式,不需要那麼繁瑣的流程控制,而更需要高計算能力,這也不需要大cache。但也因此,每個gpu的計算單元的結構是十分簡單的,因此對程式的可並行性的要求也是十分苛刻的。
這裡我們再介紹一下使用gpu計算的優缺點(摘自《深入淺出談cuda》,所以舉的例子稍微老了一點,但不影響意思哈):
使用顯示晶元來進行運算工作,和使用 cpu 相比,主要有幾個好處:
顯示晶元通常具有更大的記憶體頻寬。例如,nvidia 的 geforce 8800gtx 具有超過50gb/s 的記憶體頻寬,而目前高階 cpu 的記憶體頻寬則在 10gb/s 左右。
顯示晶元具有更大量的執行單元。例如 geforce 8800gtx 具有 128 個 「stream processors」,頻率為 1.35ghz。cpu 頻率通常較高,但是執行單元的數目則要少得多。
和高階 cpu 相比,顯示卡的**較為低廉。例如一張 geforce 8800gt 包括512mb 記憶體的**,和一顆 2.4ghz 四核心 cpu 的**相若。
當然,使用顯示晶元也有它的一些缺點:
顯示晶元的運算單元數量很多,因此對於不能高度並行化的工作,所能帶來的幫助就不大。
顯示晶元目前通常只支援 32 bits 浮點數,且多半不能完全支援 ieee 754 規格, 有些運算的精確度可能較低。目前許多顯示晶元並沒有分開的整數運算單元,因此整數運算的效率較差。
顯示晶元通常不具有分支**等複雜的流程控制單元,因此對於具有高度分支的程式,效率會比較差。
目前 gpgpu 的程式模型仍不成熟,也還沒有公認的標準。例如 nvidia 和amd/ati 就有各自不同的程式模型。
host 和 kernel:
在 cuda 的架構下,乙個程式分為兩個部份:host 端和 device 端。host 端是指在 cpu 上執行的部份,而 device 端則是在顯示晶元上執行的部份。device 端的程式又稱為 「kernel」。通常 host 端程式會將資料準備好後,複製到顯示卡的記憶體中,再由顯示晶元執行 device 端程式,完成後再由 host 端程式將結果從顯示卡的記憶體中取回。
由於 cpu 訪問顯示卡記憶體時只能透過 pci express 介面,因此速度較慢(pci express x16 的理論頻寬是雙向各 4gb/s),因此不能太常進行這類動作,以免降低效率。
thread-block-grid 結構:
在 cuda 架構下,顯示晶元執行時的最小單位是thread。數個 thread 可以組成乙個block。乙個 block 中的 thread 能訪問同一塊共享的記憶體,而且可以快速進行同步的動作。
每乙個 block 所能包含的 thread 數目是有限的。不過,執行相同程式的 block,可以組成grid。不同 block 中的 thread 無法訪問同乙個共享的記憶體,因此無法直接互通或進行同步。因此,不同 block 中的 thread 能合作的程度是比較低的。不過,利用這個模式,可以讓程式不用擔心顯示晶元實際上能同時執行的 thread 數目限制。例如,乙個具有很少量執行單元的顯示晶元,可能會把各個 block 中的 thread 順序執行,而非同時執行。不同的 grid 則可以執行不同的程式(即 kernel)。
每個 thread 都有自己的乙份 register 和 local memory 的空間。同乙個 block 中的每個thread 則有共享的乙份 share memory。此外,所有的 thread(包括不同 block 的 thread)都共享乙份 global memory、constant memory、和 texture memory。不同的 grid 則有各自的 global memory、constant memory 和 texture memory。
執行模式:
由於顯示晶元大量平行計算的特性,它處理一些問題的方式,和一般 cpu 是不同的。主要的特點包括:
記憶體訪問 latency 的問題:cpu 通常使用 cache 來減少訪問主記憶體的次數,以避免記憶體 latency 影響到執行效率。顯示晶元則多半沒有 cache(或很小),而利用並行化執行的方式來隱藏記憶體的 latency(即,當第乙個 thread 需要等待記憶體讀取結果時,則開始執行第二個 thread,依此類推)。
分支指令的問題:cpu 通常利用分支**等方式來減少分支指令造成的 pipeline bubble。顯示晶元則多半使用類似處理記憶體 latency 的方式。不過,通常顯示晶元處理分支的效率會比較差。
因此,最適合利用 cuda 處理的問題,是可以大量並行化的問題,才能有效隱藏記憶體的latency,並有效利用顯示晶元上的大量執行單元。使用 cuda 時,同時有上千個 thread 在執行是很正常的。因此,如果不能大量並行化的問題,使用 cuda 就沒辦法達到最好的效率了。
再寫下去篇幅就太長了,本篇部落格主要還是介紹了cuda的安裝以及一些基本的cuda的架構,大家趁著cuda安裝的空可以仔細看一下cuda的結構,這對後面的程式設計還是很重要的,下面我會從乙個很小的程式寫起,不斷地把上面介紹到的東西都加進去,希望能幫助到大家的學習。
CUDA程式設計(一)第乙個CUDA程式
cuda compute unified device architecture 是顯示卡廠商nvidia推出的運算平台。是一種通用平行計算架構,該架構使gpu能夠解決複雜的計算問題。說白了就是我們可以使用gpu來並行完成像神經網路 影象處理演算法這些在cpu上跑起來比較吃力的程式。通過gpu和高並...
CUDA 第乙個CUDA程式 addVector
本文主要通過對兩個浮點陣列中的資料進行相加,並將其結果放入第三個陣列中。其演算法分別在cpu gpu上分別執行,並比較了所需時間,強烈感受到gpu的平行計算能力。這裡,每個陣列的元素大小為30000000個。include include include include for the cuda r...
cuda筆記 第乙個cuda程式
釋放gpu中的記憶體cudafree cuda函式的定義 global 定義在gpu上,可以在cpu上呼叫的函式 device 定義在gpu上,由gpu呼叫函式 host 在cpu上定義的函式,一般與 device 一起用 在gpu上開闢空間 cudamalloc devptr,byte size ...