CPU+GPU異構計算程式設計簡介
分享一下我老師大神的人工智慧教程!零基礎,通俗易懂!http://blog.csdn.net/jiangjunshow
也歡迎大家轉載本篇文章。分享知識,造福人民,實現我們中華民族偉大復興!
異構計算(CPU + GPU)程式設計簡介
1. 概念
所謂異構計算,是指CPU+ GPU或者CPU+ 其它裝置(如FPGA等)協同計算。一般我們的程式,是在CPU上計算。但是,當大量的資料需要計算時,CPU顯得力不從心。那麼,是否可以找尋其它的方法來解決計算速度呢?那就是異構計算。例如可利用CPU(Central Processing Unit)、GPU(Graphic Processing Unit)、甚至APU(Accelerated Processing Units, CPU與GPU的融合)等計算裝置的計算能力從而來提高系統的速度。異構系統越來越普遍,對於支援這種環境的計算而言,也正受到越來越多的關注。
2. 異構計算的實現
目前異構計算使用最多的是利用GPU來加速。主流GPU都採用了統一架構單元,憑藉強大的可程式設計流處理器陣容,GPU在單精度浮點運算方面將CPU遠遠甩在身後。英特爾Core i7 965處理器,在預設情況下,它的浮點計算能力只有NVIDIA GeForce GTX 280 的1/13,與AMD Radeon HD 4870相比差距就更大。
3.基於GPU程式設計
不同廠商通常僅僅提供對於自己裝置程式設計的實現。對於異構系統一般很難用同種風格的程式語言來實現機構程式設計,而且將不同的裝置作為統一的計算單元來處理的難度也是非常大的。基於GPU程式設計的,目前主要兩大廠商提供:一個是NVidia,其提供的GPU程式設計為CUDA,目前使用的CUDA SDK
4. OpenCL簡介
那麼有沒有可能讓他們統一起來,簡化程式設計呢?有,那就是由蘋果公司發起並最後被業界認可的OpenCL,目前版本1.2。
開放式計算語言(Open Computing Language:OpenCL),旨在滿足這一重要需求。通過定義一套機制,來實現硬體獨立的軟體開發環境。利用OpenCL可以充分利用裝置的並行特性,支援不同級別的並行,並且能有效對映到由CPU,GPU, FPGA(Field-Programmable Gate Array)和將來出現的裝置所組成的同構或異構,單裝置或多裝置的系統。OpenCL定義了執行時, 允許用來管理資源,將不同型別的硬體結合在同種執行環境中,並且很有希望在不久的將來,以更加自然的方式支援動態地平衡計算、功耗和其他資源。
5. DirectCompute簡介
作為軟體行業的老大—微軟在這方面又做了什麼呢?微軟也沒閒著,微軟推出DirectCompute,與OpenCL抗衡。DirectCompute整合在DX中,目前版本DX11,其中就包括DirectCompute。由於微軟的地位,所以大多數廠商也都支援DirectCompute。
6. GPU計算模型
核心是執行模型的核心,能在裝置上執行。當一個核心執行之前,需要指定一個 N-維的範圍(NDRange)。一個NDRange是一個一維、二維或三維的索引空間。 還需要指定全域性工作節點的數目,工作組中節點的數目。如圖NDRange所示,全域性工作節點的範圍為{12, 12},工作組的節點範圍為{4, 4},總共有9個工作組。
如果定義向量為1024維,特別地,我們可以定義全域性工作節點為1024,工作組中節點為128,則總共有8個組。定義工作組主要是為有些僅需在組內交換資料的程式提供方便。當然工作節點數目的多少要受到裝置的限制。如果一個裝置有1024個處理節點,則1024維的向量,每個節點計算一次就能完成。而如果一個裝置僅有128個處理節點,那麼每個節點需要計算8次。合理設定節點數目,工作組數目能提高程式的並行度。
7. 程式例項
不論是OpenCL還是DirectCompute,其程式設計風格都基本差不多,程式是分成兩部分的:一部分是在裝置上執行的(對於我們,是GPU),另一部分是在主機上執行的(對於我們,是CPU)。在裝置上執行的程式或許是你比較關注的。它是OpenCL和DirectCompute產生神奇力量的地方。為了能在裝置上執行程式碼,OpenCL程式設計師需要寫一個特殊的函式(kernel函式)放在專用檔案中(.cl),這個函式需要使用OpenCL語言編寫。OpenCL語言採用了C語言的一部分加上一些約束、關鍵字和資料型別。在主機上執行的程式提供了API,所以可以管理你在裝置上執行的程式。主機程式可以用C或者C++編寫,它控制OpenCL的環境(上下文,指令佇列…)。DirectCompute程式設計師需要寫Shader檔案(.hlsl),在這個檔案中寫函式。Shader檔案的格式可以查MSDN。
在寫程式時,先初始化裝置,然後編譯需要在GPU上執行的程式(執行在GPU上的程式是在應用程式執行時編譯的)。然後對映需要在GPU上執行的函式名字,OpenCL 呼叫clEnqueueNDRangeKernel執行kernel函式,DirectCompute呼叫ID3D11DeviceContext:: Dispatch執行Shader函式。函式是併發執行的。
執行在GPU上的函式一般都很簡單。以求和為例:
用CPU運算
void vector_add_cpu (const float* fIn1, const float* fIn2, float* fOut, const int iNum){ for (int i = 0; i < iNum; i++) { fOut [i] = fIn1[i] + fIn2[i]; }}
以下是OPenCL的kernel函式
//在GPU上,邏輯就會有一些不同。我們使每個執行緒計算一個元素的方法來代替cpu程式中的迴圈計算。每個執行緒的index與要計算的向量的index相同。__kernel void vector_add_gpu (__global const float* fIn1, __global const float* fIn2, __global float* fOut, const int iNum){ /* get_global_id(0) 返回正在執行的這個執行緒的ID。 許多執行緒會在同一時間開始執行同一個kernel, 每個執行緒都會收到一個不同的ID,所以必然會執行一個不同的計算。*/ const int idx = get_global_id(0); /* 每個work-item都會檢查自己的id是否在向量陣列的區間內。 如果在,work-item就會執行相應的計算。*/ if (idx < iNum) { fOut [idx] = fIn1[idx] + fIn2[idx]; }}
有一些需要注意的地方:
1. Kernel關鍵字定義了一個函式是kernel函式。Kernel函式必須返回void。
2. Global關鍵字位於引數前面。它定義了引數記憶體的存放位置。
另外,所有kernel都必須寫在“.cl”檔案中,“.cl”檔案必須只包含OpenCL程式碼。
Shader函式
#define NUM_THREAD 16StructuredBuffer<float> fInput1 : register( t0 );StructuredBuffer<float> fInput2 : register( t0 );StructuredBuffer<float> fOutput : register( u0 );[numthreads(NUM_THREAD, 1, 1)]void vector_add_gpu( uint3 Gid : SV_GroupID, uint3 DTid : SV_DispatchThreadID, uint3 GTid : SV_GroupThreadID, uint GI : SV_GroupIndex ){ fOutput[DTid.x] = fInput1[DTid.x] + fInput2[DTid.x] ;}
影象旋轉是指把定義的影象繞某一點以逆時針或順時針方向旋轉一定的角度,通常是指繞影象的中心以逆時針方向旋轉。假設影象的左上角為(l, t), 右下角為(r, b),則影象上任意點(x, y) 繞其中心(xcenter, ycenter)逆時針旋轉θ角度後, 新的座標位置(x',y')的計算公式為:
x′ = (x - xcenter) cosθ - (y - ycenter) sinθ + xcenter,
y′ = (x - xcenter) sinθ + (y - ycenter) cosθ + ycenter.
C程式碼:
void image_rotate( unsigned int* iInbuf, unsigned int* iOutbuf, int iWidth, int iHeight, float fSinTheta, float fCosTheta){ int i, j; int iXc = iWidth /2; int iYc = iHeight /2; for(i = 0; i < iHeight; i++) { for(j=0; j< iWidth; j++) { int iXpos = (j- iXc)*fCosTheta - (i - iYc) * fSinTheta + iXc; int iYpos = (j- iXc)*fSinTheta + (i - iYc) * fCosTheta + iYc; if(iXpos >=0&& iYpos >=0&& iXpos < iWidth && iYpos < iHeight) iOutbuf[iYpos * iWidth + iXpos] = iInbuf[i* iWidth +j]; } }}
CL程式碼:
__kernel void image_rotate( __global int * iInbuf, __global int * iOutbuf, //Data in global memory int iWidth , int iHeight, //Image Dimensions float fSinTheta, float fCosTheta ) //Rotation Parameters{ const int ix = get_global_id(0); const int iy = get_global_id(1); int iXc = iWidth /2; int iYc = iHeight /2; int iXpos = ( ix- iXc)*fCosTheta - (iy- iYc)*fSinTheta+ iXc; int iYpos = (ix- iXc)*fSinTheta + ( iy- iYc)*fCosTheta+ iYc; if ((iXpos >=0) && (iXpos < iWidth ) && (iYpos >=0) && (iYpos < iHeight)) iOutbuf[iYpos * iWidth + iXpos]= iInbuf[iy* iWidth +ix];}
不論是OpenCL還是 DirectCompute,其程式設計還是有些複雜,特別是對於裝置的初始化,以及資料交換,非常麻煩。對於初學者難度相當大。那麼有沒有更簡單的程式設計方法呢?
8. C++AMP
還要提到微軟,因為我們基本上都使用微軟的東東。微軟也不錯,推出了C++AMP,這是個開放標準,嵌入到VS2012中(VS2012目前還是預覽版),在Win7和Win8系統中才能使用,其使用就簡單多了:
void CppAmpMethod(){ int aCPP[] = {1, 2, 3, 4, 5}; int bCPP[] = {6, 7, 8, 9, 10}; int sumCPP[size]; // Create C++ AMP objects. array_view<const int, 1> a(size, aCPP); array_view<const int, 1> b(size, bCPP); array_view<int, 1> sum(size, sumCPP); sum.discard_data(); parallel_for_each( // Define the compute domain, which is the set of threads that are created. sum.extent, // Define the code to run on each thread on the accelerator. [=](index<1> idx) restrict(amp) { sum[idx] = a[idx] + b[idx]; } );}
就這麼簡單,只需要包含一個頭檔案,使用一個名稱空間,包含庫檔案,一切就OK,在呼叫時,只是一個函式parallel_for_each。(有點類似OpenMP)。
9. 應用
MATLAB 2010b中Parallel Computing Toolbox與MATLAB Distributed Computing Server的最新版本可利用NVIDIA的CUDA平行計算架構在NVIDIA計算能力1.3以上的GPU上處理資料,執行GPU加速的MATLAB運算,將使用者自己的CUDA Kernel函式整合到MATLAB應用程式當中。另外,通過在臺式機上使用Parallel Computing Toolbox以及在計算叢集上使用MATLAB Distributed Computing Server來執行多個MATLAB worker程式,從而可在多顆NVIDIA GPU上進行計算。AccelerEyes公司開發的Jacket外掛也能夠使MATLAB利用GPU進行加速計算。Jacket不僅提供了GPU API(應用程式介面),而且還集成了GPU MEX功能。在一定程度說,Jacket是一個完全對使用者透明的系統,能夠自動的進行記憶體分配和自動優化。Jacket使用了一個叫“on-the- fly”的編譯系統,使MATLAB互動式格式的程式能夠在GPU上執行。目前,Jacket只是基於NVIDIA的CUDA技術,但能夠執行在各主流作業系統上。
從Photoshop CS4開始,Adobe將GPU通用計算技術引入到自家的產品中來。GPU可提供對影象旋轉、縮放和放大平移這些常規瀏覽功能的加速,還能夠實現2D/3D合成,高質量抗鋸齒,HDR高動態範圍貼圖,色彩轉換等。而在Photoshop CS5中,更多的演算法和濾鏡也開始支援GPU加速。另外,Adobe的其他產品如Adobe After Effects CS4、Adobe Premiere Pro CS4也開始使用GPU進行加速。這些軟體藉助的也是NVIDIA的CUDA技術。
Windows 7 的核心組成部分包括了支援GPU通用計算的Directcompute API,為視訊處理、動態模擬等應用進行加速。Windows 7藉助Directcompute增加了對由GPU支援的高清播放的in-the-box支援,可以流暢觀看,同時CPU佔用率很低。Internet Explorer 9加入了對Directcompute技術的支援,可以呼叫GPU對網頁中的大計算量元素做加速計算;Excel2010、Powerpoint2010也開始提供對Directcompute技術的支援。
比利時安特衛普大學,通用電氣醫療集團,西門子醫療,東芝中風研究中心和紐約州立大學水牛城分校的都針對GPU加速CT重建進行了各自的研究,不僅如此,西門子醫療用GPU實現了加速MRI中的GRAPPA自動校準,完成MR重建,快速MRI網格化,隨機擴散張量磁共振影象(DT-MRI)連通繪圖等演算法。其他的一些研究者則把醫學成像中非常重要的二維與三維影象中器官分割(如Level Set演算法),不同來源影象的配準,重建體積影象的渲染等也移植到GPU上進行計算。
10 .不足
異構平行計算變得越來越普遍,然而對於現今存在的OpenCL和DirectCompute版本來說,的確還存在很多不足,例如編寫核心,需要對問題的並行情況做較為深入的分析,對於記憶體的管理還是需要程式設計師來顯式地申明、顯式地在主存和裝置的儲存器之間進行移動,還不能完全交給系統自動完成。從這些方面,OpenCL和 DirectCompute的確還需加強,要使得人們能高效而又靈活地開發應用,還有很多工作要完成。
http://blog.csdn.net/iddialog/article/details/8078747