OpenCL编程
https://blog.csdn.net/huayunhualuo/article/details/102575789
2008年,蘋(píng)果公司向Khronos Group提交了一份關(guān)于跨平臺(tái)計(jì)算框架的草案,該草案由蘋(píng)果公司開(kāi)發(fā),并與AMD、IBM、Intel和NVIDIA公司合作逐步晚上。這個(gè)跨平臺(tái)計(jì)算框架就是OpenCL。2008年12月8日,OpenCL1.0技術(shù)規(guī)范發(fā)布。2010年6月14日,OpenCL1.1發(fā)布,2011年11月19日,OpenCL1.2發(fā)布,2013年11月19日,OpenCL2.0分布。
OpenCL是一個(gè)異構(gòu)并行計(jì)算平臺(tái)編寫(xiě)程序的工作標(biāo)準(zhǔn),此異構(gòu)計(jì)算可映射到CPU、GPU、DSP和FPGA等計(jì)算設(shè)備。OpenCL提供了底層硬件結(jié)構(gòu)的抽象模型,旨在提供一個(gè)通用的開(kāi)發(fā)的API。開(kāi)發(fā)人員可以編寫(xiě)在GPU上運(yùn)行的通用計(jì)算程序,而無(wú)需將其算法映射到OpenGL或DirectX的3D圖形的API上。
為了描述OpenCL設(shè)計(jì)的核心,Khronos Group將OpenCL異構(gòu)并行計(jì)算架構(gòu)劃分為平臺(tái)模型(platform model)、存儲(chǔ)器模型(memory model)、執(zhí)行模型(execution model)和編程模型(programming model)。這些模型相互獨(dú)立、有相互聯(lián)系,組成了OpenCL的有機(jī)整體。
OpenCL平臺(tái)模型:
平臺(tái)模型是關(guān)于OpenCL如何看待硬件的一個(gè)抽象描述。OpenCL平臺(tái)模型由主機(jī)以及相連的一個(gè)或多個(gè)OpenCL設(shè)備組成。如下圖所示,通常主機(jī)是包含X86或ARM處理器的計(jì)算平臺(tái)。OpenCL設(shè)計(jì)可以是CPU、GPU、DSP、FPGA或硬件商提供,OpenCL開(kāi)發(fā)商支持的任何其他處理器。每個(gè)OpenCL設(shè)備有一個(gè)或者多個(gè)計(jì)算單元(Compute Units,CU),而每個(gè)計(jì)算單元又可以由一個(gè)或多個(gè)單元(Processing Elements,PE)組成,處理單元是設(shè)備上執(zhí)行數(shù)據(jù)計(jì)算的最小單元
由于OpenCL的平臺(tái)模型包括了至少兩種處理器,如何連接著兩種處理器就和在這兩種處理器之間傳輸信息的性能密切相關(guān),目前主要通過(guò)PCI-e總線和主機(jī)相連接。
OpenCL執(zhí)行模型:
OpenCL程序包含主機(jī)端程序和設(shè)備端內(nèi)核(kernel)程序。主機(jī)端程序運(yùn)行在主機(jī)處理器上,主機(jī)端程序以命令方式將內(nèi)核程序從主機(jī)提交到OpenCL設(shè)備,OpenCL設(shè)備在處理單元上執(zhí)行計(jì)算。
內(nèi)核在OpenCL設(shè)備上執(zhí)行,完成OpenCL應(yīng)用的具體工作,內(nèi)核通常是一些計(jì)算量大,邏輯比較簡(jiǎn)單的函數(shù),OpenCL設(shè)備通過(guò)內(nèi)核將輸入數(shù)據(jù)就散處理后輸出到主機(jī)。在OpenCL中定義了三類內(nèi)核:
- OpenCL內(nèi)核:用OpenCL C編程語(yǔ)言編寫(xiě),用OpenCL C編譯器編譯的函數(shù),所有OpenCL實(shí)現(xiàn)都必須支持OpenCL內(nèi)核和OpenCL C編程語(yǔ)言。
- 原生內(nèi)核:OpenCL之外創(chuàng)建的函數(shù),在OpenCL中可以通過(guò)一個(gè)函數(shù)指針來(lái)訪問(wèn),執(zhí)行原生內(nèi)核是OpenCL的一個(gè)可選功能,原生內(nèi)核的語(yǔ)義依賴于具體OpenCL實(shí)現(xiàn)。
- 內(nèi)建內(nèi)核:被綁定到特定設(shè)備,并不需要編碼編譯成程序?qū)ο蟮暮瘮?shù)。常見(jiàn)用法是針對(duì)公開(kāi)固定函數(shù)硬件或固件,將它們關(guān)聯(lián)到一個(gè)特定的OpenCL設(shè)備或自定義設(shè)備。內(nèi)建內(nèi)核OpenCL擴(kuò)展功能,內(nèi)建內(nèi)核語(yǔ)義依賴于具體OpenCL實(shí)現(xiàn)。
由于OpenCL設(shè)備通常沒(méi)有I/O處理功能,因此I/O操作通常由主機(jī)承擔(dān),這意味著程序開(kāi)始執(zhí)行時(shí),數(shù)據(jù)通常在主機(jī)上,故OpenCL設(shè)備需要從主機(jī)上獲得數(shù)據(jù),在OpenCL設(shè)備計(jì)算完成后,有需要將數(shù)據(jù)從OpenCL設(shè)備復(fù)制回主機(jī)。
對(duì)于OpenCL執(zhí)行模型來(lái)說(shuō),最重要的就是上下文、命令列隊(duì)和內(nèi)核三個(gè)概念。
上下文
OpenCL程序的計(jì)算工作是OpenCL設(shè)備上執(zhí)行的,不過(guò)主機(jī)在OpenCL應(yīng)用中扮演著重要的角色。主機(jī)使用OpenCL API來(lái)創(chuàng)建和管理上下文,內(nèi)核在此上下文中執(zhí)行。包含了:
- 設(shè)備:OpenCL平臺(tái)包含的一個(gè)或多個(gè)設(shè)備
- 內(nèi)核對(duì)象:在OpenCL設(shè)備上運(yùn)行的OpenCL內(nèi)核函數(shù)
- 程序?qū)ο?/strong>:實(shí)現(xiàn)整個(gè)內(nèi)核程序的源代碼和目標(biāo)二進(jìn)制碼
- 存儲(chǔ)器對(duì)象:對(duì)主機(jī)和OpenCL設(shè)備可見(jiàn)的對(duì)象,內(nèi)核執(zhí)行時(shí)操作這些對(duì)象的實(shí)例。
OpenCL是一套跨平臺(tái)計(jì)算框架,對(duì)OpenCL開(kāi)發(fā)人員來(lái)說(shuō),你可能并不知道OpenCL應(yīng)用最終會(huì)在CPU平臺(tái)、GPU平臺(tái)還是FPGA平臺(tái)上運(yùn)行的,只知道目標(biāo)符合OpenCL規(guī)范。解決的方法就是主機(jī)程序根據(jù)上下文中設(shè)備特性,在運(yùn)行時(shí)從代碼中構(gòu)建程序?qū)ο蟆榱嗽谛阅芎推脚_(tái)無(wú)關(guān)之間均衡,OpenCL提供了兩種方式從代碼中構(gòu)建對(duì)象,一種是從程序源代碼中構(gòu)建,另外一種是從源代碼中已經(jīng)編譯好的代碼上構(gòu)建。
命令隊(duì)列:
OpenCL沒(méi)有定義主機(jī)代碼如何工作的具體細(xì)節(jié),知識(shí)定義了它通過(guò)命令隊(duì)列與OpenCL設(shè)備如何交互。命令隊(duì)列有主機(jī)或運(yùn)行在設(shè)備中的內(nèi)核(該功能需要支持OpenCL2.0的設(shè)備)提交給命令隊(duì)列。命令會(huì)在命令隊(duì)列中等待,直至別調(diào)度到OpenCL設(shè)備上執(zhí)行。放入命令隊(duì)列中的命令分為下列三種類型:
- 內(nèi)核入隊(duì)命令:將一個(gè)內(nèi)核隊(duì)列關(guān)聯(lián)到同一個(gè)OpenCL設(shè)備的命令隊(duì)列中。
- 存儲(chǔ)器入隊(duì)命令:將在主機(jī)和設(shè)備內(nèi)存對(duì)象間傳輸數(shù)據(jù),或者將內(nèi)存對(duì)象映射到主機(jī)地址空間,或從主機(jī)地址空間取消映射提交給命令隊(duì)列。
- 同步命令:對(duì)命令隊(duì)列中需要執(zhí)行的命令施加執(zhí)行順序約束,如只有某個(gè)命令執(zhí)行完成其他命令才能開(kāi)始執(zhí)行。
命令可以異步方式執(zhí)行,在這種方式下主機(jī)或運(yùn)行在設(shè)備中內(nèi)核向命令隊(duì)列提交命令,然后繼續(xù)工作,而不必等待命令完成。如果有必要等待一個(gè)命令完成,可以利用命令執(zhí)行相關(guān)的同步機(jī)制進(jìn)行同步。
一個(gè)命令隊(duì)列中命令執(zhí)行時(shí)可以有以下兩種模式:
- 按序(in-order)執(zhí)行:命令按其排入命令隊(duì)列中的先后順序執(zhí)行,并按順序完成。
- 亂序(out-of-order)執(zhí)行:命令以任意順序執(zhí)行,通過(guò)顯示的同步點(diǎn)或顯示事件依賴項(xiàng)來(lái)約束順序。
OpenCL平臺(tái)都支持按序模式,但亂序模式是可選的。
與內(nèi)核相關(guān)的計(jì)算,在運(yùn)行時(shí)只作為內(nèi)核實(shí)例執(zhí)行。常規(guī)方法上,可以通過(guò)主機(jī)程序多次啟動(dòng)內(nèi)核實(shí)例來(lái)執(zhí)行,但這樣會(huì)顯著地增加開(kāi)銷或加重應(yīng)用程序控制流。一個(gè)方便有效的方法就是從內(nèi)核內(nèi)部嵌套內(nèi)核命令隊(duì)列。對(duì)支持OpenCL2.0的設(shè)備,可以在設(shè)備上入隊(duì)內(nèi)核,不需要主機(jī)程序參與,這稱為設(shè)備端隊(duì)列,實(shí)現(xiàn)了設(shè)備端隊(duì)列,實(shí)現(xiàn)了嵌套并行。
內(nèi)核在OpenCL設(shè)備上執(zhí)行:
主機(jī)發(fā)出一個(gè)命令,提交一個(gè)內(nèi)核到OpenCL設(shè)備上執(zhí)行,OpenCL運(yùn)行時(shí)將會(huì)創(chuàng)建一個(gè)整數(shù)索引空間。索引空間是OpenCL支持的一個(gè)N維值的網(wǎng)格,稱為NDRange,其中N為1,2,或3。三個(gè)長(zhǎng)度為N的數(shù)據(jù)確定了NDRange的一下特征。
- 每個(gè)維度索引空間(或全局大小)的范圍
- 一個(gè)偏移指數(shù)F表明每個(gè)維度的初始索引值(默認(rèn)為0)
- 一個(gè)工作組(局部大小)每個(gè)維度大小。
OpenCL存儲(chǔ)器模型:
OpenCL異構(gòu)平臺(tái)由主機(jī)端和設(shè)計(jì)端構(gòu)成,存儲(chǔ)器區(qū)域包括主機(jī)和設(shè)備的內(nèi)存。
- 主機(jī)內(nèi)存(host memory):主機(jī)直接可用的內(nèi)存,OPenCL并未定義主機(jī)內(nèi)存的具體的行為。通過(guò)OpenCL API或者共享虛擬存儲(chǔ)接口,實(shí)現(xiàn)存儲(chǔ)器對(duì)象在主機(jī)和設(shè)備之間的傳輸。
- 全局存儲(chǔ)器(global memory):這個(gè)存儲(chǔ)器區(qū)域允許上下文任何設(shè)備中所有工作組的所有工作項(xiàng)的讀寫(xiě),工作項(xiàng)可以讀寫(xiě)存儲(chǔ)器對(duì)象中的任意元素。全局存儲(chǔ)器的讀寫(xiě)能力可能會(huì)被緩存,取決于設(shè)備的能力。
- 常量存儲(chǔ)器(constant memory):全局存儲(chǔ)器中的一塊區(qū)域,在內(nèi)核實(shí)例執(zhí)行期間其保存的數(shù)據(jù)保持不變。主機(jī)負(fù)責(zé)該存儲(chǔ)器對(duì)象的分配和初始化。
- 局部存儲(chǔ)器(lcoal memory):這個(gè)存儲(chǔ)器區(qū)域?qū)ぷ鹘M是局部可見(jiàn)的,用來(lái)分配該工作組中所有工作項(xiàng)共享的變量。
- 私有存儲(chǔ)器(private memory):這個(gè)存儲(chǔ)器區(qū)域是一個(gè)工作項(xiàng)的私有區(qū)域。
存儲(chǔ)器對(duì)象
全局存儲(chǔ)器中的數(shù)據(jù)內(nèi)容通過(guò)存儲(chǔ)器對(duì)象來(lái)表示,一個(gè)存儲(chǔ)器對(duì)象就是對(duì)全局存儲(chǔ)器區(qū)域的一個(gè)引用,主要有三種不同的類型:
- 緩沖(buffer):內(nèi)核可用一個(gè)連續(xù)的存儲(chǔ)器區(qū)域,編程人員可以將內(nèi)建數(shù)據(jù)類型、矢量數(shù)據(jù)類型或用戶自定義的數(shù)據(jù)結(jié)構(gòu)(符合OpenCL編程規(guī)范)映射到這個(gè)緩沖區(qū),,內(nèi)核通過(guò)指針來(lái)訪問(wèn)緩存區(qū)。
- 圖像(image):圖像對(duì)象用戶存儲(chǔ)基于標(biāo)準(zhǔn)格式的圖像,圖像對(duì)象是一個(gè)不透明的數(shù)據(jù)結(jié)構(gòu),使用OpenCL API函數(shù)來(lái)管理。通常不允許OpenCL內(nèi)核對(duì)單個(gè)圖像同時(shí)進(jìn)行讀和寫(xiě)。在OpenCL2.0之后,提供了同步和柵欄操作來(lái)放寬這個(gè)限制。
- 管道(pipe):管道存儲(chǔ)器是數(shù)據(jù)項(xiàng)有序的隊(duì)列,管道有兩個(gè)端點(diǎn):一個(gè)是寫(xiě)端點(diǎn),用于插入數(shù)據(jù)項(xiàng),另外一個(gè)讀端點(diǎn),數(shù)據(jù)項(xiàng)從讀端點(diǎn)被移除。同一個(gè)時(shí)刻,僅有一個(gè)內(nèi)核實(shí)例可向一個(gè)管道寫(xiě)入數(shù)據(jù),同時(shí)僅有一個(gè)內(nèi)核實(shí)例從一個(gè)管道讀出數(shù)據(jù)。
主機(jī)和設(shè)備間有三種交互方式:讀、寫(xiě)、映射和解映射以及拷貝。
OpenCL編程基礎(chǔ):
OpenCL編程中重要的概念:
- Host:CPU以及系統(tǒng)的內(nèi)存統(tǒng)稱為Host,我們?cè)贖ost上定義控制OpenCl的環(huán)境與執(zhí)行邏輯。
- Device:將GPU以及GPU本身的顯示內(nèi)存稱為Device,在Host上定義的計(jì)算邏輯最終要GPU進(jìn)行計(jì)算。
- Kernel:OpenCl程序分為Host(CPU)和Device(GPU)兩個(gè)部分,主程序由CPU執(zhí)行,當(dāng)需要并行計(jì)算數(shù)據(jù)時(shí),OpenCL就會(huì)將程序編譯成GPU能執(zhí)行的程序,發(fā)送給GPU,這個(gè)程序就是Kernel。Kernel是Device程序執(zhí)行的入口點(diǎn)。
- SIMT:單指令多線程(SINGLE INSTRUCTION MULTI THREAD)的簡(jiǎn)寫(xiě),相同的代碼在不同線程中并行執(zhí)行,每個(gè)線程使用不同的數(shù)據(jù)來(lái)執(zhí)行同一段代碼。
- Work-item:工作項(xiàng),與CUDA中的Thread是一樣的,是最小執(zhí)行單元。一個(gè)GPU并行程序會(huì)被許多個(gè)threads來(lái)執(zhí)行。
- Work-group:工作組,與CUDA中Block是一樣的。Work-group的存在是為了允許work-item之間的通信與協(xié)作。數(shù)個(gè)Work-item會(huì)被群組成一個(gè)Work-group,同一個(gè)Work-group中的Work-item可以同步并通過(guò)shared memory通信。(work-group是以N維網(wǎng)格形式組織的,N=1,2或3)
- ND-Range:ND-Range是Work-group下一個(gè)組織級(jí)別,定義了work-group的組織形式,與CUDA中的grid是一樣的。(ND-Rang以N維網(wǎng)格形式組織的,N=1,2或3)
- Global memory:通俗意義上的設(shè)備內(nèi)存。
- Constant memory:位于設(shè)備內(nèi)存,常量?jī)?nèi)存用于保存在核函數(shù)執(zhí)行期間不會(huì)發(fā)生變化的數(shù)據(jù)。變量的訪問(wèn)限制為只讀。常量?jī)?nèi)存采取了不同于標(biāo)準(zhǔn)全局內(nèi)存的處理方式。在某些情況下,用常量?jī)?nèi)存替換全局內(nèi)存能有效地減少內(nèi)存帶寬,從而提升性能。
- Local memory:局部?jī)?nèi)存,隸屬于一個(gè)工作組的內(nèi)存區(qū)域。它可以用來(lái)分配一些變量, 這些變量由此工作組中的所有工作項(xiàng)共享。在OpenCL設(shè)備上,可能會(huì)將 其實(shí)現(xiàn)成一塊專有的內(nèi)存區(qū)域,也可能將其映射到全局內(nèi)存中。
- Private memory:私有內(nèi)存,隸屬于一個(gè)工作項(xiàng)的內(nèi)存區(qū)域。 一個(gè)工作項(xiàng)的私有內(nèi)存中所定義的變量對(duì)另外一個(gè)工作項(xiàng)來(lái)說(shuō)是不可見(jiàn)的
注意點(diǎn):
- 核函數(shù)是以Work-group為單位執(zhí)行
- Host程序中包括:Plantform(平臺(tái))、Context(上下文,即OpenCl環(huán)境,包括OpenCL kernel、設(shè)備、內(nèi)存管理、命令隊(duì)列等)、Command-Queue(指令隊(duì)列,多個(gè)指令隊(duì)列允許應(yīng)用程序在不需要同步的情況下執(zhí)行多條無(wú)關(guān)聯(lián)的指令)
- Work-item、Work-group和ND-range都使用一維、二維和三維的索引進(jìn)行標(biāo)識(shí)
OpenCL編程入門(mén):
OpenCL編程的步驟。 首先給出實(shí)現(xiàn)流程,然后給出實(shí)現(xiàn)圖像旋轉(zhuǎn)的C循環(huán)實(shí)現(xiàn)和OpenCL C kernel實(shí)現(xiàn)。
OpenCL C kernel代碼:
#pragma OPENCL EXTENSION cl_amd_printf : enable __kernel void image_rotate(__global uchar * src_data,__global uchar * dest_data, //Data in global memoryint W, int H, //Image Dimensionsfloat sinTheta, float cosTheta ) //Rotation Parameters {const int ix = get_global_id(0);const int iy = get_global_id(1);int xc = W/2;int yc = H/2;int xpos = ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc;int ypos = (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc;if ((xpos>=0) && (xpos< W) && (ypos>=0) && (ypos< H))dest_data[ypos*W+xpos]= src_data[iy*W+ix]; }?
?
總結(jié)
- 上一篇: 注册.NET Framework
- 下一篇: JAVA多线程之先行发生原则