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