AMD OpenCL大學課程是非常好的入門級OpenCL教程，通過看教程中的PPT，我們能夠很快的了解OpenCL機制以及編程方法。下載地址：http://developer.amd.com/zones/OpenCLZone/universities/Pages/default.aspx ???? 教程中的英文很簡單，我相信學OpenCL的人都能看得懂，而且看原汁原味的英文表述，更有利于我們了解各種術語的來龍去脈。 ???? 我把這些教程翻譯成自己的中文表述，主要是強化理解需要，其實我的英文很爛。 一、并行計算概述 ??? 在計算機術語中，并行性是指：把一個復雜問題,分解成多個能同時處理的子問題的能力。要實現并行計算，首先我們要有物理上能夠實現并行計算的硬件設備，比如多核CPU，每個核能同時實現算術或邏輯運算。 ??? 通常，我們通過GPU實現兩類并行計算： ????? 任務并行：把一個問題分解為能夠同時執行的多個任務。 ????? 數據并行：同一個任務內，它的各個部分同時執行。 ?? 下面我們通過一個農場主雇傭工人摘蘋果的例子來描述不同種類的并行計算。

• 摘蘋果的工人就是硬件上的并行處理單元（process elements)。
• 樹就是要執行的任務。
• 蘋果就是要處理的數據。

? 串行的任務處理就如下圖所示，一個工人背著梯子摘完所有樹上的蘋果（一個處理單元處理完所有任務的數據）。

? ? 數據并行就好比農場主雇傭了好多工人來摘完一個樹上的蘋果（多個處理單元并行完成一個任務中的數據），這樣就能很快摘完一顆樹上的蘋果。

?? 農場主也可以為每棵樹安排一個工人，這就好比任務并行。在每個任務內，由于只有一個工人，所以是串行執行的，但任務之間是并行的。

對一個復雜問題，影響并行計算的因素很多。通常，我們都是通過分解問題的方式來實施并算法行。

這又包括兩方面內容：

• 任務分解：把算法分解成很多的小任務，就像前面的例子中，把果園按蘋果樹進行劃分，這時我們并不關注數據，也就是說不關注每個樹上到底有多少個蘋果。
• 數據分解：就是把很多數據，分成不同的、離散的小塊，這些數據塊能夠被并行執行，就好比前面例子中的蘋果。

?? 通常我們按照算法之間的依賴關系來分解任務，這樣就形成了一個任務關系圖。一個任務只有沒有依賴任務的時候，才能夠被執行。

??? 這有點類似于數據結構中的有向無環圖，兩個沒有連通路徑的任務之間可以并行執行。下面再給一個烤面包的例子，如果所示，預熱烤箱和購買面粉糖兩個任務之間可以并行執行。

??

???? 對大多數科學計算和工程應用來說，數據分解一般都是基于輸出數據,例如：

• 在一副圖像中，對一個滑動窗口（例如：3*3像素)內的像素實施濾波操作，可以得到一個輸出像素的卷積。
• 第一個輸入矩陣的第i行乘以第二個輸入矩陣的第j列，得到的向量和即為輸出矩陣第i行，第j列的元素。

這種方法對于輸入和輸出數據是一對一，或者多對一的對應關系比較有效。

??? 也有的數據分解算法是基于輸入數據的，這時，輸入數據和輸出數據一般是一對多的關系，比如求圖像的直方圖，我們要把每個像素放到對應的槽中(bins，對于灰度圖，bin數量通常是256）。一個搜索函數，輸入可能是多個數據，輸出卻只有一個值。對于這類應用，我們一般用每個線程計算輸出的一部分，然后通過同步以及原子操作得到最終的值，OpenCL中求最小值的kernel函數就是典型代表[可以看下ATI Stream Computing OpenCL programming guide第二章中求最小值的kernel例子]。

? ?? 通常來說，怎樣分解問題和具體算法有關，而且還要考慮自己使用的硬件和軟件，比如AMD GPU平臺和Nvdia GPU平臺的優化就有很多不同。

二、常用基于硬件和軟件的并行

??? 在上個實際90年代，并行計算主要研究如何在cpu上實施指自動的指令級并行。

• 同時發射多條指令（之間沒有依賴關系），并行執行這些指令。
• 在本教程中，我么不講述自動的硬件級并行，感興趣的話，可以看看計算機體系結構的教程。

??? 高層的并行，比如線程級別的并行，一般很難自動化，需要程序員告訴計算機，該做什么，不該做什么。這時，程序員還要考慮硬件的具體指標，通常特定硬件都是適應于某一類并行編程，比如多核cpu就適合基于任務的并行編程，而GPU更適應于數據并行編程。

Hardware type	Examples	Parallelism
Multi-core superscalar processors	Phenom II CPU	Task
Vector or SIMD processors	SSE units (x86 CPUs)	Data
Multi-core SIMD processors	Radeon 5870 GPU	Data

?

??? 現代的GPU有很多獨立的運算核（processor）組成，在AMD GPU上就是stream core，這些core能夠執行SIMD操作（單指令，多數據），所以特別適合數據并行操作。通常GPU上執行一個任務，都是把任務中的數據分配到各個獨立的core中執行。

???? 在GPU上，我們一般通過循環展開，Loop strip mining 技術，來把串行代碼改成并行執行的。比如在CPU上，如果我們實現一個向量加法，代碼通常如下：

1: for(i = 0; i < n; i++) 2: { 3: C[i] = A[i] + B[i]; 4: }

在GPU上，我們可以設置n個線程，每個線程執行一個加法，這樣大大提高了向量加法的并行性。

1: __kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int n) 2: { 3: int i = get_global_id(0); 4: c[i] = a[i] + b[i]; 5: }

??? 上面這個圖展示了向量加法的SPMD（單指令多線程）實現，從圖中可以看出如何實施Loop strip mining 操作的。

??? GPU的程序一般稱作Kernel程序，它是一種SPMD的編程模型（the Single Program Multiple Data ）。SPMD執行同一段代碼的多個實例，每個實例對數據的不同部分進行操作。

???? 在數據并行應用中，用loop strip mining來實現SPMD是最常用的方法：

• • 在分布式系統中，我們用Message Passing Interface (MPI)來實現SPMD。
  • 在共享內存并行系統中，我們用POSIX線程來實現SPMD。
  • 在GPU中，我們就是用Kernel來顯現SPMD。

??? 在現代的CPU上，創建一個線程的開銷還是很大的，如果要在CPU上實現SPMD，每個線程處理的數據塊就要盡量大點，做更多的事情，以便減少平均線程開銷。但在GPU上，都是輕量級的線程，創建、調度線程的開銷比較小，所以我們可以做到把循環完全展開，一個線程處理一個數據。

GPU上并行編程的硬件一般稱作SIMD。通常，發射一條指令后，它要在多個ALU單元中執行（ALU的數量即使simd的寬度），這種設計減少了控制流單元以級ALU相關的其他硬件數量。

SIMD的硬件如下圖所示：

?

?

??? 在向量加法中，寬度為4的SIMD單元，可以把整個循環分為四個部分同時執行。在工人摘蘋果的例子中，工人的雙手類似于SIMD的寬度為2。另外，我們要知道，現在的GPU硬件上都是基于SIMD設計，GPU硬件隱式的把SPMD線程映射到SIMD core上。對開發有人員來說，我們并不需要關注硬件執行結果是否正確，我們只需要關注它的性能就OK了。

??? CPU一般都支持并行級的原子操作，這些操作保證不同的線程讀寫數據，相互之間不會干擾。有些GPU支持系統范圍的并行操作，但會有很大開銷，比如Global memory的同步。

1、OpenCL架構

?? OpenCL可以實現混合設備的并行計算，這些設備包括CPU，GPU，以及其它處理器，比如Cell處理器，DSP等。使用OpenCL編程，可以實現可移植的并行加速代碼。[但由于各個OpenCL device不同的硬件性能，可能對于程序的優化還要考慮具體的硬件特性]。

?? 通常OpenCL架構包括四個部分：

• 平臺模型(Platform Model)
• 執行模型(Execution Model)
• 內存模型(Memory Model)
• 編程模型(Programming Model)

2、OpenCL平臺模型

?? 不同廠商的OpenCL實施定義了不同的OpenCL平臺，通過OpenCL平臺，主機能夠和OpenCL設備之間進行交互操作。現在主要的OpenCL平臺有AMD、Nvida，Intel等。OpenCL使用了一種Installable Client Driver模型，這樣不同廠商的平臺就能夠在系統中共存。在我的計算機上就安裝有AMD和Intel兩個OpenCL Platform[現在的OpenCL driver模型不允許不同廠商的GPU同時運行]。

??? OpenCL平臺通常包括一個主機(Host)和多個OpenCL設備（device），每個OpenCL設備包括一個或多個CU(compute units)，每個CU包括又一個或多個PE（process element）。 每個PE都有自己的程序計數器(PC)。主機就是OpenCL運行庫宿主設備，在AMD和Nvida的OpenCL平臺中，主機一般都指x86 CPU。

?? 對AMD平臺來說，所有的CPU是一個設備，CPU的每一個core就是一個CU，而每個GPU都是獨立的設備。

??

3、OpenCL編程的一般步驟

? 下面我們通過一個實例來了解OpenCL編程的步驟，假設我們用的是AMD OpenCL平臺（因為本人的GPU是HD5730），安裝了AMD Stream SDK 2.6,并在VS2008中設置好了include，lib目錄等。

??? 首先我們建立一個控制臺程序，最初的代碼如下：

1: #include "stdafx.h" 2: #include <CL/cl.h> 3: #include <stdio.h> 4: #include <stdlib.h> 5:? 6: #pragma comment (lib,"OpenCL.lib") 7:? 8: int main(int argc, char* argv[]) 9: { 10: return 0; 11: }

?

第一步，我們要選擇一個OpenCL平臺，所用的函數就是

??? 通常，這個函數要調用2次，第一次得到系統中可使用的平臺數目，然后為（Platform）平臺對象分配空間，第二次調用就是查詢所有的平臺，選擇自己需要的OpenCL平臺。代碼比較長，具體可以看下AMD Stream SDK 2.6中的TemplateC例子，里面描述如何構建一個robust的最小OpenCL程序。為了簡化代碼，使程序看起來不那么繁瑣，我直接調用該函數，選取系統中的第一個OpenCL平臺，我的系統中安裝AMD和Intel兩家的平臺，第一個平臺是AMD的。另外，我也沒有增加錯誤檢測之類的代碼，但是增加了一個status的變量，通常如果函數執行正確，返回的值是0。

1: #include "stdafx.h" 2: #include <CL/cl.h> 3: #include <stdio.h> 4: #include <stdlib.h> 5:? 6: #pragma comment (lib,"OpenCL.lib") 7:? 8: int main(int argc, char* argv[]) 9: { 10: cl_uint status; 11: cl_platform_id platform; 12:? 13: status = clGetPlatformIDs( 1, &platform, NULL ); 14:? 15: return 0; 16: }

第二步是得到OpenCL設備，

???? 這個函數通常也是調用2次，第一次查詢設備數量，第二次檢索得到我們想要的設備。為了簡化代碼，我們直接指定GPU設備。

?

1: #include "stdafx.h" 2: #include <CL/cl.h> 3: #include <stdio.h> 4: #include <stdlib.h> 5:? 6: #pragma comment (lib,"OpenCL.lib") 7:? 8: int main(int argc, char* argv[]) 9: { 10: cl_uint status; 11: cl_platform_id platform; 12:? 13: status = clGetPlatformIDs( 1, &platform, NULL ); 14:? 15: cl_device_id device; 16:? 17: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 18: 1, 19: &device, 20: NULL); 21:? 22: return 0; 23: }

下面我們來看下OpenCL中Context的概念：

通常，Context是指管理OpenCL對象和資源的上下文環境。為了管理OpenCL程序，下面的一些對象都要和Context關聯起來：

?

—設備（Devices）:執行Kernel程序對象。

—程序對象（Program objects）: kernel程序源代碼

—Kernels:運行在OpenCL設備上的函數。

—內存對象（Memory objects）: device處理的數據對象。

—命令隊列（Command queues）: 設備之間的交互機制。

• ?

注意：創建一個Context的時候，我們必須把一個或多個設備和它關聯起來。對于其它的OpenCL資源，它們創建時候，也要和Context關聯起來，一般創建這些資源的OpenCL函數的輸入參數中，都會有context。

這個函數中指定了和context關聯的一個或多個設備對象,properties參數指定了使用的平臺，如果為NULL,廠商選擇的缺省值被使用，這個函數也提供了一個回調機制給用戶提供錯誤報告。

現在的代碼如下：

1: #include "stdafx.h" 2: #include <CL/cl.h> 3: #include <stdio.h> 4: #include <stdlib.h> 5:? 6: #pragma comment (lib,"OpenCL.lib") 7:? 8: int main(int argc, char* argv[]) 9: { 10: cl_uint status; 11: cl_platform_id platform; 12:? 13: status = clGetPlatformIDs( 1, &platform, NULL ); 14:? 15: cl_device_id device; 16:? 17: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 18: 1, 19: &device, 20: NULL); 21: cl_context context = clCreateContext( NULL, 22: 1, 23: &device, 24: 25:? 26: return 0; 27: }

接下來，我們要看下命令隊列。在OpenCL中，命令隊列就是主機的請求，在設備上執行的一種機制。

• 在Kernel執行前，我們一般要進行一些內存拷貝的工作，比如把主機內存中的數據傳輸到設備內存中。

另外要注意的幾點就是：對于不同的設備，它們都有自己的獨立的命令隊列；命令隊列中的命令(kernel函數）可能是同步的，也可能是異步的，它們的執行順序可以是有序的，也可以是亂序的。

命令隊列在device和context之間建立了一個連接。

命令隊列properties指定以下內容:

• 是否亂序執行（在AMD GPU中，好像現在還不支持亂序執行）
• 是否啟動profiling。Profiling通過事件機制來得到kernel執行時間等有用的信息，但它本身也會有一些開銷。

?

如下圖所示，命令隊列把設備和context聯系起來，盡管它們之間不是物理連接。

添加命令隊列后的代碼如下：

1: #include "stdafx.h" 2: #include <CL/cl.h> 3: #include <stdio.h> 4: #include <stdlib.h> 5:? 6: #pragma comment (lib,"OpenCL.lib") 7:? 8: int main(int argc, char* argv[]) 9: { 10: cl_uint status; 11: cl_platform_id platform; 12:? 13: status = clGetPlatformIDs( 1, &platform, NULL ); 14:? 15: cl_device_id device; 16:? 17: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 18: 1, 19: &device, 20: NULL); 21: cl_context context = clCreateContext( NULL, 22: 1, 23: &device, 24: NULL, NULL, NULL); 25:? 26: cl_command_queue queue = clCreateCommandQueue( context, 27: device, 28: CL_QUEUE_PROFILING_ENABLE, NULL ); 29:? 30: return 0; 31: }

?

OpenCL內存對象：

??? OpenCL內存對象就是一些OpenCL數據，這些數據一般在設備內存中，能夠被拷入也能夠被拷出。OpenCL內存對象包括buffer對象和image對象。

buffer對象：連續的內存塊----順序存儲，能夠通過指針、行列式等直接訪問。

image對象：是2維或3維的內存對象，只能通過read_image() 或 write_image()來讀取。image對象可以是可讀或可寫的，但不能同時既可讀又可寫。

??? 該函數會在指定的context上創建一個buffer對象，image對象相對比較復雜，留在后面再講。

flags參數指定buffer對象的讀寫屬性，host_ptr可以是NULL，如果不為NULL，一般是一個有效的host buffer對象，這時，函數創建OpenCL buffer對象后，會把對應host buffer的內容拷貝到OpenCL buffer中。

???? 在Kernel執行之前，host中原始輸入數據必須顯式的傳到device中，Kernel執行完后，結果也要從device內存中傳回到host內存中。我們主要通過函數clEnqueue{Read|Write}{Buffer|Image}來實現這兩種操作。從host到device，我們用clEnqueueWrite，從device到host，我們用clEnqueueRead。clEnqueueWrite命令包括初始化內存對象以及把host 數據傳到device內存這兩種操作。當然，像前面一段說的那樣，也可以把host buffer指針直接用在CreateBuffer函數中來實現隱式的數據寫操作。

???? 這個函數初始化OpenCL內存對象，并把相應的數據寫到OpenCL內存關聯的設備內存中。其中，blocking_write參數指定是數拷貝完成后函數才返回還是數據開始拷貝后就立即返回（阻塞模式于非阻塞模式）。Events參數指定這個函數執行之前，必須要完成的Event（比如先要創建OpenCL內存對象的Event）。

?

OpenCL程序對象：

?? 程序對象就是通過讀入Kernel函數源代碼或二進制文件，然后在指定的設備上進行編譯而產生的OpenCL對象。

???? 這個函數通過源代碼(strings)，創建一個程序對象，其中counts指定源代碼串的數量，lengths指定源代碼串的長度（為NULL結束的串時，可以省略）。當然，我們還必須自己編寫一個從文件中讀取源代碼串的函數。

???? 對context中的每個設備，這個函數編譯、連接源代碼對象，產生device可以執行的文件，對GPU而言就是設備對應shader匯編。如果device_list參數被提供，則只對這些設備進行編譯連接。options參數主要提供一些附加的編譯選項，比如宏定義、優化開關標志等等。

???? 如果程序編譯失敗，我們能夠根據返回的狀態，通過調用clGetProgramBuildInfo來得到錯誤信息。

加上創建內存對象以及程序對象的代碼如下:

1:? 2: #include "stdafx.h" 3: #include <CL/cl.h> 4: #include <stdio.h> 5: #include <stdlib.h> 6: #include <time.h> 7: #include <iostream> 8: #include <fstream> 9:? 10: using namespace std; 11: #define NWITEMS 262144 12:? 13: #pragma comment (lib,"OpenCL.lib") 14:

Kernel對象：

??? Kernel就是在程序代碼中的一個函數，這個函數能在OpenCL設備上執行。一個Kernel對象就是kernel函數以及其相關的輸入參數。

?

Kernel對象通過程序對象以及指定的函數名字創建。注意：函數必須是程序源代碼中存在的函數。

運行時編譯：

??? 在運行時，編譯程序和創建kernel對象是有時間開銷的，但這樣比較靈活，能夠適應不同的OpenCL硬件平臺。程序動態編譯一般只需一次，而Kernel對象在創建后，可以反復調用。

?

創建Kernel后，運行Kernel之前，我們還要為Kernel對象設置參數。我們可以在Kernel運行后，重新設置參數再次運行。

arg_index指定該參數為Kernel函數中的第幾個參數(比如第一個參數為0，第二個為1,…)。內存對象和單個的值都可以作為Kernel參數。下面是2個設置Kernel參數的例子：

clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&d_iImage);

clSetKernelArg(kernel, 1, sizeof(int), (void*)&a);

在Kernel運行之前，我們先看看OpenCL中的線程結構：

大規模并行程序中，通常每個線程處理一個問題的一部分，比如向量加法，我們會把兩個向量中對應的元素加起來，這樣，每個線程可以處理一個加法。

下面我看一個16個元素的向量加法：兩個輸入緩沖A、B，一個輸出緩沖C

在這種情況下，我們可以創建一維的線程結構去匹配這個問題。

每個線程把自己的線程id作為索引，把相應元素加起來。

???? OpenCL中的線程結構是可縮放的，Kernel的每個運行實例稱作WorkItem(也就是線程），WorkItem組織在一起稱作WorkGroup，OpenCL中，每個Workgroup之間都是相互獨立的。

通過一個global id(在索引空間，它是唯一的）或者一個workgroup id和一個work group內的local id，我就能標定一個workitem。

在kernel函數中，我們能夠通過API調用得到global id以及其他信息：

get_global_id(dim)

get_global_size(dim)

這兩個函數能得到每個維度上的global id。

get_group_id(dim)

get_num_groups(dim)

get_local_id(dim)

get_local_size(dim）

這幾個函數用來計算group id以及在group內的local id。

get_global_id(0) = column, get_global_id(1) = row

get_num_groups(0) * get_local_size(0) == get_global_size(0)

?

OpenCL內存模型

??? OpenCL的內存模型定義了各種各樣內存類型，各種內存模型之間有層級關系。各種內存之間的數據傳輸必須是顯式進行的，比如從host memory到device memory，從global memory到local memory等等。

??? WorkGroup被映射到硬件的CU上執行（在AMD 5xxx系列顯卡上，CU就是simd，一個simd中有16個pe，或者說是stream core），OpenCL并不提供各個workgroup之間的一致性，如果我們需要在各個workgroup之間共享數據或者通信之類的，要自己通過軟件實現。

Kernel函數的寫法

每個線程（workitem）都有一個kenerl函數的實例。下面我們看下kernel的寫法：

1: __kernel void vecadd(__global const float* A, __global const float* B, __global float* C) 2: { 3: int id = get_global_id(0); 4: C[id] = A[id] + B[id]; 5: }

每個Kernel函數都必須以__kernel開始，而且必須返回void。每個輸入參數都必須聲明使用的內存類型。通過一些API，比如get_global_id之類的得到線程id。

內存對象地址空間標識符有以下幾種：

__global – memory allocated from global address space

__constant – a special type of read-only memory

__local – memory shared by a work-group

__private – private per work-item memory

__read_only/__write_only – used for images

Kernel函數參數如果是內存對象，那么一定是__global,__local或者constant。

?

運行Kernel

?? 首先要設置線程索引空間的維數以及workgroup大小等。

?? 我們通過函數clEnqueueNDRangeKerne把Kernel放在一個隊列里，但不保證它馬上執行，OpenCL driver會管理隊列，調度Kernel的執行。注意：每個線程執行的代碼都是相同的，但是它們執行數據卻是不同的。

?

?? 該函數把要執行的Kernel函數放在指定的命令隊列中，globald大小（線程索引空間）必須指定，local大小（work group）可以指定，也可以為空。如果為空，則系統會自動根據硬件選擇合適的大小。event_wait_list用來選定一些events，只有這些events執行完后，該kernel才可能被執行，也就是通過事件機制來實現不同kernel函數之間的同步。

?? 當Kernel函數執行完畢后，我們要把數據從device memory中拷貝到host memory中去。

釋放資源：

??? 大多數的OpenCL資源都是指針，不使用的時候需要釋放掉。當然，程序關閉的時候這些對象也會被自動釋放掉。

??? 釋放資源的函數是：clRelase{Resource} ，比如: clReleaseProgram(), clReleaseMemObject()等。

?

錯誤捕捉：

??? 如果OpenCL函數執行失敗，會返回一個錯誤碼，一般是個負值，返回0則表示執行成功。我們可以根據該錯誤碼知道什么地方出錯了，需要修改。錯誤碼在cl.h中定義，下面是幾個錯誤碼的例子.

CL_DEVICE_NOT_FOUND -1

CL_DEVICE_NOT_AVAILABLE -2

CL_COMPILER_NOT_AVAILABLE -3

CL_MEM_OBJECT_ALLOCATION_FAILURE -4

…

下面是一個OpenCL機制的示意圖

程序模型

??? 數據并行：work item和內存對象元素之間是一一映射關系；workgroup可以顯示指定，也可以隱式指定。

??? 任務并行：kernel的執行獨立于線程索引空間；用其他方法表示并行，比如把不同的任務放入隊列，用設備指定的特殊的向量類型等等。

??? 同步：workgroup內work item之間的同步；命令隊列中不同命令之間的同步。

完整代碼如下：

1: #include "stdafx.h" 2: #include <CL/cl.h> 3: #include <stdio.h> 4: #include <stdlib.h> 5: #include <time.h> 6: #include <iostream> 7: #include <fstream> 8:? 9: using namespace std; 10: #define NWITEMS 262144 11:? 12: #pragma comment (lib,"OpenCL.lib") 13:? 14: //把文本文件讀入一個string中 15: int convertToString(const char *filename, std::string& s) 16: { 17: size_t size; 18: char* str; 19:? 20: std::fstream f(filename, (std::fstream::in | std::fstream::binary)); 21:? 22: if(f.is_open()) 23: { 24: size_t fileSize; 25: f.seekg(0, std::fstream::end); 26: size = fileSize = (size_t)f.tellg(); 27: f.seekg(0, std::fstream::beg); 28:? 29: str = new char[size+1]; 30: if(!str) 31: { 32: f.close(); 33: return NULL; 34: } 35:? 36: f.read(str, fileSize); 37: f.close(); 38: str[size] = '\0'; 39: 40: s = str; 41: delete[] str; 42: return 0; 43: } 44: printf("Error: Failed to open file %s\n", filename); 45: return 1; 46: } 47:? 48: int main(int argc, char* argv[]) 49: { 50: //在host內存中創建三個緩沖區 51: float *buf1 = 0; 52: float *buf2 = 0; 53: float *buf = 0; 54: 55: buf1 =(float *)malloc(NWITEMS * sizeof(float)); 56: buf2 =(float *)malloc(NWITEMS * sizeof(float)); 57: buf =(float *)malloc(NWITEMS * sizeof(float)); 58:? 59: //初始化buf1和buf2的內容 60: int i; 61: srand( (unsigned)time( NULL ) ); 62: for(i = 0; i < NWITEMS; i++) 63: buf1[i] = rand()%65535; 64:? 65: srand( (unsigned)time( NULL ) +1000); 66: for(i = 0; i < NWITEMS; i++) 67: buf2[i] = rand()%65535; 68:? 69: for(i = 0; i < NWITEMS; i++) 70: buf[i] = buf1[i] + buf2[i]; 71:? 72: cl_uint status; 73: cl_platform_id platform; 74:? 75: //創建平臺對象 76: status = clGetPlatformIDs( 1, &platform, NULL ); 77:? 78: cl_device_id device; 79:? 80: //創建GPU設備 81: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 82: 1, 83: &device, 84: NULL); 85: //創建context 86: cl_context context = clCreateContext( NULL, 87: 1, 88: &device, 89: NULL, NULL, NULL); 90: //創建命令隊列 91: cl_command_queue queue = clCreateCommandQueue( context, 92: device, 93: CL_QUEUE_PROFILING_ENABLE, NULL ); 94: //創建三個OpenCL內存對象，并把buf1的內容通過隱式拷貝的方式 95: //拷貝到clbuf1,buf2的內容通過顯示拷貝的方式拷貝到clbuf2 96: cl_mem clbuf1 = clCreateBuffer(context, 97: CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 98: NWITEMS*sizeof(cl_float),buf1, 99: NULL ); 100:? 101: cl_mem clbuf2 = clCreateBuffer(context, 102: CL_MEM_READ_ONLY , 103: NWITEMS*sizeof(cl_float),NULL, 104: NULL ); 105:? 106: status = clEnqueueWriteBuffer(queue, clbuf2, 1, 107: 0, NWITEMS*sizeof(cl_float), buf2, 0, 0, 0); 108:? 109: cl_mem buffer = clCreateBuffer( context, 110: CL_MEM_WRITE_ONLY, 111: NWITEMS * sizeof(cl_float), 112: NULL, NULL ); 113:? 114: const char * filename = "add.cl"; 115: std::string sourceStr; 116: status = convertToString(filename, sourceStr); 117: const char * source = sourceStr.c_str(); 118: size_t sourceSize[] = { strlen(source) }; 119:? 120: //創建程序對象 121: cl_program program = clCreateProgramWithSource( 122: context, 123: 1, 124: &source, 125: sourceSize, 126: NULL); 127: //編譯程序對象 128: status = clBuildProgram( program, 1, &device, NULL, NULL, NULL ); 129: if(status != 0) 130: { 131: printf("clBuild failed:%d\n", status); 132: char tbuf[0x10000]; 133: clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL); 134: printf("\n%s\n", tbuf); 135: return -1; 136: } 137:? 138: //創建Kernel對象 139: cl_kernel kernel = clCreateKernel( program, "vecadd", NULL ); 140: //設置Kernel參數 141: cl_int clnum = NWITEMS; 142: clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1); 143: clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2); 144: clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &buffer); 145:? 146: //執行kernel 147: cl_event ev; 148: size_t global_work_size = NWITEMS; 149: clEnqueueNDRangeKernel( queue, 150: kernel, 151: 1, 152: NULL, 153: &global_work_size, 154: NULL, 0, NULL, &ev); 155: clFinish( queue ); 156:? 157: //數據拷回host內存 158: cl_float *ptr; 159: ptr = (cl_float *) clEnqueueMapBuffer( queue, 160: buffer, 161: CL_TRUE, 162: CL_MAP_READ, 163: 0, 164: NWITEMS * sizeof(cl_float), 165: 0, NULL, NULL, NULL ); 166: //結果驗證，和cpu計算的結果比較 167: if(!memcmp(buf, ptr, NWITEMS)) 168: printf("Verify passed\n"); 169: else printf("verify failed"); 170:? 171: if(buf) 172: free(buf); 173: if(buf1) 174: free(buf1); 175: if(buf2) 176: free(buf2); 177:? 178: //刪除OpenCL資源對象 179: clReleaseMemObject(clbuf1); 180: clReleaseMemObject(clbuf2); 181: clReleaseMemObject(buffer); 182: clReleaseProgram(program); 183: clReleaseCommandQueue(queue); 184: clReleaseContext(context); 185: return 0; 186: } 187:?

也可以在http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode1.zip&can=2&q=#makechanges上下載完整版本。

GPU架構

內容包括：

1.OpenCLspec和多核硬件的對應關系

• AMD GPU架構
• Nvdia GPU架構
• Cell Broadband Engine

2.一些關于OpenCL的特殊主題

• OpenCL編譯系統
• Installable client driver

?

首先我們可能有疑問，既然OpenCL具有平臺無關性，我們為什么還要去研究不同廠商的特殊硬件設備呢?

• 了解程序中的循環和數據怎樣映射到OpenCL Kernel中，便于我們提高代碼質量，獲得更高的性能。
• 了解AMD和Nvdia顯卡的區別。
• 了解各種硬件的區別，可以幫助我們使用基于這些硬件的一些特殊的OpenCL擴展，這些擴展在后面課程中會講到。

3、傳統的CPU架構

• ??? 對單個線程來說，CPU優化能獲得最小時延，而且CPU也適合處理控制流密集的工作，比如if、else或者跳轉指令比較多的任務。
• 控制邏輯單元在芯片中占用的面積要比ALU單元多。
• 多層次的cache設計被用來隱藏時延（可以很好的利用空間和時間局部性原理）
• 有限的寄存器數量使得同時active的線程不能太多。
• 控制邏輯單元記錄程序的執行、提供指令集并行（ILP）以及最小化CPU管線的空置周期（stalls，在該時鐘周期，ALU沒做什么事）。

4、現代的GPGPU架構

?

• 對于現代的GPU，通常的它的控制邏輯單元比較簡單（和cpu相比），cache也比較小
• 線程切換開銷比較小，都是輕量級的線程。
• GPU的每個“核”有大量的ALU以及很小的用戶可管理的cache。[這兒的核應該是指整個GPU]。
• 內存總線都是基于帶寬優化的。150GB/s的帶寬可以使得大量ALU同時進行內存操作。

5、AMD GPU硬件架構

現在我們簡單看下AMD 5870顯卡(cypress)的架構

• 20個simd引擎，每個simd引擎包含16個simd。
• 每個simd包含16個stream core
• 每個stream core都是5路的乘法-加法運算單元（VLIW processing）。
• 單精度運算可以達到 Teraflops。
• 雙精度運算可以達到544Gb/s

上圖為一個simd引擎的示意圖，每個simd引擎由一系列的stream core組成。

• 每個stream core是一個5路的VLIW處理器，在一個VLIW指令中，可以最多發射5個標量操作。標量操作在每個pe上執行。
• CU（8xx系列cu對應硬件的simd）內的stream core執行相同的VLIW指令。
• 在CU（或者說simd）內同時執行的work item放在一起稱作一個wave，它是cu中同時執行的線程數目。在5870中wave大小是64，也就是說一個cu內，最多有64個work item在同時執行。

注：5路的運算對應(x,y,z,w),以及T（超越函數），在cayman中，已經取消了T，改成四路了。

?

我們現在看下AMD GPU硬件在OpenCL中的對應關系：

• 一個workitme對應一個pe，pe就是單個的VLIW core
• 一個cu對應多個pe，cu就是simd引擎。

上圖是AMD GPU的內存架構（原課件中的圖有點小錯誤，把Global memory寫成了LDS)

• 對每個cu來說，它使用的內存包括onchip的LDS以及相關寄存器。在5870中，每個LDS是32K，共32個bank，每個bank 1k，讀寫單位4 byte。
• 對沒給cu來說，有8K的L1 cache。（for 5870）
• 各個cu之間共享的L2 cache，在5870中是512K。
• fast Path只能執行32位或32位倍數的內存操作。
• complete path能夠執行原子操作以及小于32位的內存操作。

AMD GPU的內存架構和OpenCL內存模型之間的對應關系：

• LDS對應local memeory，主要用來在一個work group內的work times之間共享數據。steam core訪問LDS的速度要比Global memory快一個數量級。
• private memory對應每個pe的寄存器。
• constant memory主要是利用了L1 cache

注意：對AMD CPU，constant memory的訪問包括三種方式：Direct-Addressing Patterns，這種模式要求不包括行列式，它的值都是在kernel函數初始化的時候就決定了，比如傳入一個固定的參數。Same Index Patterns，所有的work item都訪問相同的索引地址。Globally scoped constant arrays，行列式會被初始化，如果小于16K，會使用L1 cache，從而加快訪問速度。

當所有的work item訪問不同的索引地址時候，不能被cache，這時要在global memory中讀取。

?

?

6、Nvdia GPU Femi架構

?

GTX480-Compute 2.0 capability：

• 有15個core或者說SM（Streaming Multiprocessors ）。
• 每個SM,一般有32 cuda處理器。
• 共480個cuda處理器。
• 帶ECC的global memory
• 每個SM內的線程按32個單位調度執行，稱作warp。每個SM內有2個warp發射單元。
• 一個cuda核由一個ALU和一個FPU組成，FPU是浮點處理單元。

SIMT和SIMD

SIMT是指單指令、多線程。

• 硬件決定了多個ALU之間要共享指令。
• 通過預測來處理多個線程間的Diverage(是指同一個warp中的指令執行路徑產生不同）。
• NV把一個warp中執行的指令當作一個SIMT。SIMT指令指定了一個線程的執行以及分支行為。

SIMD指令可以得到向量的寬度，這點和X86 SSE向量指令比較類似。

SIMD的執行和管線相關

• 所有的ALU執行相同的指令。
• 根據指令可以管線分為不同的階段。當第一條指令完成的時候（4個周期），下條指令開始執行。

Nvida GPU內存機制：

• 每個SM都有L1 cache，通過配置，它可以支持shared memory，也可以支持global memory。
• 48 KB Shared / 16 KB of L1 cache，16 KB Shared / 48 KB of L1 cache
• work item之間數據共享通過shared memory
• 每個SM有32K的register bank
• L2(768K)支持所有的操作，比如load,store等等
• Unified path to global for loads and stores?

和AMD GPU類似，Nv的GPU 內存模型和OpenCL內存模型的對應關系是：

• shared memory對應local memory
• 寄存器對應private memory

7、Cell Broadband Engine

?

由索尼，東芝，IBM等聯合開發，可用于嵌入式平臺，也可用于高性能計算（SP3次世代游戲主機就用了cell處理器）。

• Bladecenter servers提供OpenCL driver支持
• 如圖所示，cell處理器由一個Power Processing Element (PPE) 和多個Synergistic Processing Elements (SPE)組成。
• Uses the IBM XL C for OpenCL compiler 11
• Cell Power/VMX CPU 的設備類型是CL_DEVICE_TYPE_CPU，Cell SPU 的設備類型是CL_DEVICE_TYPE_ACCELERATOR。
• OpenCL Accelerator設備和CPU共享內存總線。
• 提供一些擴展，比如Device Fission、Migrate Objects來指定一個OpenCL對象駐留在什么位置。
• 不支持OpenCL image對象，原子操作，sampler對象以及字節內存地址。

8、OpenCL編譯系統

• LLVM-底層的虛擬機
• Kernel首先在front-end被編譯成LLVM IR
• LLVM是一個開源的編譯器，具有平臺獨立性，可以支持不同廠商的back_end編譯,網址：http://llvm.org

9、Installable Client Driver

• ICD支持不同廠商的OpenCL實施在系統中共存。
• 代碼緊被鏈接接到libOpenCL.so
• 應用程序可在運行時選擇不同的OpenCL實施（就是選擇不同platform）
• 現在的GPU驅動還不支持跨廠商的多個GPU設備同時工作。
• 通過clGetPlatformIDs() 和clGetPlatformInfo() 來檢測不同廠商的OpenCL平臺。

本節主要講述GPU的memory架構。優化基于GPU device的kernel程序時，我們需要了解很多GPU的memory知識，比如內存合并，bank conflit（沖突）等等，這樣才能針對具體算法做一些優化工作。

1、GPU總線尋址介紹

?

?? 假定X是一個指向整數（32位整數）數組的指針，數組的首地址為0x00001232。一個線程要訪問元素X[0],

?? int tmp = X[0];

???

??? 假定memory總線寬度為256位(HD5870就是如此,即為32字節），因為基于字節地址的總線要訪問memeory，必須和總線寬度對齊，也就是說按必須32字節對齊來訪問memory，比如訪問0x00000000,0x00000020,0x00000040,…等，所以我們要得到地址0x00001232中的數據，比如訪問地址0x00001220,這時，它會同時得到0x00001220到 0x0000123F 的所有數據。因為我們只是取的一個32位整數，所以有用的數據是4個字節，其它28的字節的數據都被浪費了，白白消耗了帶寬。

???

?

2、合并內存訪問

??? 為了利用總線帶寬，GPU通常把多個線程的內存訪問盡量合并到較少的內存請求命令中去。

??? 假定下面的OpenCL kernel代碼：int tmp = X[get_global_id(0)];

數組X的首地址和前面例子一樣，也是0x00001232，則前16個線程將訪問地址：0x00001232 到 0x00001272。假設每個memory訪問請求都單獨發送的話，則有16個request，有用的數據只有64字節，浪費掉了448字節（16*28）。

??? 假定多個線程訪問32個字節以內的地址，它們的訪問可以通過一個memory request完成，這樣可以大大提高帶寬利用率，在專業術語描述中這樣的合并訪問稱作coalescing。

?? 例如上面16個線程訪問地址0x00001232 到 0x00001272，我們只需要3次memory requst。

?? 在HD5870顯卡中，一個wave中16個連續線程的內存訪問會被合并，稱作quarter-wavefront，是重要的硬件調度單位。

?? 下面的圖是HD5870中，使用memory訪問合并以及沒有使用合并的bandwidth比較：

?? 下圖是GTX285中的比較：

3、Global memory的bank以及channel訪問沖突

?? 我們知道內存由bank，channel組成，bank是實際存儲數據的單元，一個mc可以連接多個channel，形成單mc，多channel的連接方式。在物理上，不同bank的數據可以同時訪問，相同的bank的數據則必須串行訪問，channel也是同樣的道理。但由于合并訪問的緣故，對于global memory來說，bank conflit影響要小很多，除非是非合并問，不同線程訪問同一個bank。理想情況下，我們應該做到不同的workgroup訪問的不同的bank，同一個group內，最好用合并操作。

?? 下面我簡單的畫一個圖，不知道是否準確，僅供參考：

?

???

???? 在HD5870中，memory地址的低8位表示一個bank中的數據，接下來的3位表示channel（共8個channel），bank位的多少依賴于顯存中bank的多少。

4、local memory的bank conflit

?? bank訪問沖突對local memory操作有更大的影響（相比于global memory），連續的local memory訪問地址，應該映射到不同的bank上，

???? 在AMD顯卡中，一個產生bank訪問沖突wave將會等待所有的local memory訪問完成，硬件不能通過切換到另一個wave來隱藏local memory訪問時延。所以對local memory訪問的優化就很重要。HD5870顯卡中，每個cu（simd）有32bank，每個bank 1k，按4字節對齊訪問。如果沒有bank conflit，每個bank能夠沒有延時的返回一個數據，下面的圖就是這種情況。

?? 如果多個memory訪問對應到一個bank上，則conflits的數量決定時延的大小。下面的訪問方式將會有3倍的時延。

? 但是，如果所有訪問都映射到一個bank上，則系統會廣播數據訪問，不會產生額外時延。

GPU線程及調度

???? 本節主要講述OpenCL中的Workgroup如何在硬件設備中被調度執行。同時也會講一下同一個workgroup中的workitem，如果它們執行的指令發生diverage（就是執行指令不一致）對性能的影響。學習OpenCL并行編程，不僅僅是對OpenCL Spec本身了解，更重要的是了解OpenCL硬件設備的特性，現階段來說，主要是了解GPU的的架構特性，這樣才能針對硬件特性優化算法。

???? 現在OpenCL的Spec是1.1，隨著硬件的發展，相信OpenCL會支持更多的并行計算特性?；贠penCL的并行計算才剛剛起步，…

1、workgroup到硬件線程

???? 在OpenCL中，Kernel函數被workgroup中的workitem（線程，我可能混用這兩個概念）執行。在硬件層次，workgroup被映射到硬件的cu（compute unit）單元來執行具體計算，而cu一般由更多的SIMT（單指令，線程）pe（processing elements）組成。這些pe執行具體的workitem計算，它們執行同樣的指令，但操作的數據不一樣，用simd的方式完成最終的計算。

??? 由于硬件的限制，比如cu中pe數量的限制，實際上workgroup中線程并不是同時執行的，而是有一個調度單位，同一個workgroup中的線程，按照調度單位分組，然后一組一組調度硬件上去執行。這個調度單位在nv的硬件上稱作warp,在AMD的硬件上稱作wavefront，或者簡稱為wave。

? 上圖顯示了workgroup中，線程被劃分為不同wave的分組情況。wave中的線程同步執行相同的指令，但每個線程都有自己的register狀態，可以執行不同的控制分支。比如一個控制語句

if(A)

{

… //分支A

}

else

{

? … //分支B

}

??? 假設wave中的64個線程中，奇數線程執行分支A，偶數線程執行分支B，由于wave中的線程必須執行相同的指令，所以這條控制語句被拆分為兩次執行[編譯階段進行了分支預測]，第一次分支A的奇數線程執行，偶數線程進行空操作，第二次偶數線程執行，奇數線程空操作。硬件系統有一個64位mask寄存器，第一次是它為01…0101,第二次會進行反轉操作10…1010，根據mask寄存器的置位情況，來選擇執行不同的線程?？梢妼τ诜种Ф嗟膋ernel函數，如果不同線程的執行發生diverage的情況太多，會影響程序的性能。

2、AMD wave調度

? ? AMD GPU的線程調度單位是wave，每個wave的大小是64。指令發射單元發射5路的VLIW指令，每個stream core(SC)執行一條VLIW指令，16個stream core在一個時鐘周期執行16條VLIW指令。每個時鐘周期，1/4wave被完成，整個wave完成需要四個連續的時鐘周期。

??? 另外還有以下幾點值得我們了解：

• 發生RAW hazard情況下，整個wave必須stall 4個時鐘周期，這時，如果其它的wave可以利用，ALU會執行其它的wave以便隱藏時延，8個時鐘周期后，如果先前等待wave已經準備好了，ALU會繼續執行這個wave。
• 兩個wave能夠完全隱藏RAW時延。第一個wave執行時候，第二個wave在調度等待數據，第一個wave執行完時，第二個wave可以立即開始執行。

3、nv warp調度

???? work group以32個線程為單位，分成不同warp，這些warp被SM調度執行。每次warp中一半的線程被發射執行，而且這些線程能夠交錯執行?？梢杂玫膚arp數量依賴于每個block的資源情況。除了大小不一樣外，wave和warp在硬件特性上很相似。

4、Occupancy開銷

??? 在每個cu中，同時激活的wave數量是受限制的，這和每個線程使用register和local memory大小有關，因為對于每個cu，register和local memory總量是一定的。

??? 我們用術語Occupancy來衡量一個cu中active wave的數量。如果同時激活的wave越多，能更好的隱藏時延，在后面性能優化的章節中，我們還會更具體討論Occupancy。

5、控制流和分支預測(prediction)

?? 前面我說了if else的分支執行情況，當一個wave中不同線程出現diverage的時候，會通過mask來控制線程的執行路徑。這種預測（prediction）的方式基于下面的考慮：

• 分支的代碼都比較短
• 這種prediction的方式比條件指令更高效。
• 在編譯階段，編譯器能夠用predition替換switch或者if else。

? prediction 可以定義為：根據判斷條件，條件碼被設置為true或者false。

__kernel void test() {int tid= get_local_id(0) ;if( tid %2 == 0) Do_Some_Work() ;else Do_Other_Work() ; }

例如上面的代碼就是可預測的，

Predicate = True for threads 0,2,4….

Predicate = False for threads 1,3,5….

下面在看一個控制流diverage的例子

• 在case1中，所有奇數線程執行DoSomeWork2(),所有偶數線程執行DoSomeWorks,但是在每個wave中，if和else代碼指令都要被發射。
• 在case2中，第一個wave執行if，其它的wave執行else，這種情況下，每個wave中，if和else代碼只被發射一個。

?? 在prediction下，指令執行時間是if，else兩個代碼快執行時間之和。

6、Warp voting

?? warp voting是一個warp內的線程之間隱式同步的機制。

??? 比如一個warp內線程同時寫Local meory某個地址，在線程并發執行時候，warp voting機制可以保證它們的前后順序正確。更詳細的warp voting大家可以參考cuda的資料。

??

??? 在OpenCL編程中，由于各種硬件設備不同，導致我們必須針對不同的硬件進行優化，這也是OpenCL編程的一個挑戰，比如warp和wave數量的不同，使得我們在設計workgroup大小時候，必須針對自己的平臺進行優化，如果選擇32，對于AMD GPU，可能一個wave中32線程是空操作，而如果選擇64，對nv GPU來說，可能會出現資源競爭的情況加劇，比如register以及local meomory的分配等等。這兒還不說混合CPU device的情況,OpenCL并行編程的道路還很漫長，期待新的OpenCL架構的出現。

性能優化

1、線程映射

?? 所謂線程映射是指某個線程訪問哪一部分數據，其實就是線程id和訪問數據之間的對應關系。

合適的線程映射可以充分利用硬件特性，從而提高程序的性能，反之，則會降低performance。

?? 請參考Static Memory Access Pattern Analysis on a Massively Parallel GPU這篇paper，文中講述線程如何在算法中充分利用線程映射。這是我在google中搜索到的下載地址：http://www.ece.neu.edu/~bjang/patternAnalysis.pdf

?? 使用不同的線程映射，同一個線程可能訪問不同位置的數據。下面是幾個線程映射的例子：

????? 我們考慮一個簡單的串行矩陣乘法：這個算法比較適合輸出數據降維操作，通過創建N*M個線程，我們移去兩層外循環，這樣每個線程執行P個加法乘法操作?，F在需要我們考慮的問題是，線程索引空間究竟應該是M*N還是N*M？

??? 當我們使用M*N線程索引空間時候，Kernel如下圖所示：

?? 而使用N*M線程索引空間時候，Kernel如下圖所示：

??? 使用兩種映射關系，程序執行結果是一樣的。下面是在nv的卡GeForce 285 and 8800 GPUs上的執行結果?？梢钥吹接成?（及N*M線程索引空間），程序的performance更高。

??? performance差異主要是因為在兩種映射方式下，對global memory訪問的方式有所不同。在行主序的buffer中，數據都是按行逐個存儲，為了保證合并訪問，我們應該把一個wave中連續的線程映射到矩陣的列（第二維），這樣在A*B=C的情況下，會把矩陣B和C的內存讀寫實現合并訪問，而兩種映射方式對A沒有影響（A又i3決定順序）。

?? 完整的源代碼請從：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amduniCourseCode4.zip&can=2&q=#makechanges下載，程序中我實現了兩種方式的比較。結果確實第二種方式要快一些。

?? 下面我們再看一個矩陣轉置的例子，在例子中，通過改變映射方式，提高了global memory訪問的效率。

?? 矩陣轉置的公式是:Out(x,y) = In(y,x)

?? 從上圖可以看出，無論才去那種映射方式，總有一個buffer是非合并訪問方式（注：在矩陣轉置時，必須要把輸入矩陣的某個元素拷貝到臨時位置，比如寄存器，然后才能拷貝到輸出矩陣）。我們可以改變線程映射方式，用local memory作為中間元素，從而實現輸入，輸出矩陣都是global memory合并訪問。

? 下面是AMD 5870顯卡上，兩種線程映射方式實現的矩陣轉置性能比較：

??? 完整代碼：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amduniCourseCode5.zip&can=2&q=#makechanges

2、Occupancy

??? 前面的教程中，我們提到過Occupancy的概念，它主要用來描述CU中資源的利用率。

??? OpenCL中workgroup被映射到硬件的CU中執行，在一個workgroup中的所有線程執行完之后，這個workgroup才算執行結束。對一個特定的cu來說，它的資源（比如寄存器數量，local memory大小，最大線程數量等）是固定的，這些資源都會限制cu中同時處于調度狀態的workgroup數量。如果cu中的資源數量足夠的的話，映射到同一個cu的多個workgroup能同時處于調度狀態，其中一個workgroup的wave處于執行狀態，當處于執行狀態的workgroup所有wave因為等待資源而切換到等待狀態的話，不同workgroup能夠從就緒狀態切換到ALU執行，這樣隱藏memory訪問時延。這有點類似操作系統中進程之間的調度狀態。我簡單畫個圖，以供參考：

• 對于一個比較長的kernel，寄存器是主要的資源瓶頸。假設kernel需要的最大寄存器數目為35，則workgroup中的所有線程都會使用35個寄存器，而一個CU（假設為5870）的最大寄存器數目為16384，則cu中最多可有16384/35=468線程，此時，一個workgroup中的線程數目（workitem）不可能超過468，
• 考慮另一個問題，一個cu共16384個寄存器，而workgroup固定為256個線程，則使用的寄存器數量可達到64個。

??? 每個CU的local memory也是有限的，對于AMD HD 5XXX顯卡，local memory是32K，NV的顯卡local memory是32-48K（具體看型號）。和使用寄存器的情況相似，如果kernel使用過多的local memory，則workgroup中的線程數目也會有限制。

?? GPU硬件還有一個CU內的最大線程數目限制：AMD顯卡256，nv顯卡512。

?? NV的顯卡對于每個CU內的激活線程有數量限制，每個cu 8個或16個warp，768或者1024個線程。

?? AMD顯卡對每個CU內的wave數量有限制，對于5870，最多496個wave。

?? 這些限制都是因為有限的資源競爭引起的，在nv cuda中，可以通過可視化的方式查看資源的限制情況。

3、向量化

?? 向量化允許一個線程同時執行多個操作。我們可以在kernel代碼中，使用向量數據類型，比如float4來獲得加速。向量化在AMD的GPU上效果更為明顯，這是因為AMD的顯卡的stream core是(x,y,z,w)這樣的向量運算單元。

?? 下圖是在簡單的向量賦值運算中，使用float和float4的性能比較。

??? kernel代碼為：

??? 完整代碼請從：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amduniCourseCode6.zip&can=2&q=#makechanges下載

分類: OpenCL 綠色通道：好文要頂關注我收藏該文與我聯系 邁克老狼2012
關注 - 7
粉絲 - 127 +加關注 1 0 (請您對文章做出評價) ?上一篇：AMD OpenCL大學課程(10)
?下一篇：AMD OpenCL大學課程(12) 性能優化案例NBody

posted on 2012-01-31 19:26 邁克老狼2012 閱讀(881) 評論(3)編輯 收藏

評論

#1樓??

“我們應該把一個wave中連續的線程映射到矩陣的列（第二維）”

這句話我覺得重點是“連續”而不是“列”。 支持(0)反對(0) 2013-07-26 22:20 | 撥浪鼓兒

#2樓??

在行主序的buffer中，數據都是按行逐個存儲，為了保證合并訪問，我們應該把一個wave中連續的線程映射到矩陣的列（第二維）。
這句話每次讀都不順口，實際是這樣的，比如B【4】【4】，存的時候是B[0][0],B[0][1],B[0][2]...這樣存儲的，在每一行中是按照第二維遞增的，也就是按列遞增，存完第一行然后存的第二行。因此應該把一個wave中連續的線程映射到矩陣的列。Opencl線程是以（0，0），（1，0）（2，0）...這樣變換的，也就是先是get_global_id(0)變換，然后是get_global_id(1)變換，因此將get_global_id(0)對應到列。 支持(0)反對(0) 2013-12-02 10:52 | 撥浪鼓兒

#3樓??

“我們可以改變線程映射方式，用local memory作為中間元素，從而實現輸入，輸出矩陣都是global memory合并訪問。” 這里雖然會實現合并訪問，但是有可能會出現local memory上的bank conflict. 本節主要介紹NBody算法的OpenCL性能優化。

1、NBody

??? NBody系統主要用來通過粒子之間的物理作用力來模擬星系系統。每個粒子表示一個星星，多個粒子之間的相互作用，就呈現出星系的效果。

?

?? 上圖為一個粒子模擬星系的圖片：Source: THE GALAXY-CLUSTER-SUPERCLUSTER CONNECTION，http://www.casca.ca/ecass/issues/1997-DS/West/west-bil.html

?? 由于每個粒子之間都有相互作用的引力，所以這個算法的復雜度是N²的。下面我們主要探討如何優化算法以及在OpenCL基礎上優化算法。

2、NBody算法

?? 假設兩個粒子之間通過萬有引力相互作用，則任意兩個粒子之間的相互作用力F公式如下：

?? 最笨的方法就是計算每個粒子和其它粒子的作用力之和，這個方法通常稱作N-Pair的NBody模擬。

?? 粒子之間的萬有引力和它們之間的距離成反比，對于一個粒子而言（假設粒子質量都一樣），遠距離粒子的作用力有時候很小，甚至可以忽略。Barnes Hut 把3D空間按八叉樹進行分割，只有在相鄰cell的粒子才直接計算它們之間的引力，遠距離cell中的粒子當作一個整體來計算引力。

3、OpenCL優化Nbody

???? 在本節中，我們不考慮算法本身的優化，只是通過OpenCL機制來優化N-Pair的NBody模擬。

???? 最簡單的實施方法就是每個例子的作用力相加，代碼如下：

for(i=0; i<n; i++) { ax = ay = az = 0;// Loop over all particles "j” for (j=0; j<n; j++) {//Calculate Displacement dx=x[j]-x[i]; dy=y[j]-y[i]; dz=z[j]-z[i];// small eps is delta added for dx,dy,dz = 0 invr= 1.0/sqrt(dx*dx+dy*dy+dz*dz +eps); invr3 = invr*invr*invr; f=m[ j ]*invr3;// Accumulate acceleration ax += f*dx; ay += f*dy; az += f*dx; }// Use ax, ay, az to update particle positions }

我們對每個粒子計算作用在它上面的合力，然后求在合力作用下，delta時間內粒子的新位置，并把這個新位置當作下次計算的輸入參數。

沒有優化的OpenCL kernel代碼如下：

__kernel void nbody_sim_notile( __global float4* pos , __global float4* vel,int numBodies,float deltaTime,float epsSqr, __local float4* localPos, __global float4* newPosition, __global float4* newVelocity){unsigned int tid = get_local_id(0);unsigned int gid = get_global_id(0);unsigned int localSize = get_local_size(0);// position of this work-item float4 myPos = pos[gid]; float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);// load one tile into local memoryint idx = tid * localSize + tid; localPos[tid] = pos[idx];// calculate acceleration effect due to each body// a[i->j] = m[j] * r[i->j] / (r^2 + epsSqr)^(3/2)for(int j = 0; j < numBodies; ++j) {// Calculate acceleartion caused by particle j on particle i localPos[tid] = pos[j]; float4 r = localPos[j] - myPos;float distSqr = r.x * r.x + r.y * r.y + r.z * r.z;float invDist = 1.0f / sqrt(distSqr + epsSqr);float invDistCube = invDist * invDist * invDist;float s = localPos[j].w * invDistCube;// accumulate effect of all particles acc += s * r; }float4 oldVel = vel[gid];// updated position and velocity float4 newPos = myPos + oldVel * deltaTime + acc * 0.5f * deltaTime * deltaTime; newPos.w = myPos.w;float4 newVel = oldVel + acc * deltaTime;// write to global memory newPosition[gid] = newPos; newVelocity[gid] = newVel; }

在這種實現中，每次都要從global memory中讀取其它粒子的位置，速度，內存訪問= N reads*N threads=N²

我們可以通過local memory進行優化，一個粒子數據讀進來以后，可以被p*p個線程共用，p*p即為workgroup的大小，對于每個粒子，我們通過迭代p*p的tile，累積得到最終結果。

優化后的kernel代碼如下：

__kernel void nbody_sim(__global float4* pos ,__global float4* vel,int numBodies,float deltaTime,float epsSqr,__local float4* localPos, __global float4* newPosition, __global float4* newVelocity){unsigned int tid = get_local_id(0);unsigned int gid = get_global_id(0);unsigned int localSize = get_local_size(0);// Number of tiles we need to iterateunsigned int numTiles = numBodies / localSize;// position of this work-itemfloat4 myPos = pos[gid];float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);for(int i = 0; i < numTiles; ++i){// load one tile into local memoryint idx = i * localSize + tid;localPos[tid] = pos[idx];// Synchronize to make sure data is available for processingbarrier(CLK_LOCAL_MEM_FENCE);// calculate acceleration effect due to each body// a[i->j] = m[j] * r[i->j] / (r^2 + epsSqr)^(3/2)for(int j = 0; j < localSize; ++j){// Calculate acceleartion caused by particle j on particle ifloat4 r = localPos[j] - myPos;float distSqr = r.x * r.x + r.y * r.y + r.z * r.z;float invDist = 1.0f / sqrt(distSqr + epsSqr);float invDistCube = invDist * invDist * invDist;float s = localPos[j].w * invDistCube;// accumulate effect of all particlesacc += s * r;}// Synchronize so that next tile can be loadedbarrier(CLK_LOCAL_MEM_FENCE);}float4 oldVel = vel[gid];// updated position and velocityfloat4 newPos = myPos + oldVel * deltaTime + acc * 0.5f * deltaTime * deltaTime;newPos.w = myPos.w;float4 newVel = oldVel + acc * deltaTime;// write to global memorynewPosition[gid] = newPos;newVelocity[gid] = newVel; }

下面是在AMD, NV兩個平臺上性能測試結果：

AMD GPU = 5870 Stream SDK 2.2

Nvidia GPU = GTX 480 with CUDA 3.1

另外，在程序中，也嘗試了循環展開，通過展開內循環，從而減少GPU執行分支指令，我的測試中，使用展開四次，得到的FPS比沒展開前快了30%。（AMD 5670顯卡）。具體實現可以看kernel代碼中的__kernel void nbody_sim_unroll函數。在AMD平臺上，使用向量化也可以提高10%左右的性能。

最后提供2篇NBody優化的文章：

—Nvidia GPU Gems

—http://http.developer.nvidia.com/GPUGems3/gpugems3_ch31.html

—Brown Deer Technology

—http://www.browndeertechnology.com/docs/BDT_OpenCL_Tutorial_NBody.html

第二個可能地址需要fan墻。

完整的代碼可從：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amduniCourseCode7.zip&can=2&q=#makechanges 下載。

1、OpenCL擴展

???? OpenCL擴展是指device支持某種特性，但這中特性并不是OpenCL標準的一部分。通過擴展，廠商可以給device增加一些新的功能，而不用考慮兼容性問題?，F在各個廠商在OpenCL的實現中或多或少的使用了自己的擴展。

???? 擴展的類型分為三種：

• Khronos OpenCL工作組批準的擴展，這種要經過一致性測試，可能會被增加到新版本的OpenCL規范中。這種擴展都以cl_khr作為擴展名。
• 外部擴展, 以cl_ext為擴展名。這種擴展是由2個或2個以上的廠商發起，并不需要進行一致性測試。比如cl_ext_device_fission擴展。
• 某個廠商自己的擴展，比如AMD的擴展printf

2、使用擴展

????? OpenCL中，要使用擴展，我們必須打開擴展，在默認狀態下，所有的擴展都是禁止的。

?????? #pragma OPENCL EXTENSION extension_name : enable

?????? 對于OpenCL，一個函數只有在運行時，才知道其是否可用，所以要確定某個擴展是否可用，是程序員的責任，我們必須在使用前查詢它的狀態。下面是查詢擴展是否可用的代碼：

3、一些Khronos批準的擴展

?? 原子操作，它可以保證函數只在一個device上實施原子操作，比如：

—cl_khr_{global | local}_int32_base_atomics

—cl_khr_{global | local}_int32_extended_atomics

—cl_khr_int64_base_atomics

—cl_khr_int64_extended_atomics

注意：原子操作能夠保證操作結果正確，但不保證操作的順序。

?????? 雙精度和half精度擴展cl_khr_fp64，在一些物理模擬或者科學計算中，需要雙精度支持。AMD的64位擴展用cl_amd_fp64，對于cl_khr_fp64是部分支持，NV支持cl_khr_fp64擴展。但half精度擴展cl_khr_fp16，這兩家廠商現在都還不支持。

?????? 在OpenCL中，Byte addressable store 也是一個擴展，對于sub 32的寫，比如char，需要該擴展的支持。例如AMD 直方圖的例子中，每個bin用一個byte來存儲。

?????? 3D Image Write Extensions，在OpenCL標準中，支持2D圖像的讀寫，3D圖形的寫就需要通過擴展來操作。

?????? The extension cl_KHR_gl_sharing 允許應用程序使用OpenGL buffer，紋理等。

4、AMD擴展

???? cl_ext_device_fission擴展，通過該擴展把一個設備分成多個子設備，每一個設備都有自己的隊列，主要是多核cpu以及Cell Broadband Engine使用，該擴展由AMD，Apple，Intel以及IBM四家聯合提出。

???? fission設備可能的用途包括：

• 保留一部分設備處理高優先級、低時延的任務。
• Control for the assignment of work to individual compute units
• Subdivide compute devices along some shared hardware feature like a cache

???? 對于每個子設備，都有自己的queue，比如下面的圖中，我們把不同任務發送到兩個子設備。值得注意的是：要把設備拆分為子設備，首先我們要了解該設備的架構，然后根據任務及device架構進行拆分。

?????? GPU printf 擴展，主要用來debug kernel代碼。cl_amd_media_ops擴展，主要用于一些多媒體操作。The AMD device query extension 主要用于查詢和事件處理。

??????

? 5、NV擴展

• Compiler Options
• Interoperability Extensions
• Device Query Extension

6、Cell Broadband Engine Extensions

????? cell處理器用的不多，就不詳細說了，使用的人可以查詢其相關手冊。

，??? 在本節，我們主要介紹OpenCL中buffer的使用，同時提供了2個完整的例子，一個是圖像的旋轉，一個是矩陣乘法（非常簡單，沒有分塊優化）。

?

1、創建OpenCL設備緩沖(buffer)

?? OpenCL設備使用的數據都存放在設備的buffer中[其實就是device memory中]。我們用下面的代碼創建buffer對象：

cl_mem bufferobj = clCreateBuffer ( cl_context context, //Context name cl_mem_flags flags, //Memory flags size_t size, //Memory size allocated in buffervoid *host_ptr, //Host data cl_int *errcode) //Returned error code

??? 如果host_ptr指向一個有效的host指針，則創建一個buffer對象的同時會實現隱式的數據拷貝（會在kernel函數進入隊列時候，把host_prt中的數據從host memory拷貝到設備內存對象bufferobj中)。

??? 我們可以通過flags參數指定buffer對象的屬性。

?

?? 函數clEnqueueWriteBuffer()用來實現顯示的數據拷貝，即把host memory中的數據拷貝到device meomory中。

cl_int clEnqueueWriteBuffer ( cl_command_queue queue, //Command queue to device cl_mem buffer, //OpenCL Buffer Object cl_bool blocking_read, //Blocking/Non-Blocking Flag size_t offset, //Offset into buffer to write to size_t cb, //Size of datavoid *ptr, //Host pointer cl_uint num_in_wait_list, //Number of events in wait listconst cl_event * event_wait_list, //Array of events to wait for cl_event *event) //Event handler for this function

2、圖像旋轉的例子

?? 下面是一個完整的OpenCL例子，實現圖像的旋轉。在這個例子中，我把美麗的lenna旋轉了90度。

下面是原始圖像和旋轉后的圖像（黑白）

在這個例子中，我使用FreeImage庫，可以從FreeImage網站或者我的code工程中下載。

http://code.google.com/p/imagefilter-opencl/downloads/detail?name=Dist.rar&can=2&q=#makechanges

?? 圖像旋轉是指把定義的圖像繞某一點以逆時針或順時針方向旋轉一定的角度，通常是指繞圖像的中心以逆時針方向旋轉。

假設圖像的左上角為（left, top),右下角為（right, bottom)，則圖像上任意點（x0, y0）繞其中心（xcenter, ycenter)逆時針旋轉angle角度后，新的坐標位置（x′, y′）的計算公式為：

xcenter = (right － left ＋ 1) ／ 2 ＋ left;
ycenter = (bottom － top ＋ 1) ／ 2 ＋ top;
x′ = (x0 － xcenter) cosθ － (y0 － ycenter) sinθ ＋ xcenter;
y′ = (x0 － xcenter) sinθ ＋ (y0 － ycenter) cosθ ＋ ycenter;

下面給出kernel的代碼：

1: __kernel void image_rotate( __global uchar * src_data, __global uchar * dest_data, //Data in global memory 2: int W, int H, //Image Dimensions 3: float sinTheta, float cosTheta ) //Rotation Parameters 4: { 5: //Thread gets its index within index space 6: const int ix = get_global_id(0); 7: const int iy = get_global_id(1); 8:? 9: int xc = W/2; 10: int yc = H/2; 11:? 12: int xpos = ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc; 13: int ypos = (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc; 14:? 15: if ((xpos>=0) && (xpos< W) && (ypos>=0) && (ypos< H)) //Bound Checking 16: { 17: dest_data[ypos*W+xpos]= src_data[iy*W+ix]; 18: } 19: } 20:?

src_data為原始圖像（灰度圖）數據，dest_data為旋轉后的圖像數據。W、H分別為圖像的高度和寬度。sinTheta和cosTheta是旋轉參數。我在代碼中實現了旋轉90度，所以sinTheta為1，cosTheta為0，大家可以嘗試其它的值。

下面是程序的流程圖：

在前面向量加法的例子中，我已經介紹了OpenCL一些基本的步驟。

創建platform對象

創建GPU設備

創建contex

創建命令隊列

創建緩沖對象，代碼如下： 1: cl_mem d_ip = clCreateBuffer( 2: context, CL_MEM_READ_ONLY, 3: mem_size, 4: NULL, NULL); 5: l_mem d_op = clCreateBuffer( 6: context, CL_MEM_WRITE_ONLY, 7: mem_size, 8: NULL, NULL); 9: status = clEnqueueWriteBuffer ( 10: queue , d_ip, CL_TRUE, 11: 0, mem_size, (void *)src_image, 12: 0, NULL, NULL);

創建程序對象

編譯程序對象

創建Kernel對象

設置kernel參數

執行kernel

數據拷貝回host memory，我采用映射memory的方式。

1: unsigned char *op_data=0; 2: //op_data =(unsigned char *)malloc(mem_size); 3: status = clEnqueueReadBuffer( 4: //queue, d_op, 5: //CL_TRUE, //Blocking Read Back 6: //0, mem_size,(void*)op_data, NULL, NULL, NULL); 7: op_data = (cl_uchar *) clEnqueueMapBuffer( queue, 8: d_op, 9: CL_TRUE, 10: CL_MAP_READ, 11: 0, 12: mem_size, 13: 0, NULL, NULL, NULL );

?? kernel執行時間的計算后面教程會有詳細介紹，但在本節中，我們會給出通過事件機制來得到kernel執行時間，首先要在創建隊列時候，使用CL_QUEUE_PROFILING_ENABLE參數，否則計算的kernel運行時間是0。

?? 下面是代碼：

1: //計算kerenl執行時間 2: cl_ulong startTime, endTime; clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START, 4: sizeof(cl_ulong), &startTime, NULL); 5: clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END, 6: sizeof(cl_ulong), &endTime, NULL); 7: cl_ulong kernelExecTimeNs = endTime-startTime; 8: printf("kernal exec time :%8.6f ms\n ", kernelExecTimeNs*1e-6 );

?

完整的程序代碼：

1: #include "stdafx.h" 2: #include <CL/cl.h> 3: #include <stdio.h> 4: #include <stdlib.h> 5: #include <time.h> 6: #include <iostream> 7: #include <fstream> 8:? 9: #include "gFreeImage.h" 10:? 11: using namespace std; 12: #define NWITEMS 4 13: #pragma comment (lib,"OpenCL.lib") 14: #pragma comment(lib,"FreeImage.lib") 15:? 16: //把文本文件讀入一個string中 17: int convertToString(const char *filename, std::string& s) 18: { 19: size_t size; 20: char* str; 21:? 22: std::fstream f(filename, (std::fstream::in | std::fstream::binary)); 23:? 24: if(f.is_open()) 25: { 26: size_t fileSize; 27: f.seekg(0, std::fstream::end); 28: size = fileSize = (size_t)f.tellg(); 29: f.seekg(0, std::fstream::beg); 30:? 31: str = new char[size+1]; 32: if(!str) 33: { 34: f.close(); 35: return NULL; 36: } 37:? 38: f.read(str, fileSize); 39: f.close(); 40: str[size] = '\0'; 41: 42: s = str; 43: delete[] str; 44: return 0; 45: } 46: printf("Error: Failed to open file %s\n", filename); 47: return 1; 48: } 49:? 50: //CPU旋轉圖像 51: void cpu_rotate(unsigned char* inbuf, unsigned char* outbuf, int w, int h,float sinTheta, float cosTheta) 52: { 53: int i, j; 54: int xc = w/2; 55: int yc = h/2; 56:? 57: for(i = 0; i < h; i++) 58: { 59: for(j=0; j< w; j++) 60: { 61: int xpos = ( j-xc)*cosTheta - (i-yc)*sinTheta+xc; 62: int ypos = (j-xc)*sinTheta + ( i-yc)*cosTheta+yc; 63:? 64: if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h) 65: outbuf[ypos*w + xpos] = inbuf[i*w+j]; 66: } 67: } 68: } 69:? 70: int main(int argc, char* argv[]) 71: { 72: //裝入圖像 73: unsigned char *src_image=0; 74: unsigned char *cpu_image=0; 75: int W, H; 76: gFreeImage img; 77: if(!img.LoadImageGrey("lenna.jpg")) 78: { 79: printf("裝入lenna.jpg失敗\n"); 80: exit(0); 81: } 82: else 83: src_image = img.getImageDataGrey(W, H); 84:? 85: size_t mem_size = W*H; 86: cpu_image = (unsigned char*)malloc(mem_size); 87:? 88: cl_uint status; 89: cl_platform_id platform; 90:? 91: //創建平臺對象 92: status = clGetPlatformIDs( 1, &platform, NULL ); 93:? 94: cl_device_id device; 95:? 96: //創建GPU設備 97: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 98: 1, 99: &device, 100: NULL); 101: //創建context 102: cl_context context = clCreateContext( NULL, 103: 1, 104: &device, 105: NULL, NULL, NULL); 106: //創建命令隊列 107: cl_command_queue queue = clCreateCommandQueue( context, 108: device, 109: CL_QUEUE_PROFILING_ENABLE, NULL ); 110:? 111: //創建三個OpenCL內存對象，并把buf1的內容通過隱式拷貝的方式 112: //拷貝到clbuf1,buf2的內容通過顯示拷貝的方式拷貝到clbuf2 113: cl_mem d_ip = clCreateBuffer( 114: context, CL_MEM_READ_ONLY, 115: mem_size, 116: NULL, NULL); 117: cl_mem d_op = clCreateBuffer( 118: context, CL_MEM_WRITE_ONLY, 119: mem_size, 120: NULL, NULL); 121: status = clEnqueueWriteBuffer ( 122: queue , d_ip, CL_TRUE, 123: 0, mem_size, (void *)src_image, 124: 0, NULL, NULL); 125:? 126: const char * filename = "rotate.cl"; 127: std::string sourceStr; 128: status = convertToString(filename, sourceStr); 129: const char * source = sourceStr.c_str(); 130: size_t sourceSize[] = { strlen(source) }; 131:? 132: //創建程序對象 133: cl_program program = clCreateProgramWithSource( 134: context, 135: 1, 136: &source, 137: sourceSize, 138: NULL); 139: //編譯程序對象 140: status = clBuildProgram( program, 1, &device, NULL, NULL, NULL ); 141: if(status != 0) 142: { 143: printf("clBuild failed:%d\n", status); 144: char tbuf[0x10000]; 145: clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL); 146: printf("\n%s\n", tbuf); 147: return -1; 148: } 149:? 150:? 151: //創建Kernel對象 152: //Use the “image_rotate” function as the kernel 153:? 154: //創建Kernel對象 155: cl_kernel kernel = clCreateKernel( program, "image_rotate", NULL ); 156:? 157: //設置Kernel參數 158: float sintheta = 1, costheta = 0; 159: clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_ip); 160: clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&d_op); 161: clSetKernelArg(kernel, 2, sizeof(cl_int), (void *)&W); 162: clSetKernelArg(kernel, 3, sizeof(cl_int), (void *)&H); 163: clSetKernelArg(kernel, 4, sizeof(cl_float), (void *)&sintheta); 164: clSetKernelArg(kernel, 5, sizeof(cl_float), (void *)&costheta); 165:? 166: //Set local and global workgroup sizes 167: size_t localws[2] = {16,16} ; 168: size_t globalws[2] = {W, H};//Assume divisible by 16 169:? 170: cl_event ev; 171: //執行kernel 172: clEnqueueNDRangeKernel( 173: queue ,kernel, 174: 2, 0, globalws, localws, 175: 0, NULL, &ev); 176:? 177: clFinish( queue ); 178:? 179: //計算kerenl執行時間 180: cl_ulong startTime, endTime; 181: clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START, 182: sizeof(cl_ulong), &startTime, NULL); 183: clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END, 184: sizeof(cl_ulong), &endTime, NULL); 185: cl_ulong kernelExecTimeNs = endTime-startTime; 186: printf("kernal exec time :%8.6f ms\n ", kernelExecTimeNs*1e-6 ); 187:? 188: //數據拷回host內存 189: // copy results from device back to host 190: unsigned char *op_data=0; 191: //op_data =(unsigned char *)malloc(mem_size); 192: // status = clEnqueueReadBuffer( 193: //queue, d_op, 194: //CL_TRUE, //Blocking Read Back 195: //0, mem_size,(void*)op_data, NULL, NULL, NULL); 196: op_data = (cl_uchar *) clEnqueueMapBuffer( queue, 197: d_op, 198: CL_TRUE, 199: CL_MAP_READ, 200: 0, 201: mem_size, 202: 0, NULL, NULL, NULL ); 203:? 204: int i; 205: cpu_rotate(src_image,cpu_image, W, H, 1, 0); 206: for(i = 0; i < mem_size; i++) 207: { 208: src_image[i] =cpu_image[i]; 209: } 210: img.SaveImage("cpu_lenna_rotate.jpg"); 211: for(i = 0; i < mem_size; i++) 212: { 213: src_image[i] =op_data[i]; 214: } 215: img.SaveImage("lenna_rotate.jpg"); 216: 217: if(cpu_image) 218: free(cpu_image); 219:? 220: //刪除OpenCL資源對象 221: clReleaseMemObject(d_ip); 222: clReleaseMemObject(d_op); 223: clReleaseProgram(program); 224: clReleaseCommandQueue(queue); 225: clReleaseContext(context); 226: return 0; 227: } 228:?

感興趣的朋友可以從http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode2.zip&can=2&q=#makechanges下載完整代碼。

注意代碼運行后，會在程序目錄生成lenna_rotate.jpg,這時gpu執行的結果，另外還有一個cpu_lenna_rotate.jpg這是CPU執行的結果。

3、一個矩陣乘法的例子

??? 在amd的slides中，本節還講了一個簡單的，沒有優化的矩陣乘法，一共才2兩頁ppt，所以我也不在這兒詳細講述了，…,但簡單介紹還是需要的。

1: for(int i = 0; i < Ha; i++) 2: for(int j = 0; j < Wb; j++){ 3: c[i][j] = 0; 4: for(int k = 0; k < Wa; k++) 5: c[i][j] += a[i][k] + b[k][j] 6: }

上面的代碼是矩陣乘法的例子，有三重循環，下面我們只給出kernel代碼，完整程序請從：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicodeCode3.zip&can=2&q=#makechanges下載。

1: __kernel void simpleMultiply( 2: __global float* c, int Wa, int Wb, 3: __global float* a, __global float* b) 4: { 5:? 6: //Get global position in Y direction 7: int row = get_global_id(1); 8: //Get global position in X direction 9: int col = get_global_id(0); 10: float sum = 0.0f; 11: //Calculate result of one element 12: for (int i = 0; i < Wa; i++) 13: {

從今天開始學習OpenCL……

??????? 因為老狼的顯卡是AMD 5xx的redwood,所以下面先介紹OpenCL APP(Accelerated Parallel processing)的安裝。

下載地址：http://developer.amd.com/tools/hc/AMDAPPSDK/downloads/Pages/default.aspx

安裝注意事項：http://developer.amd.com/tools/hc/AMDAPPSDK/assets/AMD_APP_SDK_Installation_Notes.pdf

有用的資料：http://developer.amd.com/tools/hc/AMDAPPSDK/documentation/Pages/default.aspx

其中最有用的是下面幾個文檔：

???? AMD最新顯卡Tahiti的ISA介紹，對OpenCL編程優化有用。

http://developer.amd.com/wordpress/media/2012/12/AMD_Southern_Islands_Instruction_Set_Architecture.pdf

?

????? 這本書是最好的OpenCL教程，好過市面上的任何一本OpenCL書，其中包括很多優化Kernel代碼的內容，我計劃以后就按照這本書的內容來學習opencl。

http://developer.amd.com/download/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide.pdf

?

????? 再就是OpenCL 1.2的spec了，下載地址：OpenCL? 1.2 Specification (revision 15) ，相對于1.1來說，1.2中還是有一些變化的，比如我以前寫的程序中CreateImage2D函數發現在1.2中沒有了，spec其實就是一個函數手冊，偶爾用來查詢一下而已。

??? 另外一個比較好的初級教程，就是我翻譯的AMD OpenCL大學教程了，http://www.cnblogs.com/mikewolf2002/archive/2012/01/30/2332356.html

現在，我們開始寫一個簡單的OpenCL程序，計算兩個數組相加的和，放到另一個數組中去。程序用cpu和gpu分別計算，最后驗證它們是否相等。OpenCL程序的流程大致如下：

下面是source code中的主要代碼：

?

int main(int argc, char* argv[])
??? {
??? //在host內存中創建三個緩沖區
??? float *buf1 = 0;
??? float *buf2 = 0;
??? float *buf = 0;

??? buf1 =(float *)malloc(BUFSIZE * sizeof(float));
??? buf2 =(float *)malloc(BUFSIZE * sizeof(float));
??? buf =(float *)malloc(BUFSIZE * sizeof(float));

??? //用一些隨機值初始化buf1和buf2的內容
??? int i;
??? srand( (unsigned)time( NULL ) );
??? for(i = 0; i < BUFSIZE; i++)
??????? buf1[i] = rand()%65535;

??? srand( (unsigned)time( NULL ) +1000);
??? for(i = 0; i < BUFSIZE; i++)
??????? buf2[i] = rand()%65535;

??? //cpu計算buf1,buf2的和
??? for(i = 0; i < BUFSIZE; i++)
??????? buf[i] = buf1[i] + buf2[i];

??? cl_uint status;
??? cl_platform_id platform;

??? //創建平臺對象
??? status = clGetPlatformIDs( 1, &platform, NULL );

????? 注意：如果我們系統中安裝不止一個opencl平臺，比如我的os中，有intel和amd兩家opencl平臺，用上面這行代碼，有可能會出錯，因為它得到了intel的opencl平臺，而intel的平臺只支持cpu，而我們后面的操作都是基于gpu，這時我們可以用下面的代碼，得到AMD的opencl平臺。
?

cl_uint numPlatforms;std::string platformVendor; status = clGetPlatformIDs(0, NULL, &numPlatforms);if(status != CL_SUCCESS){return 0;}if (0 < numPlatforms) {cl_platform_id* platforms = new cl_platform_id[numPlatforms];status = clGetPlatformIDs(numPlatforms, platforms, NULL);char platformName[100];for (unsigned i = 0; i < numPlatforms; ++i) {status = clGetPlatformInfo(platforms[i],CL_PLATFORM_VENDOR,sizeof(platformName),platformName,NULL);platform = platforms[i];platformVendor.assign(platformName);if (!strcmp(platformName, "Advanced Micro Devices, Inc.")) {break;}}std::cout << "Platform found : " << platformName << "\n";delete[] platforms;}

??? cl_device_id device;

??? //創建GPU設備
??? clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,
??????? 1,
??????? &device,
??????? NULL);
??? //創建context
??? cl_context context = clCreateContext( NULL,
??????? 1,
??????? &device,
??????? NULL, NULL, NULL);
??? //創建命令隊列
??? cl_command_queue queue = clCreateCommandQueue( context,
??????? device,
??????? CL_QUEUE_PROFILING_ENABLE, NULL );
??? //創建三個OpenCL內存對象，并把buf1的內容通過隱式拷貝的方式
??? //buf1內容拷貝到clbuf1,buf2的內容通過顯示拷貝的方式拷貝到clbuf2
??? cl_mem clbuf1 = clCreateBuffer(context,
??????? CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
??????? BUFSIZE*sizeof(cl_float),buf1,
??????? NULL );

??? cl_mem clbuf2 = clCreateBuffer(context,
??????? CL_MEM_READ_ONLY ,
??????? BUFSIZE*sizeof(cl_float),NULL,
??????? NULL );

?? cl_event writeEvt;

??? status = clEnqueueWriteBuffer(queue, clbuf2, 1,
??????? 0, BUFSIZE*sizeof(cl_float), buf2, 0, 0, 0);

??? 上面這行代碼把buf2中的內容拷貝到clbuf2,因為buf2位于host端，clbuf2位于device端，所以這個函數會執行一次host到device的傳輸操作，或者說一次system memory到video memory的拷貝操作，所以我在該函數的后面放置了clFush函數，表示把command queue中的所有命令提交到device(注意：該命令并不保證命令執行完成),所以我們調用函數waitForEventAndRelease來等待write緩沖的完成，waitForEventAndReleae是一個用戶定義的函數，它的內容如下，主要代碼就是通過event來查詢我們的操作是否完成，沒完成的話，程序就一直block在這行代碼處，另外我們也可以用opencl中內置的函數clWaitForEvents來代替clFlush和waitForEventAndReleae。

//等待事件完成 int waitForEventAndRelease(cl_event *event){cl_int status = CL_SUCCESS;cl_int eventStatus = CL_QUEUED;while(eventStatus != CL_COMPLETE){status = clGetEventInfo(*event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int),&eventStatus,NULL);}status = clReleaseEvent(*event);return 0;}

???? status = clFlush(queue);
???? //等待數據傳輸完成再繼續往下執行
???? waitForEventAndRelease(&writeEvt);

??? cl_mem buffer = clCreateBuffer( context,
??????? CL_MEM_WRITE_ONLY,
??????? BUFSIZE * sizeof(cl_float),
??????? NULL, NULL );

????? kernel文件中放的是gpu中執行的代碼，它被放在一個單獨的文件add.cl中，本程序中kernel代碼非常簡單，只是執行兩個數組相加。kernel的代碼為：

__kernel void vecadd(__global const float* A, __global const float* B, __global float* C) {int id = get_global_id(0);C[id] = A[id] + B[id]; }

?? //kernel文件為add.cl
??? const char * filename? = "add.cl";
??? std::string? sourceStr;
??? status = convertToString(filename, sourceStr);

convertToString也是用戶定義的函數，該函數把kernel源文件讀入到一個string中，它的代碼如下：

//把文本文件讀入一個string中,用來讀入kernel源文件 int convertToString(const char *filename, std::string& s){size_t size;char* str;std::fstream f(filename, (std::fstream::in | std::fstream::binary));if(f.is_open()){size_t fileSize;f.seekg(0, std::fstream::end);size = fileSize = (size_t)f.tellg();f.seekg(0, std::fstream::beg);str = new char[size+1];if(!str){f.close();return NULL;}f.read(str, fileSize);f.close();str[size] = '\0';s = str;delete[] str;return 0;}printf("Error: Failed to open file %s\n", filename);return 1;}

??? const char * source??? = sourceStr.c_str();
??? size_t sourceSize[]??? = { strlen(source) };

??? //創建程序對象
??? cl_program program = clCreateProgramWithSource(
??????? context,
??????? 1,
??????? &source,
??????? sourceSize,
??????? NULL);
??? //編譯程序對象
??? status = clBuildProgram( program, 1, &device, NULL, NULL, NULL );
??? if(status != 0)
??????? {
??????? printf("clBuild failed:%d\n", status);
??????? char tbuf[0x10000];
??????? clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL);
??????? printf("\n%s\n", tbuf);
??????? return -1;
??????? }

??? //創建Kernel對象
??? cl_kernel kernel = clCreateKernel( program, "vecadd", NULL );
??? //設置Kernel參數
??? cl_int clnum = BUFSIZE;
??? clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1);
??? clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2);
??? clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &buffer);

注意：在執行kernel時候，我們只設置了global work items數量，沒有設置group size，這時候，系統會使用默認的work group size，通?？赡苁?56之類的。

??? //執行kernel,Range用1維，work itmes size為BUFSIZE
??? cl_event ev;
??? size_t global_work_size = BUFSIZE;
??? clEnqueueNDRangeKernel( queue,
??????? kernel,
??????? 1,
??????? NULL,
??????? &global_work_size,
??????? NULL, 0, NULL, &ev);
?? status = clFlush( queue );
?? waitForEventAndRelease(&ev);

??? //數據拷回host內存
??? cl_float *ptr;

??? cl_event mapevt;
??? ptr = (cl_float *) clEnqueueMapBuffer( queue,
??????? buffer,
??????? CL_TRUE,
??????? CL_MAP_READ,
??????? 0,
??????? BUFSIZE * sizeof(cl_float),
??????? 0, NULL, NULL, NULL );

?? status = clFlush( queue );
?? waitForEventAndRelease(&mapevt);

???
??? //結果驗證，和cpu計算的結果比較
??? if(!memcmp(buf, ptr, BUFSIZE))
??????? printf("Verify passed\n");
??? else printf("verify failed");

??? if(buf)
??????? free(buf);
??? if(buf1)
??????? free(buf1);
??? if(buf2)
??????? free(buf2);

????? 程序結束后，這些opencl對象一般會自動釋放，但是為了程序完整，養成一個好習慣，這兒我加上了手動釋放opencl對象的代碼。

??? //刪除OpenCL資源對象
??? clReleaseMemObject(clbuf1);
??? clReleaseMemObject(clbuf2);
??? clReleaseMemObject(buffer);
??? clReleaseProgram(program);
??? clReleaseCommandQueue(queue);
??? clReleaseContext(context);
??? return 0;
??? }

程序執行后的界面如下：

完整的代碼請參考：

工程文件gclTutorial1

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip

在教程2中，我們通過函數convertToString，把kernel源文件讀到一個string串中，然后用函數clCreateProgramWithSource裝入程序對象，再調用函數clBuildProgram編譯程序對象。其實我們也可以直接調用二進制kernel文件，這樣，當不想把kernel文件給別人看的時候，起到一定的保密作用。在本教程中，我們會把讀入的源文件存儲一個二進制文件中，并且還會建立一個計時器類，用來記錄數組加法在cpu和gpu端分別執行的時間。

???? 首先我們建立工程文件gclTutorial2,在其中增加類gclFile,該類主要用來讀取文本kernel文件，或者讀寫二進制kernel文件。

class gclFile
{
public:
??? gclFile(void);
??? ~gclFile(void);

??? //打開opencl kernel源文件(文本模式)
??? bool open(const char* fileName);

??? //讀寫二進制kernel文件
??? bool writeBinaryToFile(const char* fileName, const char* birary, size_t numBytes);
??? bool readBinaryFromFile(const char* fileName);

…

}

gclFile中三個讀寫kernel文件的函數代碼為：

bool gclFile::writeBinaryToFile(const char* fileName, const char* birary, size_t numBytes) {FILE *output = NULL;output = fopen(fileName, "wb");if(output == NULL)return false;fwrite(birary, sizeof(char), numBytes, output);fclose(output);return true; }bool gclFile::readBinaryFromFile(const char* fileName) {FILE * input = NULL;size_t size = 0;char* binary = NULL;input = fopen(fileName, "rb");if(input == NULL){return false;}fseek(input, 0L, SEEK_END); size = ftell(input);//指向文件起始位置rewind(input);binary = (char*)malloc(size);if(binary == NULL){return false;}fread(binary, sizeof(char), size, input);fclose(input);source_.assign(binary, size);free(binary);return true; }bool gclFile::open(const char* fileName) //!< file name {size_t size;char* str;//以流方式打開文件std::fstream f(fileName, (std::fstream::in | std::fstream::binary));// 檢查是否打開了文件流if (f.is_open()){size_t sizeFile;// 得到文件sizef.seekg(0, std::fstream::end);size = sizeFile = (size_t)f.tellg();f.seekg(0, std::fstream::beg);str = new char[size + 1];if (!str){f.close();return false;}// 讀文件f.read(str, sizeFile);f.close();str[size] = '\0';source_ = str;delete[] str;return true;}return false; }

現在，在main.cpp中，我們就可以用gclFile類的open函數來讀入kernel源文件了：

//kernel文件為add.cl

gclFile kernelFile;
if(!kernelFile.open("add.cl"))
??? {
??? printf("Failed to load kernel file \n");
??? exit(0);
??? }
const char * source = kernelFile.source().c_str();
size_t sourceSize[] = {strlen(source)};
//創建程序對象
cl_program program = clCreateProgramWithSource(
??? context,
??? 1,
??? &source,
??? sourceSize,
??? NULL);

??? 編譯好kernel后，我們可以通過下面的代碼，把編譯好的kernel存儲在一個二進制文件addvec.bin中，在教程4種，我們將會直接裝入這個二進制的kernel文件。

//存儲編譯好的kernel文件 char **binaries = (char **)malloc( sizeof(char *) * 1 ); //只有一個設備 size_t *binarySizes = (size_t*)malloc( sizeof(size_t) * 1 );status = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES,sizeof(size_t) * 1, binarySizes, NULL); binaries[0] = (char *)malloc( sizeof(char) * binarySizes[0]); status = clGetProgramInfo(program, CL_PROGRAM_BINARIES,sizeof(char *) * 1, binaries, NULL); kernelFile.writeBinaryToFile("vecadd.bin", binaries[0],binarySizes[0]);

??? 我們還會建立一個計時器類gclTimer,用來統計時間，這個類主要用QueryPerformanceFrequency得到時鐘頻率，用QueryPerformanceCounter得到流逝的ticks數，最終得到流逝的時間。函數非常簡單，

class gclTimer
{
public:
??? gclTimer(void);
??? ~gclTimer(void);

private:

??? double _freq;
??? double _clocks;
??? double _start;
public:
??? void Start(void); // 啟動計時器
??? void Stop(void); //停止計時器
??? void Reset(void); //復位計時器
??? double GetElapsedTime(void); //計算流逝的時間
};

下面我們在cpu端執行數組加法時，增加計時器的代碼：

gclTimer clTimer;
clTimer.Reset();
clTimer.Start();
//cpu計算buf1,buf2的和
for(i = 0; i < BUFSIZE; i++)
??? buf[i] = buf1[i] + buf2[i];
clTimer.Stop();
printf("cpu costs time:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );

同理在gpu執行kernel代碼，以及copy gpu結果到cpu時候，增加計時器代碼：

//執行kernel,Range用1維，work itmes size為BUFSIZE, cl_event ev; size_t global_work_size = BUFSIZE;clTimer.Reset(); clTimer.Start(); clEnqueueNDRangeKernel( queue,kernel,1,NULL,&global_work_size,NULL, 0, NULL, &ev); status = clFlush( queue ); waitForEventAndRelease(&ev);//clWaitForEvents(1, &ev);clTimer.Stop(); printf("kernal total time:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );//數據拷回host內存 cl_float *ptr; clTimer.Reset(); clTimer.Start(); cl_event mapevt; ptr = (cl_float *) clEnqueueMapBuffer( queue,buffer,CL_TRUE,CL_MAP_READ,0,BUFSIZE * sizeof(cl_float),0, NULL, &mapevt, NULL ); status = clFlush( queue ); waitForEventAndRelease(&mapevt);//clWaitForEvents(1, &mapevt);clTimer.Stop(); printf("copy from device to host:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );

最終程序執行界面如下，在bufsize為262144時，在我的顯卡上gpu還有cpu快呢…，在程序目錄，我們可以看到也產生了vecadd.bin文件了。

完整的代碼請參考：

工程文件gclTutorial2

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip

本教程中，我們使用上一篇教程中產生的二進制kernel文件vecadd.bin作為輸入來創建程序對象，程序代碼如下：

//kernel文件為vecadd.bin
gclFile kernelFile;
if(!kernelFile.readBinaryFromFile("vecadd.bin"))
??? {
??? printf("Failed to load binary file \n");
??? exit(0);
??? }
const char * binary = kernelFile.source().c_str();
size_t binarySize = kernelFile.source().size();

cl_program program = clCreateProgramWithBinary(context,
??? 1,
??? &device,
??? (const size_t *)&binarySize,
??? (const unsigned char**)&binary,
??? NULL,
??? NULL);

程序執行的界面和教程3中一摸一樣…

完整的代碼請參考：

工程文件gclTutorial3

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip

在本教程中，我們使用二維NDRange來設置workgroup，這樣在opencl中，workitme的組織形式是二維的，Kernel中 的代碼也要做相應的改變，我們先看一下clEnqueueNDRangeKernel函數的變化。首先我們指定了workgroup size為localx*localy,通常這個值為64的倍數，但最好不要超過256。

//執行kernel,Range用2維，work itmes size為width*height,
cl_event ev;
size_t globalThreads[] = {width, height};
size_t localx, localy;
if(width/8 > 4)
??? localx = 16;
else if(width < 8)
??? localx = width;
else localx = 8;

if(height/8 > 4)
??? localy = 16;
else if (height < 8)
??? localy = height;
else localy = 8;

size_t localThreads[] = {localx, localy};// localx*localy應該是64的倍數
printf("global_work_size =(%d,%d), local_work_size=(%d, %d)\n",width,height,localx,localy);

clTimer.Reset();
clTimer.Start();
clEnqueueNDRangeKernel( queue,
??? kernel,
??? 2,
??? NULL,
??? globalThreads,
??? localThreads, 0, NULL, &ev);

注意：在上面代碼中，定義global threads以及local threads數量，都是通過二維數組的方式進行的。

??? 新的Kernel代碼如下：

#pragma OPENCL EXTENSION cl_amd_printf : enable__kernel void vecadd(__global const float* a, __global const float* b, __global float* c) {int x = get_global_id(0);int y = get_global_id(1);int width = get_global_size(0);int height = get_global_size(1);if(x == 1 && y ==1)printf("%d, %d,%d,%d,%d,%d\n",get_local_size(0),get_local_size(1),get_local_id(0),get_local_id(1),get_group_id(0),get_group_id(1));c[x + y * width] = a[x + y * width] + b[x + y * width];}

????? 我們在kernel中增加了#pragma OPENCL EXTENSION cl_amd_printf : enable，以便在kernel中通過printf函數進行debug，這是AMD的一個擴展。printf還可以直接打印出float4這樣的向量，比如printf(“%v4f”, vec)。

????? 另外，在main.cpp中增加一行代碼：

//告訴driver dump il和isa文件
_putenv("GPU_DUMP_DEVICE_KERNEL=3");
????? 我們可以在程序目錄dump出il和isa形式的kernel文件，對于熟悉isa匯編的人，這是一個很好的調試performance的方法。

???? 在最新的app sdk 2.7中，在kernel中使用printf的時候，這個程序會hang在哪兒，以前沒這種情況。

程序執行界面。

完整的代碼請參考：

工程文件gclTutorial4

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip

在本教程中，我們學習用opencl進行簡單的圖像處理，對一個圖片進行旋轉。圖片讀入、保存等工作，我們使用開源的FreeImage，下載地址： http://freeimage.sourceforge.net/

????? 首先我們建立一個gFreeImage類，用來裝入圖像，該類主要調用FreeImage的函數，首先會初始化FreeImage庫，然后根據文件名猜測圖像文件格式，最終load圖像文件到變量FIBITMAP *bitmap中去。同時，我們還定義了2個緩沖

unsigned char *imageData;
unsigned char *imageData4;

用來存放圖像數據，之所以定義imageData4,是因為通常的圖片文件，比如jpg,bmp都是3個通道，沒有包括alpha通道，但是在gpu中處理數據時候，通常以vector4或者vector的形式進行，不能以vector3進行，所以我們裝入圖像后，會把imageData指向圖像數據，同時生成包括alpha通道的圖像數據imageData4。

???? 另外，我們還定義了一個函數LoadImageGrey，該函數用來裝入灰度圖，灰度圖一個像素用一個uchar表示。

在main.cpp中，我們首先定義一個cpu處理圖像旋轉的函數：

//CPU旋轉圖像
void cpu_rotate(unsigned char* inbuf, unsigned char* outbuf, int w, int h,float sinTheta, float cosTheta)
??? {
??? int i, j;
??? int xc = w/2;
??? int yc = h/2;

??? for(i = 0; i < h; i++)
??????? {
??????? for(j=0; j< w; j++)
??????????? {
??????????? int xpos =? ( j-xc)*cosTheta - (i-yc)*sinTheta+xc;???
??????????? int ypos =? (j-xc)*sinTheta + ( i-yc)*cosTheta+yc;

??????????? if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h)
??????????????? outbuf[ypos*w + xpos] = inbuf[i*w+j];
??????????? }
??????? }
??? }

??? 在main函數中，我們首先會裝入圖像文件，代碼如下：

int W, H; gFreeImage img; if(!img.LoadImageGrey("lenna.jpg")){printf("can‘t load lenna.jpg\n");exit(0);} elsesrc_image = img.getImageDataGrey(W, H);size_t mem_size = W*H; cpu_image = (unsigned char*)malloc(mem_size);

??? 之后，定義2個cl memory對象，一個用來放原始圖像，一個用來放旋轉后的圖像。

//創建2個OpenCL內存對象
cl_mem d_ip = clCreateBuffer(
??? context, CL_MEM_READ_ONLY,
??? mem_size,
??? NULL, NULL);
cl_mem d_op = clCreateBuffer(
??? context, CL_MEM_WRITE_ONLY,
??? mem_size,
??? NULL, NULL);

cl_event writeEvt;
status = clEnqueueWriteBuffer (???
??? queue , d_ip, CL_TRUE,
??? 0, mem_size, (void *)src_image,
??? 0, NULL, &writeEvt);
//等待數據傳輸完成再繼續往下執行
status = clFlush(queue);
waitForEventAndRelease(&writeEvt);
//clWaitForEvents(1, &writeEvt);

?? 旋轉kernel函數需要傳入6個參數：

//創建Kernel對象
cl_kernel kernel = clCreateKernel( program, "image_rotate", NULL );
//設置Kernel參數
float sintheta = 1, costheta = 0;
clSetKernelArg(kernel, 0, sizeof(cl_mem),? (void *)&d_ip);
clSetKernelArg(kernel, 1, sizeof(cl_mem),? (void *)&d_op);
clSetKernelArg(kernel, 2, sizeof(cl_int),? (void *)&W);
clSetKernelArg(kernel, 3, sizeof(cl_int),? (void *)&H);
clSetKernelArg(kernel, 4, sizeof(cl_float), (void *)&sintheta);
clSetKernelArg(kernel, 5, sizeof(cl_float), (void *)&costheta);
kernel執行的代碼為：

//執行kernel,Range用2維，work itmes size為W*H,
cl_event ev;
size_t globalThreads[] = {W, H};
size_t localThreads[] = {16, 16}; // localx*localy應該是64的倍數
printf("global_work_size =(%d,%d), local_work_size=(16, 16)\n",W,H);

clTimer.Reset();
clTimer.Start();
clEnqueueNDRangeKernel( queue,
??? kernel,
??? 2,
??? NULL,
??? globalThreads,
??? localThreads, 0, NULL, &ev);
//沒有設置local group size時候，系統將會自動設置為 (256,1)
status = clFlush( queue );
waitForEventAndRelease(&ev);
//clWaitForEvents(1, &ev);

clTimer.Stop();
printf("kernal total time:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );

kernel函數代碼為：

#pragma OPENCL EXTENSION cl_amd_printf : enable __kernel void image_rotate( __global uchar * src_data, __global uchar * dest_data, //源圖像和輸出圖像都放在global memory中int W, int H, //圖像sizefloat sinTheta, float cosTheta ) //旋轉角度 { 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]; } }

gpu執行完畢后，旋轉后的圖像保存在lenna_rotate.jpg,我們還會用cpu rotate函數執行一次旋轉，同時把生成的圖像保存到cpu_lenna_rotate.jpg。

完整的代碼請參考：

工程文件gclTutorial5

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip

histogram翻譯成中文就是直方圖，在計算機圖像處理和視覺技術中，通常用histogram來進行圖像匹配，從而完成track，比如meanshift跟蹤算法中，經常要用到圖像的直方圖。

???? 灰度圖的histogram計算，首先要選擇bin(中文可以稱作槽)的數量，對于灰度圖，像素的范圍通常是[0-255],所以bin的數目就是256，然后我們循環整幅圖像，統計出每種像素值出現的次數，放到對應的bin中。比如bin[0]中放的就是整幅圖像中灰度值為0的像素個數，bin[1]中放的就是整幅圖像中灰度值為1的像素個數……

???? 下面的直方圖就是灰度圖lenna對應的直方圖。

????? 灰度圖直方圖的cpu計算特別簡單，定義一個數組hostBin[256],初始化所有數組元素為0，然后循環整幅圖像，得到直方圖，代碼如下：

//cpu求直方圖
void cpu_histgo()
??? {
??? int i, j;
??? for(i = 0; i < height; ++i)
??????? {
??????? for(j = 0; j < width; ++j)
??????????? {
??????????? //printf("data: %d\n", data[i * width + j] );
??????????? hostBin[data[i * width + j]]++;
??????????? //printf("hostbin %d=%d\n", data[i * width + j], hostBin[data[i * width + j]]);
??????????? }
??????? }
??? }

??? 如何使用opencl,來計算灰度圖，就沒有那么簡單了。我們知道gpu的優勢是并行計算，如何把圖像分塊，來并行計算直方圖，是我們討論的重點。下面是一副512*512圖像的thread，workgroup劃分：

???? 我們設定圖像的寬度是bins的整數倍，即256的倍數，高度是workgroup size(本程序中，設置為128)的倍數，如果圖像高寬不是bins和workgroup size的倍數，則我們通過下面的公式把圖像的寬度和高度變成它們的倍數：

//width是binSize的整數倍，height是groupsize的整數倍
width = (width / binSize ? width / binSize: 1) * binSize;
height = (height / groupSize ? height / groupSize: 1) * groupSize;

???? 則512*512的圖像可以分為8個work group，每個workgroup包括128個thread，每個thread計算256個像素的直方圖，并把結果放到該thread對應的local memroy空間，在kenrel代碼結束前，合并一個workgroup中所有thread的直方圖，生成一個workgroup塊的直方圖，最后在host端，合并8個workgroup塊的直方圖，產生最終的直方圖。

??? openCL的memory對象主要有3個，dataBuffer用來傳入圖像數據，而minDeviceBinBuf大小是workgroup number *256, 即每個workgroup對應一個bin，另外一個kernel函數的第二個參數，它的大小為workgroup size*256, 用于workgroup中的每個線程存放自己256個像素的直方圖結果。

//創建2個OpenCL內存對象
dataBuf = clCreateBuffer(
??? context,
??? CL_MEM_READ_ONLY,
??? sizeof(cl_uchar) * width? * height,
??? NULL,
??? 0);

//該對象存放每個block塊的直方圖結果
midDeviceBinBuf = clCreateBuffer(
??? context,
??? CL_MEM_WRITE_ONLY,
??? sizeof(cl_uint) * binSize * subHistgCnt,
??? NULL,
??? 0);

?? …

??? status = clSetKernelArg(kernel, 1, groupSize * binSize * sizeof(cl_uchar), NULL);//local memroy size, lds for amd

下面看看kernel代碼是如何計算workgroup塊的直方圖。

__kernel
void histogram256(__global const uchar* data,
????????????????? __local uchar* sharedArray,
????????????????? __global uint* binResult)
{
??? size_t localId = get_local_id(0);
??? size_t globalId = get_global_id(0);
??? size_t groupId = get_group_id(0);
??? size_t groupSize = get_local_size(0);

下面這部分代碼初始化每個thread對應的local memory，也就是對應的256個bin中計數清零。sharedArray大小是workgroup size * 256 = 128 * 256

??? //初始化共享內存
??? for(int i = 0; i < BIN_SIZE; ++i)
??????? sharedArray[localId * BIN_SIZE + i] = 0;

通過barrier設置workgroup中所有thread的同步點，保證所有thread都完成初始化操作。

??? barrier(CLK_LOCAL_MEM_FENCE);
???

下面的代碼，計算thread中，256個像素的直方圖，比如對于workgroup 0中的thread 0它計算的256個像素為綠條的部分像素，注意：每個thread的包含的像素并不是連續的。

??? //計算thread直方圖
??? for(int i = 0; i < BIN_SIZE; ++i)
??? {
??????? uint value = (uint)data[groupId * groupSize * BIN_SIZE + i * groupSize + localId];
??????? sharedArray[localId * BIN_SIZE + value]++;
??? }
??? 通過fence，保證每個thread都完成各自的直方圖計算。
??? barrier(CLK_LOCAL_MEM_FENCE);?
??? 下面是合并各個thread的直方圖形成整個workgroup像素塊的直方圖，每個thread合并2個bin，比如thread 0，合并bin0和bin128。

?? //合并workgroup中所有線程的直方圖，產生workgroup直方圖
??? for(int i = 0; i < BIN_SIZE / groupSize; ++i)
??? {
??????? uint binCount = 0;
??????? for(int j = 0; j < groupSize; ++j)
??????????? binCount += sharedArray[j * BIN_SIZE + i * groupSize + localId];
???????????
??????? binResult[groupId * BIN_SIZE + i * groupSize + localId] = binCount;
??? }
}

最終在host端，我們還要把每個workgroup塊的直方圖合并成得到整個圖像的直方圖，主要代碼如下：

// 合并子塊直方圖值

for(i = 0; i < subHistgCnt; ++i)
??? {
??? for( j = 0; j < binSize; ++j)
??????? {
??????? deviceBin[j] += midDeviceBin[i * binSize + j];
??????? }
??? }

完整的代碼請參考：

工程文件gclTutorial7

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip

現在我們利用上一篇教程的方法，來統計一副RGBA圖像中，有多少個像素點（該像素點滿足R, G, B, A任意分量>=5)？我考慮的方法是建立256 bin的直方圖，對于一個像素，求max(R, G,B,A),用該值決定該像素點進入那個bin，這樣求出直方圖后，width*height - hostBin[0] - hostBin[1] - hostBin[2] - hostBin[3] - hostBin[4]，即為我們要的結果。

???? 本教程代碼基本上和上一篇教程中代碼是一樣的，區別主要包括以下2點：

1、我們裝入的是RGBA 彩色圖。

//裝入圖像
unsigned char *src_image=0;
gFreeImage img;
if(!img.LoadImage("../lenna.jpg"))
??? {
??? printf("can‘t load lenna.jpg\n");
??? exit(0);
??? }
src_image = img.getImageData(width,height);

2、在kernel代碼中，有點小變化，kernel代碼如下：

#define LINEAR_MEM_ACCESS #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable #define BIN_SIZE 256/** * 計算直方圖,bins是256 * data 輸入數據 * 一個workgroup內所有thread共享的內存， * 每個workgroup的直方圖 */__kernel void histogram256(__global const uchar4* data,__local uchar* sharedArray,__global uint* binResult) {size_t localId = get_local_id(0);size_t globalId = get_global_id(0);size_t groupId = get_group_id(0);size_t groupSize = get_local_size(0);//初始化共享內存for(int i = 0; i < BIN_SIZE; ++i)sharedArray[localId * BIN_SIZE + i] = 0;barrier(CLK_LOCAL_MEM_FENCE);uchar R, G, B, A, T;//計算thread直方圖for(int i = 0; i < BIN_SIZE; ++i){ #ifdef LINEAR_MEM_ACCESSR = (uint)data[groupId * groupSize * BIN_SIZE + i * groupSize + localId].x;G = (uint)data[groupId * groupSize * BIN_SIZE + i * groupSize + localId].y;B = (uint)data[groupId * groupSize * BIN_SIZE + i * groupSize + localId].z;A = (uint)data[groupId * groupSize * BIN_SIZE + i * groupSize + localId].w;uint value = (uint)max(max(R,G),max(B,A)); #elseuint value = data[globalId * BIN_SIZE + i]; #endif // LINEAR_MEM_ACCESSsharedArray[localId * BIN_SIZE + value]++;}barrier(CLK_LOCAL_MEM_FENCE); //合并workgroup中所有線程的直方圖，產生workgroup直方圖for(int i = 0; i < BIN_SIZE / groupSize; ++i){uint binCount = 0;for(int j = 0; j < groupSize; ++j)binCount += sharedArray[j * BIN_SIZE + i * groupSize + localId];binResult[groupId * BIN_SIZE + i * groupSize + localId] = binCount;} }

?

完整的代碼請參考：

工程文件gclTutorial8

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip

在opencl編程中，特別是基于gpu的opencl的編程，提高程序性能最主要的方法就是想法提高memory的利用率：一個是提高global memory的合并讀寫效率，另一個就是減少local memory的bank conflit。下面我們分析一下教程7中的代碼，其的memory利用率如何？

??? 首先我們用amd的opencl profiler分析一下程序性能(不會找不到用吧，點擊view-other windows-app profiler…,然后就看到了…)。

???? 下面我們來分析我們的kernel代碼中memory操作：

???? 首先是shared memory的的初始化，我們知道shared memory是local memory，被一個workgroup中的所有thread，或者說work item共享。在amd硬件系統中，local memory是LDS，它通常是為32k，分為32個bank，dword字節地址，每個bank 512個item，我們可以通過函數得到自己系統中的local memory數量：

cl_ulong DeviceLocalMemSize;
clGetDeviceInfo(device,
??? CL_DEVICE_LOCAL_MEM_SIZE,
??? sizeof(cl_ulong),
??? &DeviceLocalMemSize,
??? NULL);

lds的示意圖如下，對于每個bank，同時只能有一個讀寫請求，如果兩個thread都讀寫bank1，那個必須串行訪問，這就稱作bank conflict。

kernel初始化local memory的代碼如下：

//初始化共享內存
for(int i = 0; i < BIN_SIZE; ++i)
??? sharedArray[localId * BIN_SIZE + i] = 0;

???? 在同一時間，thread0訪問地址0(bank1),thread1,訪問地址256，也在bank1，…,這樣就有很多bank conflit，降低程序的性能。從profiler里面可以看到，lds bank conflit為13.98，很高的比例，所以此時同時運行的thread就比較少，只有總wave(每個wave 64個thread)的12%（我曾經默認lds內存分配是0，這樣我們就省去了這些代碼，但是實際上分配內存是一些隨機的值…)。

第二段memory操作的代碼為：

??? //計算thread直方圖
??? for(int i = 0; i < BIN_SIZE; ++i)
??? {
??????? uint value = (uint)data[groupId * groupSize * BIN_SIZE + i * groupSize + localId];
??????? sharedArray[localId * BIN_SIZE + value]++;
??? }

其中有lds的操作，也有global memory的操作，對于全局memory的訪問，在同一時刻，thread0訪問i=0的memory

，thread1訪問相鄰的memory單元…,這是對于global memory的訪問會采用合并讀寫的方式(coalencing)，就是一個memory請求返回16個dword，也就是一個請求滿足16個thread，提高memory利用率。此時對lds的寫是隨機的，根據value的值決定，不能控制…

?

最后一段memory讀寫的代碼：

//合并workgroup中所有線程的直方圖，產生workgroup直方圖
for(int i = 0; i < BIN_SIZE / groupSize; ++i)
{
???? uint binCount = 0;
???? for(int j = 0; j < groupSize; ++j)
???????? binCount += sharedArray[j * BIN_SIZE + i * groupSize + localId];
????????
???? binResult[groupId * BIN_SIZE + i * groupSize + localId] = binCount;
}

其中lds的讀寫如下圖，此時每個線程訪問不同的bank，因為amd lds訪問就是以32為單位，所以實際上，這段代碼不會有bank conflit。

本章學習一下在opencl中如何實現矩陣的轉置，主要的技巧還是利用好local memory，防止bank conflit以及使得全局memory的讀寫盡量是合并(coalensing)讀寫。

???? 我們的矩陣是一副二維灰度圖像256*256，矩陣的轉置也就是圖像的轉置。每個thread處理16(4*4)個pixel(uchar)，workgroup的size是(16,16)。

???? 下面直接看shader代碼：

uint wiWidth? = get_global_size(0);

uint gix_t = get_group_id(0);
uint giy_t = get_group_id(1);???
uint num_of_blocks_x = get_num_groups(0);

uint giy = gix_t;
uint gix = (gix_t+giy_t)%num_of_blocks_x;

uint lix = get_local_id(0);
uint liy = get_local_id(1);

uint blockSize = get_local_size(0);

uint ix = gix*blockSize + lix;
uint iy = giy*blockSize + liy;
int index_in = ix + (iy)*wiWidth*4;

// 通過合并讀寫把輸入數據裝入到lds中

int ind = liy*blockSize*4+lix;

block[ind]??????? = input[index_in];
block[ind+blockSize]??? = input[index_in+wiWidth];
block[ind+blockSize*2] = input[index_in+wiWidth*2];
block[ind+blockSize*3] = input[index_in+wiWidth*3];

???? 因為workgroup size是（16,16），所以lix,liy的取值范圍都是0-15，下面我們通過圖片看下，lix=0 liy=0,lix=1 liy=0時候，ind,以及index_in的值，從而得到輸入圖像數據如何映射到local memory中。

lix=0 liy=0

lix=1 liy=0

????? 下面是影射關系，(0,0) thread處理的16個pixel用血紅色表示，它們映射到lds的0bank和16bank，(1,0)thread處理的像素用綠色表示，它們映射到lds的bank1和bank17，有效的避免了bank conflit，而全局memory的訪問不同thread對應連續的全局memory空間，可以實現合并讀寫，從而提高程序性能。

? 把轉置的數據寫到全局memory中的代碼如下：

ix = giy*blockSize + lix;
iy = gix*blockSize + liy;
int index_out = ix + (iy)*wiWidth*4;

ind = lix*blockSize*4+liy;
uchar4 v0 = block[ind];
uchar4 v1 = block[ind+blockSize];
uchar4 v2 = block[ind+blockSize*2];
uchar4 v3 = block[ind+blockSize*3];

// 通過合并讀寫把lds中數據寫回到全局memory中
output[index_out]??????????? = (uchar4)(v0.x, v1.x, v2.x, v3.x);
output[index_out+wiWidth]??? = (uchar4)(v0.y, v1.y, v2.y, v3.y);
output[index_out+wiWidth*2]??? = (uchar4)(v0.z, v1.z, v2.z, v3.z);
output[index_out+wiWidth*3]??? = (uchar4)(v0.w, v1.w, v2.w, v3.w);
對應copy關系圖如下：

完整的代碼請參考：

工程文件gclTutorial9

代碼下載：

稍后提供

本篇教程中，我們學習一下如何用opencl有效實現數組求和，也就是通常所說的reduction問題。

???? 在程序中，我們設置workgroup size為256，kernel的輸入、輸出緩沖參數都用uint4的格式，這樣我們原始求和的數組大小為256*4的倍數，數據類型為uint。我們設定每個workgroup處理處理512個uint，即2048個uint

???? 為了簡便期間，我們輸出數組長度定為4096，即需要2個workgruop來處理。

???

kernel代碼如下：

__kernel void reduce(__global uint4* input, __global uint4* output, __local uint4* sdata)
{
??? // 把數據裝入lds
??? unsigned int tid = get_local_id(0);
??? unsigned int bid = get_group_id(0);
??? unsigned int gid = get_global_id(0);

??? unsigned int localSize = get_local_size(0);
??? unsigned int stride = gid * 2;
??? sdata[tid] = input[stride] + input[stride + 1];

??? barrier(CLK_LOCAL_MEM_FENCE);
??? // 在lds中進行reduction操作，得到數組求和的結果
??? for(unsigned int s = localSize >> 1; s > 0; s >>= 1)
??? {
??????? if(tid < s)
??????? {
??????????? sdata[tid] += sdata[tid + s];
??????? }
??????? barrier(CLK_LOCAL_MEM_FENCE);
??? }

?? // 把一個workgroup計算的結果輸出到輸出緩沖，是一個uint4，還需要在host端再進行一次reduction過程
??? if(tid == 0) output[bid] = sdata[0];
}

???? 在程序中，global和local的NDRange，我們都用一維的形式。下面以圖的方式看下kernel代碼是如何執行的：

??? ? 對第一個workgroup中的第一個thread的來說，它首先進行一次reduction操作，把兩個uint4相加，放到lds(shared memory)中，然后再在lds中進行reduction操作，此時要從global memory中取數據，可以看出連續的thread訪問連續的global memory，這時可以利用合并讀寫。

????? 申請的shared memory大小為groupsize*sizeof(uint4)，相加后uint4放入32bank的lds中，放置的方式應該是如下圖所示，因為放入的是uint4，所以會放入連續的4個bank中(每個bank都是dword寬)，可見只能同時有8個thread訪問lds，所以會有一定程序的bank conflit。從App profiler session，我們可以看到：

????? 接下來，kernel會通過一個for循環迭代執行reduction操作，求得一個workgroup中的uint4的和。

迭代的第一次s=128,這時會執行如下圖的兩兩相加，workgroup中同時執行的thread為128,thread local id大于等于128的線程都不會做什么事情，在每個循環的末尾，有一個barrier來同步所有thread，以便所有thread都完成這次循環后再進入下一次循環。

? ?? 第二次迭代的時候，只剩下前面128個uint4，workgroup中同時執行的thread為64。最后，當s=1時候，完成迭代reduction操作，然后把thread0（第一個thread)的結果輸出。

???? 在host段，我們還要做一次相加操作，把不同workgroup得到的uint4，拆分成uint，并相加求得最終的結果。

//在cpu reduction各個workgroup的結果以及uint4分量 reduction
output = 0;
for(int i = 0; i < numBlocks * VECTOR_SIZE; ++i)
??? output += outMapPtr[i];
printf("gpu reduction result:%d\n", output);
if(refOutput==output) printf("passed\n");

程序執行后結果如下：

完整的代碼請參考：

工程文件gclTutorial11

代碼下載：

稍后提供

?

總結

以上是生活随笔為你收集整理的AMD OpenCL 大学课程的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。