CUDA顯示卡運算程式設計菜鳥入門指南1——Hello world
第一次知道有顯示卡(GPU)程式設計這個東西,是去年比特幣最熱門的時候,看了幾篇關於比特幣的文章,說比特幣挖礦要靠顯示卡,CPU的速度與GPU根本就沒法比,於是就非常好奇,顯示卡是什麼神奇的東西?為什麼運算速度會比CPU快很多?當時也只是好奇而已,根本沒想過這東西會與自己有任何關係。
去年年底的時候,我開始研究機器學習,試著用PHP編了幾個遺傳演算法和神經網路演算法的程式,發現很有趣,功能很強大,我一直想做醫學方面的人工智慧開發,覺得機器學習這個東西很有用,但就是運算速度太慢,估算了一下按程式的執行速度,要初略的解決我的問題至少要算一兩個星期,如果要以較高的精確度解決我的問題,可能要算幾個月。仔細權衡利弊之後,我決定從頭學C語言,畢竟C是高階語言中運算速度最快的(現在後悔,當時要從C++開始學就更好了)。從初學到把我的PHP程式改寫為C,一共花了兩個星期左右,可喜的是運算速度提高了大約150倍。
後來發現神經網路用遺傳演算法來驅動遠不如用退火演算法驅動效率高,並且神經網路演算法可以有很多冗餘計算可以去掉,這樣一來,我的演算法效率前前後後又提高了一萬多倍,與最初的PHP程式相比,運算效率已經提高了大約200萬倍。 儘管如此,我還是不滿意,因為問題稍微複雜些,且要求的精度再高些,我的電腦還是要沒日沒夜轉好幾天,我希望程式運算速度能夠再有至少一個數量級的提升,在軟體上能改進的地方几乎都改進了,只能在硬體上想辦法了。
我本以為,按照摩爾定律,CPU的運算速度每一年半要提升一倍,不怕大家笑話,我現在的電腦是2007年買的,已經過去七年了,最新的CPU速度應該提升二十多倍了,於是我上網查了一下,竟然發現最新款的電腦,CPU的主頻不過比我的多10%左右,這是怎麼回事?再查才知道,原來CPU的摩爾定律到2008年左右就失效了,這麼多年來CPU主頻都沒有大的變化,主要靠不斷增加CPU的數量(雙核、四核、八核.......)進一步提高電腦效能,但這種多核結構對於我要解決的問題,幫助似乎並不明顯。
於是我就想到比特幣文章中提到的GPU運算——這個從來沒覺得會與自己有關係的東西。
查了一些文章,大致瞭解到GPU之所以比CPU快得多,並不是GPU的處理器運算速度更快——其實目前最好的顯示卡,處理器的主頻也不到CPU的一半——但是CPU只有一個處理器,而GPU的處理器少則幾十個,多則幾千個,一隻手幹活幹得再快,也沒有成百上千隻手一起幹活幹得快,就是這個道理。
GPU這個特點決定了並不是所有程式都適合用GPU來加速,只有你的問題能夠分解為若干能夠獨立執行的部分(即一部分程式的運算,並不依賴於另一部分運算的執行結果),才適合考慮用GPU來處理,這也是GPU程式設計的核心觀念:並行運算。
既然GPU是成百上千隻手一起幹活,那麼GPU的價格為什麼不是CPU的成百上千倍呢?關鍵就在於GPU的“手”與CPU的“手”相比是帶有明顯殘疾的,只有完成某些特定動作的時候效率高,完成其他動作效率就很低。具體表現在:
(1)所有的“手”(或者叫處理器、執行緒)完成相同動作時效率高,完成不同動作時效率低——這一條是顯示卡程式設計最關鍵、最核心的部分,這決定了GPU雖然有幾千隻手,但不可能像工廠工人那樣進行流水作業,你裝填,我打包,他裝箱,這對於GPU是行不通的。大家都一起打包,然後大家一起來裝箱,只有這樣的方式才適合讓GPU來處理。
(2)如果各個並行的部分(即各個執行緒),需要共享大量資料,那麼不適合用GPU進行處理。要注意,我說的是“不能共享大量資料”,這裡並不是說GPU不適合處理大量資料,相反,由於GPU具有較大的視訊記憶體空間(比較差的顯示卡也有2G視訊記憶體)和非常快的資料吞吐速度(專業術語叫頻寬),所以非常適合做大資料的處理。GPU程式設計的一個早期經典案例就是醫學影象的處理,據說有個醫生髮明瞭一種彩超快速診斷子宮肌瘤的方法,但是如果用CPU處理幾個T的彩超檢查資料的話,幾天都搞不定,後來改用GPU處理,幾十分鐘就解決了——所以處理大資料是GPU的強項。
但這是指各個執行緒獨立的讀取一部分專屬於自己的資料,那麼處理的速度可以很高。如果大量的執行緒,要相互交錯的同時訪問大量資料那麼執行效率就很低。打個比方:菜市場每一個攤位相當於一個要被訪問的資料,逛菜市場的人相當於不同的執行緒,那麼大家雜亂無章的在市場裡亂逛,執行效率就很低。如果大家排好隊進場,規定一個人只准面對一個攤位,那麼執行效率就會很高。
(3)並行處理的執行緒數不能太少。如果你的程式只能分解為幾百個並行運算的部分,就不適合用GPU來處理——至少要有幾千個並行的執行緒才能勉強體現出GPU的優勢,如果你的程式撐死了也就分解為兩三千個並行的部分,那麼還是建議你不要考慮用GPU加速了——舉個例子,我用GPU執行我的神經網路(我的顯示卡非常差,只有48個處理器),剛把程式執行通的時候,我安排了1024個執行緒進行測試,與我原來的CPU程式相比,速度只提高了20%左右,讓我大失所望。後來發發狠,把並行的執行緒數安排到8192個——速度竟然提高到了原來的八倍左右——這讓我欣喜若狂,這意味著如果我換更好的顯示卡,比如有兩千多個處理器的顯示卡,那麼我的程式速度還有四五十倍的提升空間,這樣一來我的演算法速度有望比在CPU的速度提高2-3個數量級,想想都讓人興奮。
需要說明一下,這個教程叫“菜鳥入門”,這裡的“菜鳥”是指顯示卡程式設計方面的菜鳥,而不是指程式設計菜鳥。要學顯示卡程式設計,一定要對C/C++比較熟練,對指標的運用比較自如,才可能進行顯示卡程式設計——也可以說距離顯示卡程式設計就很近了,據說很久很久以前......顯示卡程式設計是一件非常麻煩的大工程,後來一種叫CUDA的東西橫空出世——顯示卡廠商NVIDIA推出的顯示卡運算平臺——顯示卡程式設計就變得非常簡單了,只要在熟練C的前提下,學幾個簡單的語句就可以搞定。
由於我是在windows上進行顯示卡程式設計的,我簡單說一下我我在windows上安裝CUDA遇到的一些問題(詳細安裝步驟請自行搜尋相關文獻,我這裡只提一下文獻中很少提及的需要注意的問題,以幫助各位少走彎路,至於使用Linux的同學,只能是自己查閱文獻慢慢摸索了)。
(一)需要先安裝Visual Studio(其他C/C++的SDK似乎都不行),有的文獻說可以用VC EXPRESS,但我安了VC EXPRESS,死活用不了,後來裝了全套Visual Studio 才搞定。
(二)安裝好Visual Studio之後,到NVIDIA官網下載CUDA安裝包,安裝時務必選擇“自定義安裝”,務必勾選全部安裝專案。
(三)CUDA安裝好後,可以搜尋一個叫CUDA-Z的小軟體,這個軟體可以顯示你的顯示卡的狀態。下載執行後,如果告訴你顯示卡沒有執行,很可能是你的顯示卡驅動太舊了——哪怕你是剛下載的最新的安裝包,也可能出這個問題——你需要升級顯示卡驅動。另外我有個朋友遇到這樣的問題:顯示卡沒有執行,後來發現他是遠端登陸,他到本地使用的時候,顯示卡就正常了——我分析是遠端登陸後,顯示卡就休眠了。
上面步驟完成後,就可以開啟VS,新建一個CUDA工程,然後就可以開始第一個CUDA程式了:
- #include<cuda_runtime.h>
- void main()
- {printf("Hello world");
- }
慢著,這個程式和C有什麼區別?用到顯示卡了嗎?
答:沒有區別,沒用顯示卡。如果你非要用顯示卡乾點什麼事情的話,可以改成這個樣子:
- #include<cuda_runtime.h>
- __global__ void test()
- {int i=0;}
- void main()
- {test<<<1,1>>>();
- printf("Hello world");
- }
這樣,我們就用顯示卡的一個執行緒,執行了一個int i=0語句。