異構計算(CPU + GPU)編程簡介

2021-01-07 CSDN技術社區

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.2.另一個是AMD,其提供的GPU編程為AMD APP(其前身是ATI Stream),目前最新版本AMD APP 2.7。這兩個東西是不兼容的,各自為政。作為軟體開發者而言,用CUDA開發的軟體只能在NVidia相應的顯卡上運行,用AMD APP開發的軟體,只能在ATI相應的顯卡上運行。

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相同。

__kernelvoid 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函數

#defineNUM_THREAD 16

StructuredBuffer <float> fInput1: register(t0 );

StructuredBuffer <float> fInput2: register(t0 );

StructuredBuffer <float> fOutput: register(u0 );

[numthreads(NUM_THREAD, 1, 1)]

voidvector_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代碼:

voidimage_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系統中才能使用,其使用就簡單多了:

#include<amp.h>

usingnamespaceconcurrency;

constint size = 5;

voidCppAmpMethod()

{

   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高動態範圍貼圖,色彩轉換等。而在PhotoshopCS5中,更多的算法和濾鏡也開始支持GPU加速。另外,Adobe的其他產品如Adobe After EffectsCS4、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的確還需加強,要使得人們能高效而又靈活地開發應用,還有很多工作要完成。

相關焦點

  • CUDA編程之認識CPU與GPU
    CPU架構示意圖追求單線程的最高性能,對延遲敏感,指令級並行大量的電晶體用於緩存而非計算單元,而緩存並不 提供原生/峰值計算能力CPU有強大的ALU,時鐘頻率很高,但由於散熱、電晶體尺寸等影響,近年來GPU架構示意圖總的來說,CPU擅長處理邏輯複雜、串行的計算任務;而GPU擅長的是大規模的數據並行(data-parallel)的計算任務。
  • CPU+獨顯實戰異構計算
    其實就是宣稱APU內置的GPU能支持「異構計算」,實現對部分軟體的加速,使CPU與GPU協同運算,提升電腦的運算效率。   雖然GPU的「異構計算」很美好,也反映了未來發展方向,但理想與總是有的差距。在異構計算還僅僅停留在起步階段的今天,「異構計算」是否有實用價值呢?它僅僅是APU獨有的技術嗎?今天我們一起來探討一下。
  • gpu渲染和cpu渲染有什麼區別_GPU渲染有何優勢?
    gpu渲染簡介 1、gpu是圖形處理器,啟用Gpu渲染加速,就是調用gpu加速圖形的渲染和填充。         2、開啟gpu渲染加速後可以提升圖形加載速度,降低cpu處理器的負擔,使系統運行更加流暢,但是也更加耗電。 GPU與CPU有何不同?
  • 英特爾發布oneAPI計劃及其beta產品 簡化異構計算的開發編程
    北京時間11月18日,在2019超級計算大會上,英特爾發布了全新軟體行業計劃oneAPI及oneAPI beta產品,以在異構計算時代下,為開發者提供一個統一和簡化的應用程式開發編程模型。      按照英特爾的規劃,oneAPI計劃為開發者提供了一個通用、開放的編程體驗,使其可以自由選擇架構,無需在性能上作出妥協,並降低了開發者在不同代碼庫、程式語言、編程工具和工作流程中的複雜性。同時,oneAPI保留了現有軟體投資,支持現有語言,在開發人員創造更多豐富應用程式提供了靈活性。
  • 靈活控制GPU/CPU協同工作 詳解聯發科CorePilot異構計算技術
    本文引用地址:http://www.eepw.com.cn/article/201710/366435.htm  在2013年,聯發科全球首創CorePilot異構計算技術,CorePilot是聯發科為旗下多核心產品量身定製的一項新技術,CorePilot異構計算技術可以簡單的看做ARM Big.LITTLE,是一種大核心+小核心的架構技術,如今已經進化到了3.0版本,不僅可以支持大小核異構計算
  • gpu工作原理介紹
    在 2003-2004年左右,圖形學之外的領域專家開始注意到gpu與眾不同的計算能力,開始嘗試把gpu用於通用計算(即GPgpu)。之後NVIDIA發布了CUDA,amd和等公司也發布了OpenCL,gpu開始在通用計算領域得到廣泛應用,包括:數值分析,海量數據處理(排序,Map- Reduce等),金融分析等等。
  • 異構計算神器來了,它能帶來性能革命嗎
    如此一來,「僅用一種語言,編寫出自動適配多種不同計算架構的程序,並讓程序在運行時更充分的利用現代數碼設備的異構計算能力」,就突然間變得輕而易舉了。的oneAPI的確可能是當今技術最先進,使用最方便,最有可能實現應用程式「異構計算」優化,從而帶來執行效能飛躍的編程方案。
  • 從CPU到GPU 未來計算應用誰主沉浮?
    該處理器兼容x86編程模型,能夠為高性能計算進一步實現優化,通過與強大靈活高效的至強處理器平臺協同,能夠獲得10Petaflops(萬萬億次浮點計算)級別的超級計算能力。不僅能遠遠滿足以往的各種繁重數據查詢、處理和順序串行計算,而且通過多個多核處理器,也能勝任各種並行計算需求。在不遠的將來,由於並行計算的廣闊前景,CPU更多的會以多處理器架構或者CPU+GPU異構計算來實現。在這一點上,未來計算應用的核心地位,非CPU所能撼動。3後起之秀 GPU高性能需高能耗  剛才我們提到並行計算,其實這恰恰是GPU的優勢所在。
  • PowerVR 8系列GPU全新發布:異構計算+光線追蹤
    蘋果iPhone 6S、iPad Pro剛剛用上PowerVR 7系列圖形核心,SC15高性能計算大會上,Imagination就宣布了全新一代的
  • 按需混搭集合眾多技術點異構計算有望重塑產業格局
    異構計算:採用混合動力的全新計算引擎異構計算的概念並不新鮮,早在上世紀80年代,異構計算就已經誕生,不過早期的異構融合主要是基於CPU和GPU,而目前,異構計算已經擴展到更多計算架構,如CPU、GPU、ASIC、FPGA等,並且在更多的領域如人工智慧、邊緣計算等發揮重要的作用,能夠實現性能大幅提升同時降低能耗
  • 異構計算,如何讓開源有更開放的生態?
    為了達到這一特性,從硬體設計層面,就會優化計算單元跟存儲單元之間的物理位置以及訪問協議設置。第三,能效比更高的專用硬體。針對特定的業務場景採用專用硬體,很好地提供能效比。第四,新的指令集或像MindSpore這類的新編程範式。
  • 高性能計算中的異構計算是什麼鬼,竟這麼重要
    當摩爾定律還是行業的鐵律時,計算機編程幾乎一直都是串行的,絕大多數的程序只存在一個進程或線程。大家還過著「我寫個程序,性能達不到就睡個覺,等硬體工藝刷新硬體性能,性能就達標了」的美好生活。2003年以後因為工藝達到了瓶頸,你光「睡覺」是不行了。異構並行計算作為重要技術之一承擔起了技術變革的重任。
  • NVIDIA GPU 加速科學計算專場上線,主講宇宙學N體模擬在異構計算...
    並使用包括混合精度計算、通信優化在內的方法對CUBE進行了性能優化,可以在較低的內存消耗下實現較高的計算效率和可擴展性。1月6日,智東西公開課推出NVIDIA GPU 加速科學計算專場,主題為《宇宙學N體模擬在異構計算平臺的移植和優化研究》,由上海交通大學計算機系碩士研究生胡航以視頻直播形式進行講解。胡航的主要研究方向為高性能計算,協助物理與天文學院完成了宇宙學數值模擬應用在π 2.0超算平臺上的移植和優化工作。
  • 中山大學虞志益:異構計算融合通用和專業晶片優勢 引領封裝技術...
    使用通用處理器這個傳統的方法已無法滿足人工智慧的各種應用對爆發的和高計算能力的需求。因此,具有GPU、ASIC、FPGA或其它加速器(Accelerator)等高並行、高密集的計算能力的異構計算持續火熱,而異構計算也將成為支撐先進和以後更複雜AI 應用的必然的選擇。
  • 用於物聯網的新型異構計算平臺和5G通信
    異構計算和5G網絡的集成是複雜物聯網系統和服務的最佳平臺之一。本文先介紹複雜物聯網系統需要高性能異構計算和高速通信系統(如5G)的原因,再闡述異構計算賦予物聯網系統基礎設施的一些使用案例。由於傳統處理器系統的一些缺點,我們可以使用的一種解決方案是異構計算,異構計算能有效地獲取高性能計算能力,且其擴展性好、資源利用率高。25G的需求5G與4G相比,優勢主要體現在速率和延時方 面。
  • GPU成功案例:中山大學GPU高性能計算集群
    中山大學計算科學科研團隊吸引了國內外一批知名學者組成計算科學科研團隊,並成功入選廣東省首批創新科研團隊。該團隊是一支具有國際一流水平的創新科研團隊,其成員在微分方程數值解法、積分方程快速算法、隨機問題計算方法、圖像處理、高維數據處理和機器學習理論等領域做出了國際一流的研究成果。
  • OpenCL 2.0正式發布:異構計算的春天
    七月底,Khronos Group宣布了新一代通用計算標準OpenCL 2.0。如今趁著新一屆超級計算大會SC 2013召開的機會,新標準終於正式公開發布了。OpenCL 2.0是這個免費開放跨平臺並行編程規範的重大進步,引入了增強的執行模型和C11/C++11內存模型子集、同步、原子操作,擁有豐富的算法和編程模式,可輕鬆實現計算加速和更好的性能。
  • 【CTO講堂】如何用Python一門語言通吃高性能並發、GPU計算和深度...
    主持人:數據分析方面,請您談一談如何利用gpu加速數據分析。 黃勝藍:很多人可能會覺得gpu編程是一件很困難的事情,其實當真正接觸了之後就會了解到讓你的程序在gpu上運行其實並不難,真正的難點在於理解gpu的特性從而將最合適的任務交給gpu來做,同時採用合適的優化手段。
  • 飛騰cpu怎麼樣 飛騰cpu簡介及對比評測【詳解】
    市面上大家熟知的CPU處理器大多是外國品牌,但是實際上隨著我國相關技術的改進,也出現了不少諸如飛騰cpu之類表現優異的產品。它們性能參數以及實際使用評測和知名大品牌基本相同,儘管依據專家的某些測評看來,國產處理器的微結構略微遜色,但是這並不妨礙它們的發展前景。今天為大家介紹的就是關於飛騰cpu的五個板塊的內容,包括飛騰cpu的簡介等等。
  • 上遊廠商的升維之路 英特爾用oneAPI打破異構計算的次元壁
    而且隨著算力多元化和異構計算趨勢漸成,未來勢必還會有更多的企業試圖進入半導體領域,如果還是只把目光聚焦在贏得每一場戰役之上,那麼未來,英特爾積累起的資源和優勢必然會在不斷發生的小規模戰役中消耗殆盡。於其相對應的,英特爾的未來也會充滿變數。解鈴還須繫鈴人,既然英特爾今日的挑戰來自於下遊企業的集體上探,那麼想要繼續保持優勢的英特爾也必須完成自身的上探。