小男孩‘自慰网亚洲一区二区,亚洲一级在线播放毛片,亚洲中文字幕av每天更新,黄aⅴ永久免费无码,91成人午夜在线精品,色网站免费在线观看,亚洲欧洲wwwww在线观看

搜索
分享

CPU GPU異構(gòu)計(jì)算編程簡介

 敬而遠(yuǎn) 2013-12-21

異構(gòu)計(jì)算（CPU + GPU）編程簡介

1. 概念

所 謂異構(gòu)計(jì)算，是指CPU+ GPU或者CPU+ 其它設(shè)備（如FPGA等）協(xié)同計(jì)算。一般我們的程序，是在CPU上計(jì)算。但是，當(dāng)大量的數(shù)據(jù)需要計(jì)算時(shí)，CPU顯得力不從心。那么， 是否可以找尋其它的方法來解決計(jì)算速度呢？那就是異構(gòu)計(jì)算。例如可利用CPU（Central Processing Unit）、GPU（Graphic Processing Unit）、甚至APU(Accelerated Processing Units， CPU與GPU的融合)等計(jì)算設(shè)備的計(jì)算能力從而來提高系統(tǒng)的速度。異構(gòu)系統(tǒng)越來越普遍，對于支持這種環(huán)境的計(jì)算而言，也正受到越來越多的 關(guān)注。

2. 異構(gòu)計(jì)算的實(shí)現(xiàn)

目 前異構(gòu)計(jì)算使用最多的是利用GPU來加速。主流GPU都采用了統(tǒng)一架構(gòu)單元，憑借強(qiáng)大的可編程流處理器陣容，GPU在單精度浮點(diǎn)運(yùn)算方面將CPU遠(yuǎn)遠(yuǎn)甩在 身后。英特爾Core i7 965處理器，在默認(rèn)情況下，它的浮點(diǎn)計(jì)算能力只有NVIDIA GeForce GTX 280 的1/13，與AMD Radeon HD 4870相比差距就更大。

3.基于GPU編程

不 同廠商通常僅僅提供對于自己設(shè)備編程的實(shí)現(xiàn)。對于異構(gòu)系統(tǒng)一般很難用同種風(fēng)格的編程語言來實(shí)現(xiàn)機(jī)構(gòu)編程，而且將不同的設(shè)備作為統(tǒng)一的計(jì)算單元來處理的難度 也是非常大的。基于GPU編程的，目前主要兩大廠商提供：一個(gè)是NVidia，其提供的GPU編程為CUDA，目前使用的CUDA SDK 4.2.另一個(gè)是AMD，其提供的GPU編程為AMD APP （其前身是ATI Stream），目前最新版本 AMD APP 2.7。這兩個(gè)東東是不兼容的，各自為政。作為軟件開發(fā)者而言，用CUDA開發(fā)的軟件只能在NVidia相應(yīng)的顯卡上運(yùn)行，用AMD APP開發(fā)的軟件，只能在ATI相應(yīng)的顯卡上運(yùn)行。

4.  OpenCL簡介

那么有沒有可能讓他們統(tǒng)一起來，簡化編程呢？有，那就是由蘋果公司發(fā)起并最后被業(yè)界認(rèn)可的OpenCL，目前版本1.2。

開 放式計(jì)算語言（Open Computing Language:OpenCL），旨在滿足這一重要需求。通過定義一套機(jī)制，來實(shí)現(xiàn)硬件獨(dú)立的軟件開發(fā)環(huán)境。利用OpenCL可以充分利用設(shè)備的并行特 性，支持不同級別的并行，并且能有效映射到由CPU，GPU， FPGA（Field－Programmable Gate Array）和將來出現(xiàn)的設(shè)備所組成的同構(gòu)或異構(gòu)，單設(shè)備或多設(shè)備的系統(tǒng)。OpenCL定義了運(yùn)行時(shí)， 允許用來管理資源，將不同類型的硬件結(jié)合在同種執(zhí)行環(huán)境中，并且很有希望在不久的將來，以更加自然的方式支持動(dòng)態(tài)地平衡計(jì)算、功耗和其他資源。

5. DirectCompute簡介

作 為軟件行業(yè)的老大—微軟在這方面又做了什么呢？微軟也沒閑著，微軟推出DirectCompute，與OpenCL抗衡。DirectCompute集成 在DX中，目前版本DX11，其中就包括DirectCompute。由于微軟的地位，所以大多數(shù)廠商也都支持DirectCompute。

6. GPU計(jì)算模型

內(nèi)核是執(zhí)行模型的核心，能在設(shè)備上執(zhí)行。當(dāng)一個(gè)內(nèi)核執(zhí)行之前，需要指定一個(gè) N-維的范圍（NDRange）。一個(gè)NDRange是一個(gè)一維、二維或三維的索引空間。 還需要指定全局工作節(jié)點(diǎn)的數(shù)目，工作組中節(jié)點(diǎn)的數(shù)目。如圖NDRange所示，全局工作節(jié)點(diǎn)的范圍為{12, 12}，工作組的節(jié)點(diǎn)范圍為{4, 4}，總共有9個(gè)工作組。

  

如 果定義向量為1024維，特別地，我們可以定義全局工作節(jié)點(diǎn)為1024，工作組中節(jié)點(diǎn)為128，則總共有8個(gè)組。定義工作組主要是為有些僅需在組內(nèi)交換數(shù) 據(jù)的程序提供方便。當(dāng)然工作節(jié)點(diǎn)數(shù)目的多少要受到設(shè)備的限制。如果一個(gè)設(shè)備有1024個(gè)處理節(jié)點(diǎn)，則1024維的向量，每個(gè)節(jié)點(diǎn)計(jì)算一次就能完成。而如果 一個(gè)設(shè)備僅有128個(gè)處理節(jié)點(diǎn)，那么每個(gè)節(jié)點(diǎn)需要計(jì)算8次。合理設(shè)置節(jié)點(diǎn)數(shù)目，工作組數(shù)目能提高程序的并行度。

 

7. 程序?qū)嵗?/strong>

不論是OpenCL還是DirectCompute，其編程風(fēng)格都基本差不多，程序是分成兩部分的：一部分是在設(shè)備上執(zhí)行的（對于我們，是GPU），另一部分是在主機(jī)上運(yùn)行的（對于我們，是CPU）。在設(shè)備上執(zhí)行的程序或許是你比較關(guān)注的。它是OpenCL和DirectCompute產(chǎn)生神奇力量的地方。為了能在設(shè)備上執(zhí)行代碼，OpenCL程 序員需要寫一個(gè)特殊的函數(shù)（kernel函數(shù)）放在專用文件中（.cl），這個(gè)函數(shù)需要使用OpenCL語言編寫。OpenCL語言采用了C語言的一部分 加上一些約束、關(guān)鍵字和數(shù)據(jù)類型。在主機(jī)上運(yùn)行的程序提供了API，所以可以管理你在設(shè)備上運(yùn)行的程序。主機(jī)程序可以用C或者C++編寫，它控制 OpenCL的環(huán)境（上下文，指令隊(duì)列…）。DirectCompute程序員需要寫Shader文件（.hlsl），在這個(gè)文件中寫函數(shù)。Shader文件的格式可以查MSDN。

在寫程序時(shí)，先初始化設(shè)備，然后編譯需要在GPU上運(yùn)行的程序（運(yùn)行在GPU上的程序是在應(yīng)用程序運(yùn)行時(shí)編譯的）。然后映射需要在GPU上運(yùn)行的函數(shù)名字，OpenCL 調(diào)用clEnqueueNDRangeKernel執(zhí)行kernel函數(shù)，DirectCompute調(diào)用ID3D11DeviceContext:: Dispatch執(zhí)行Shader函數(shù)。函數(shù)是并發(fā)執(zhí)行的。

運(yùn)行在GPU上的函數(shù)一般都很簡單。以求和為例：

用CPU運(yùn)算

  1. void vector_add_cpu (const float* fIn1,  
  2.                      const float* fIn2,  
  3.                      float*  fOut,  
  4.                      const int iNum)  
  5. {  
  6.     for (int i = 0; i < iNum; i++)  
  7.     {  
  8.         fOut [i] = fIn1[i] + fIn2[i];  
  9.     }  
  10. }  

以下是OPenCL的kernel函數(shù)

  1. //在GPU上，邏輯就會有一些不同。我們使每個(gè)線程計(jì)算一個(gè)元素的方法來代替cpu程序中的循環(huán)計(jì)算。每個(gè)線程的index與要計(jì)算的向量的index相同。  
  2. __kernel void vector_add_gpu (__global const float* fIn1,  
  3.                      __global const float* fIn2,  
  4.                      __global float* fOut,  
  5.            const int iNum)  
  6. {  
  7.    /* get_global_id(0) 返回正在執(zhí)行的這個(gè)線程的ID。 
  8.    許多線程會在同一時(shí)間開始執(zhí)行同一個(gè)kernel， 
  9.    每個(gè)線程都會收到一個(gè)不同的ID，所以必然會執(zhí)行一個(gè)不同的計(jì)算。*/  
  10.    const int idx = get_global_id(0);  
  11.    
  12.    /* 每個(gè)work-item都會檢查自己的id是否在向量數(shù)組的區(qū)間內(nèi)。 
  13.    如果在，work-item就會執(zhí)行相應(yīng)的計(jì)算。*/  
  14.    if (idx < iNum)  
  15.     {  
  16.             fOut [idx] = fIn1[idx] + fIn2[idx];  
  17.     }  
  18. }  

有一些需要注意的地方：

1. Kernel關(guān)鍵字定義了一個(gè)函數(shù)是kernel函數(shù)。Kernel函數(shù)必須返回void。

2. Global關(guān)鍵字位于參數(shù)前面。它定義了參數(shù)內(nèi)存的存放位置。

另外，所有kernel都必須寫在“.cl”文件中，“.cl”文件必須只包含OpenCL代碼。

 

Shader函數(shù)

  1. #define NUM_THREAD 16  
  2. StructuredBuffer<float> fInput1 : register( t0 );  
  3. StructuredBuffer<float> fInput2 : register( t0 );  
  4. StructuredBuffer<float> fOutput : register( u0 );  
  5.   
  6. [numthreads(NUM_THREAD, 1, 1)]  
  7. void vector_add_gpu( uint3 Gid : SV_GroupID,  
  8.                     uint3 DTid : SV_DispatchThreadID,  
  9.                     uint3 GTid : SV_GroupThreadID,  
  10.                     uint GI : SV_GroupIndex )  
  11. {  
  12.     fOutput[DTid.x] =  fInput1[DTid.x]  + fInput2[DTid.x] ;  
  13. }  

圖像旋轉(zhuǎn)是指把定義的圖像繞某一點(diǎn)以逆時(shí)針或順時(shí)針方向旋轉(zhuǎn)一定的角度，通常是指繞圖像的中心以逆時(shí)針方向旋轉(zhuǎn)。假設(shè)圖像的左上角為(l, t), 右下角為(r, b)，則圖像上任意點(diǎn)(x, y) 繞其中心(xcenter, ycenter)逆時(shí)針旋轉(zhuǎn)θ角度后， 新的坐標(biāo)位置(x',y')的計(jì)算公式為：

x′ = (x - xcenter) cosθ - (y － ycenter) sinθ + xcenter,

y′ = (x - xcenter) sinθ + (y － ycenter) cosθ + ycenter.

 

C代碼：

  1. void image_rotate(  
  2.                   unsigned int* iInbuf,  
  3.                   unsigned int* iOutbuf,  
  4.                   int iWidth, int iHeight,  
  5.                   float fSinTheta,  
  6.                   float fCosTheta)  
  7. {  
  8.     int i, j;  
  9.     int iXc = iWidth /2;  
  10.     int iYc = iHeight /2;  
  11.     for(i = 0; i < iHeight; i++)  
  12.     {  
  13.         for(j=0; j< iWidth; j++)  
  14.         {  
  15.             int iXpos =  (j- iXc)*fCosTheta - (i - iYc) * fSinTheta + iXc;  
  16.             int iYpos =  (j- iXc)*fSinTheta + (i - iYc) * fCosTheta + iYc;  
  17.             if(iXpos >=0&& iYpos >=0&& iXpos < iWidth  && iYpos < iHeight)  
  18.                 iOutbuf[iYpos * iWidth  + iXpos] = iInbuf[i* iWidth  +j];  
  19.         }  
  20.     }  
  21. }  

CL代碼：

  1. __kernel  void image_rotate(  
  2.                             __global int * iInbuf,  
  3.                             __global int * iOutbuf,        //Data in global memory  
  4.                             int iWidth  ,    int iHeight,                  //Image Dimensions  
  5.                             float fSinTheta, float fCosTheta )   //Rotation Parameters  
  6. {  
  7.     const int ix = get_global_id(0);  
  8.     const int iy = get_global_id(1);  
  9.     int iXc = iWidth  /2;  
  10.     int iYc = iHeight /2;  
  11.     int iXpos =  ( ix- iXc)*fCosTheta - (iy- iYc)*fSinTheta+ iXc;  
  12.     int iYpos =  (ix- iXc)*fSinTheta + ( iy- iYc)*fCosTheta+ iYc;  
  13.     if ((iXpos >=0) && (iXpos < iWidth  )   && (iYpos >=0) && (iYpos < iHeight))  
  14.         iOutbuf[iYpos * iWidth  + iXpos]= iInbuf[iy* iWidth  +ix];  
  15. }  

不論是OpenCL還是 DirectCompute，其編程還是有些復(fù)雜，特別是對于設(shè)備的初始化，以及數(shù)據(jù)交換，非常麻煩。對于初學(xué)者難度相當(dāng)大。那么有沒有更簡單的編程方法呢？

8. C++AMP

還要提到微軟，因?yàn)槲覀兓旧隙际褂梦④浀臇|東。微軟也不錯(cuò)，推出了C++AMP，這是個(gè)開放標(biāo)準(zhǔn)，嵌入到VS2012中（VS2012目前還是預(yù)覽版），在Win7和Win8系統(tǒng)中才能使用，其使用就簡單多了：

  1. void CppAmpMethod()  
  2. {  
  3.     int aCPP[] = {1, 2, 3, 4, 5};  
  4.     int bCPP[] = {6, 7, 8, 9, 10};  
  5.     int sumCPP[size];  
  6.   
  7.     // Create C++ AMP objects.  
  8.     array_view<const int, 1> a(size, aCPP);  
  9.     array_view<const int, 1> b(size, bCPP);  
  10.     array_view<int, 1> sum(size, sumCPP);  
  11.     sum.discard_data();  
  12.   
  13.     parallel_for_each(  
  14.         // Define the compute domain, which is the set of threads that are created.  
  15.         sum.extent,  
  16.         // Define the code to run on each thread on the accelerator.  
  17.         [=](index<1> idx) restrict(amp)  
  18.     {  
  19.         sum[idx] = a[idx] + b[idx];  
  20.     }  
  21.     );  
  22. }  

就這么簡單，只需要包含一個(gè)頭文件，使用一個(gè)命名空間，包含庫文件，一切就OK，在調(diào)用時(shí)，只是一個(gè)函數(shù)parallel_for_each。（有點(diǎn)類似OpenMP）。


9. 應(yīng)用

MATLAB 2010b中Parallel Computing Toolbox與MATLAB Distributed Computing Server的最新版本可利用NVIDIA的CUDA并行計(jì)算架構(gòu)在NVIDIA計(jì)算能力1.3以上的GPU上處理數(shù)據(jù)，執(zhí)行GPU加速的MATLAB運(yùn) 算，將用戶自己的CUDA Kernel函數(shù)集成到MATLAB應(yīng)用程序當(dāng)中。另外，通過在臺式機(jī)上使用Parallel Computing Toolbox以及在計(jì)算集群上使用MATLAB Distributed Computing Server來運(yùn)行多個(gè)MATLAB worker程序，從而可在多顆NVIDIA GPU上進(jìn)行計(jì)算。AccelerEyes公司開發(fā)的Jacket插件也能夠使MATLAB利用GPU進(jìn)行加速計(jì)算。Jacket不僅提供了GPU API（應(yīng)用程序接口），而且還集成了GPU MEX功能。在一定程度說，Jacket是一個(gè)完全對用戶透明的系統(tǒng)，能夠自動(dòng)的進(jìn)行內(nèi)存分配和自動(dòng)優(yōu)化。Jacket使用了一個(gè)叫“on-the- fly”的編譯系統(tǒng)，使MATLAB交互式格式的程序能夠在GPU上運(yùn)行。目前，Jacket只是基于NVIDIA的CUDA技術(shù)，但能夠運(yùn)行在各主流操 作系統(tǒng)上。

從Photoshop CS4開 始，Adobe將GPU通用計(jì)算技術(shù)引入到自家的產(chǎn)品中來。GPU可提供對圖像旋轉(zhuǎn)、縮放和放大平移這些常規(guī)瀏覽功能的加速，還能夠?qū)崿F(xiàn)2D/3D合成， 高質(zhì)量抗鋸齒，HDR高動(dòng)態(tài)范圍貼圖，色彩轉(zhuǎn)換等。而在Photoshop CS5中，更多的算法和濾鏡也開始支持GPU加速。另外，Adobe的其他產(chǎn)品如Adobe After Effects CS4、Adobe Premiere Pro CS4也開始使用GPU進(jìn)行加速。這些軟件借助的也是NVIDIA的CUDA技術(shù)。

Windows 7 的核心組成部分包括了支持GPU通 用計(jì)算的Directcompute API，為視頻處理、動(dòng)態(tài)模擬等應(yīng)用進(jìn)行加速。Windows 7借助Directcompute增加了對由GPU支持的高清播放的in-the-box支持，可以流暢觀看，同時(shí)CPU占用率很低。Internet Explorer 9加入了對Directcompute技術(shù)的支持，可以調(diào)用GPU對網(wǎng)頁中的大計(jì)算量元素做加速計(jì)算；Excel2010、Powerpoint2010 也開始提供對Directcompute技術(shù)的支持。

比利時(shí)安特衛(wèi)普大學(xué)，通用電氣醫(yī)療集團(tuán)，西門子醫(yī)療，東芝中風(fēng)研究中心和紐約州立大學(xué)水牛城分校的都針對GPU加 速CT重建進(jìn)行了各自的研究，不僅如此，西門子醫(yī)療用GPU實(shí)現(xiàn)了加速M(fèi)RI中的GRAPPA自動(dòng)校準(zhǔn)，完成MR重建，快速M(fèi)RI網(wǎng)格化，隨機(jī)擴(kuò)散張量磁 共振圖像（DT-MRI）連通繪圖等算法。其他的一些研究者則把醫(yī)學(xué)成像中非常重要的二維與三維圖像中器官分割（如Level Set算法），不同來源圖像的配準(zhǔn)，重建體積圖像的渲染等也移植到GPU上進(jìn)行計(jì)算。

10 .不足

異 構(gòu)并行計(jì)算變得越來越普遍，然而對于現(xiàn)今存在的OpenCL和DirectCompute版本來說，的確還存在很多不足，例如編寫內(nèi)核，需要對問題的并行 情況做較為深入的分析，對于內(nèi)存的管理還是需要程序員來顯式地申明、顯式地在主存和設(shè)備的存儲器之間進(jìn)行移動(dòng)，還不能完全交給系統(tǒng)自動(dòng)完成。從這些方 面，OpenCL和 DirectCompute的確還需加強(qiáng)，要使得人們能高效而又靈活地開發(fā)應(yīng)用，還有很多工作要完成。

 http://blog.csdn.net/iddialog/article/details/8078747

    本站是提供個(gè)人知識管理的網(wǎng)絡(luò)存儲空間，所有內(nèi)容均由用戶發(fā)布，不代表本站觀點(diǎn)。請注意甄別內(nèi)容中的聯(lián)系方式、誘導(dǎo)購買等信息，謹(jǐn)防詐騙。如發(fā)現(xiàn)有害或侵權(quán)內(nèi)容，請點(diǎn)擊一鍵舉報(bào)。
    轉(zhuǎn)藏 分享 獻(xiàn)花（0）

    0條評論

    發(fā)表

    請遵守用戶 評論公約

    類似文章 更多