OpenCL作為一門原始碼的異構平行計算語言，設計之初就是使用一種模型來模糊各種硬體差異。作為軟體發展人員，我們關注的就是它的程式設計模型。OpenCL程式的流程大致如下：
下麵我們通過一個具體的範常式式來說明這些步驟。
使用 OpenCL API 程式設計與一般 C/C++ 引入協力廠商庫程式設計沒什麼區別。所以，首先要做的自然是 include 相關的標頭檔。由於在 MacOS X 10.6下OpenCL的標頭檔命名與其他系統不同，通常使用一個#if defined進行區分，程式碼如下：
#if defined(__APPLE__) || defined(__MACOSX)
#include <OpenCL/cl.hpp>
#else
#include <CL/cl.h>
#endif
接下來我們就進入真正的編碼流程了。
Platform
### 查詢並選擇一個 platform ### 首先我們要取得系統中所有的 OpenCL platform。所謂的 platform 指的就是硬體廠商提供的 OpenCL 框架，不同的 CPU/GPU 開發商（比如 Intel、AMD、Nvdia）可以在一個系統上分別定義自己的 OpenCL 框架。所以我們需要查詢系統中可用的 OpenCL 框架，即 platform。使用 API 函數 clGetPlatformIDs 獲取可用 platform 的數量：
cl_int status = 0;
cl_uint numPlatforms;
cl_platform_id platform = NULL;
status = clGetPlatformIDs( 0, NULL, &numPlatforms);
if(status != CL_SUCCESS){
printf("Error: Getting Platforms\n");
return EXIT_FAILURE;
}
然後根據數量來分配記憶體，並得到所有可用的 platform，所使用的 API 還是 clGetPlatformIDs。在 OpenCL 中，類似這樣的函式呼叫很常見：第一次呼叫以取得數目，便於分配足夠的記憶體；然後呼叫第二次以獲取真正的資訊。
if (numPlatforms > 0) {
cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");
return -1;
}
現在，所有的 platform 都存在了變數 platforms 中，接下來需要做的就是取得我們所需的 platform。本人的PC上配置的是 Intel 處理器和 AMD 顯卡，專業點的說法叫 Intel 的 CPU 和 AMD 的 GPU :)。所以我這兒有兩套 platform，為了體驗下 GPU 的快感，所以使用 AMD 的 platform。通過使用 clGetPlatformInfo 來獲得 platform 的資訊。通過這個 API 可以知曉 platform 的廠商資訊，以便我們選出需要的 platform。程式碼如下：
for (unsigned int i = 0; i < numPlatforms; ++i) {
char pbuff[100];
status = clGetPlatformInfo(
platforms[i],
CL_PLATFORM_VENDOR,
sizeof(pbuff),
pbuff,
NULL);
platform = platforms[i];
if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {
break;
}
}
不同的廠商資訊可以參考 OpenCL Specifications，我這兒只是簡單的篩選出 AMD 。
在 platform 上建立 context
第一步是通過 platform 得到相應的 context properties
// 如果我們能找到相應平臺，就使用它，否則返回NULL
cl_context_properties cps[3] = {
CL_CONTEXT_PLATFORM,
(cl_context_properties)platform,
0
};
cl_context_properties *cprops = (NULL == platform) ? NULL : cps;
第二步是通過 clCreateContextFromType 函數建立 context。
// 生成 context
cl_context context = clCreateContextFromType(
cprops,
CL_DEVICE_TYPE_GPU,
NULL,
NULL,
&status);
if (status != CL_SUCCESS) {
printf("Error: Creating Context.(clCreateContexFromType)\n");
return EXIT_FAILURE;
} 函數的第二個參數可以設定 context 關聯的設備類型。本例使用的是 GPU 作為OpenCL計算設備。目前可以使用的類別包括：
- CL_DEVICE_TYPE_CPU
- CL_DEVICE_TYPE_GPU
- CL_DEVICE_TYPE_ACCELERATOR
- CL_DEVICE_TYPE_DEFAULT
- CL_DEVICE_TYPE_ALL
在 context 上查詢 device
context 建立好之後，要做的就是查詢可用的 device。
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
0,
NULL,
&deviceListSize);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info device list size, clGetContextInfo)\n");
return EXIT_FAILURE;
}
cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);
if (devices == 0) {
printf("Error: No devices found.\n");
return EXIT_FAILURE;
}
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
deviceListSize,
devices,
NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info (device list, clGetContextInfo)\n");
return EXIT_FAILURE;
}
與獲取 platform 類似，我們呼叫兩次 clGetContextInfo 來完成 查詢。第一次呼叫獲取關聯 context 的 device 個數，並根據個數申請記憶體；第二次呼叫獲取所有 device 實例。如果你想瞭解每個 device 的具體資訊，可以呼叫 clGetDeviceInfo 函數來獲取，返回的資訊有設備類型、生產商以及設備對某些擴展功能的支持與否等等。詳細使用情況請參閱 OpenCL Specifications。
到此，platform 相關的程式已經準備就緒了，下麵到此的完整程式碼：
/* OpenCL_01.cpp
* (c) by keyring <keyrings@163.com>
* 2013.10.26
*/
#if defined(__APPLE__) || defined(__MACOSX)
#include <OpenCL/cl.hpp>
#else
#include <CL/cl.h>
#endif
#include <iostream>
int main(int argc, char const *argv[])
{
printf("hello OpenCL\n");
cl_int status = 0;
size_t deviceListSize;
// 得到並選擇可用平臺
cl_uint numPlatforms;
cl_platform_id platform = NULL;
status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS) {
printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");
return EXIT_FAILURE;
}
if (numPlatforms > 0) {
cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");
return -1;
}
// 搜尋所有 platform，選擇你想用的
for (unsigned int i = 0; i < numPlatforms; ++i) {
char pbuff[100];
status = clGetPlatformInfo(platforms[i],CL_PLATFORM_VENDOR,sizeof(pbuff),pbuff,
NULL);
platform = platforms[i];
if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {
break;
}
}
delete platforms;
}
// 如果我們能找到相應平臺，就使用它，否則返回NULL
cl_context_properties cps[3] = {
CL_CONTEXT_PLATFORM,
(cl_context_properties)platform,
0
};
cl_context_properties *cprops = (NULL == platform) ? NULL : cps;
// 生成 context
cl_context context = clCreateContextFromType(
cprops,
CL_DEVICE_TYPE_GPU,
NULL,
NULL,
&status);
if (status != CL_SUCCESS) {
printf("Error: Creating Context.(clCreateContexFromType)\n");
return EXIT_FAILURE;
}
// 尋找OpenCL設備
// 首先得到設備清單的長度
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
0,
NULL,
&deviceListSize);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info device list size, clGetContextInfo)\n");
return EXIT_FAILURE;
}
cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);
if (devices == 0) {
printf("Error: No devices found.\n");
return EXIT_FAILURE;
}
// 然後得到設備清單
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
deviceListSize,
devices,
NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info (device list, clGetContextInfo)\n");
return EXIT_FAILURE;
}
Running time
前面寫了一大篇，其實還沒真正進入具體的程式邏輯中，頂多算配好了 OpenCL 運行環境。真正的邏輯程式碼，即程式的任務就是運行時模組。本例的任務是在一個 4×4的二維空間上，按一定的規則給每個元素賦值，具體程式碼如下：
#define KERNEL(...)#__VA_ARGS__
const char *kernelSourceCode = KERNEL(
__kernel void hellocl(__global uint *buffer)
{
size_t gidx = get_global_id(0);
size_t gidy = get_global_id(1);
size_t lidx = get_local_id(0);
buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);
}
這一段就是我們真正的邏輯，也就是程式碼要幹的事。使用的是 OpenCL 自定的一門類C語言，具體的語法什麼的現在先不糾結。這段程式碼是直接嵌入我們的 cpp 檔的靜態字串。你也可以將 kernel 程式單獨寫成一個檔。
載入 OpenCL 核心程式並建立一個 program 物件
接下來要做的就是讀入 OpenCL kernel 程式並建立一個 program 物件。
size_t sourceSize[] = {strlen(kernelSourceCode)};
cl_program program = clCreateProgramWithSource(context,
1,
&kernelSourceCode,
sourceSize,
&status);
if (status != CL_SUCCESS) {
printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");
return EXIT_FAILURE;
}
本例中的 kernel 程式是作為靜態字串讀入的（單獨的文字檔也一樣），所以使用的是 clCreateProgramWithSource，如果你不想讓 kernel 程式讓其他人看見，可以先生成二進位檔案，再通過 clCreateProgramWithBinary 函數動態讀入二進位檔案，做一定的保密。詳細請參閱 OpenCL Specifications。
為指定的 device 編譯 program 中的 kernel
kernel 程式讀入完畢，要做的自然是使用 clBuildProgram 編譯 kernel：
status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Building Program (clBuildingProgram)\n");
return EXIT_FAILURE;
}
最終，kernel 將被相應 device 上的 OpenCL 編譯器編譯成可執行的機器碼。
建立指定名字的 kernel 對象
成功編譯後，可以通過 clCreateKernel 來建立一個 kernel 物件。
cl_kernel kernel = clCreateKernel(program, "hellocl", &status);
if (status != CL_SUCCESS) {
printf("Error: Creating Kernel from program.(clCreateKernel)\n");
return EXIT_FAILURE;
}
引號中的 hellocl 就是 kernel 物件所關聯的 kernel 函數的函數名。要注意的是，每個 kernel 物件必須關聯且只能關聯一個包含於相應 program 物件內的 kernel 程式。實際上，用戶可以在 cl 原始程式碼中寫任意多個 kernel 程式，但在執行某個 kernel 程式之前必須先建立單獨的 kernel 物件，即多次呼叫 clCreateKernel 函數。
為 kernel 建立記憶體物件
OpenCL 記憶體物件是指在 host 中建立，用於 kernel 程式的記憶體類型。按維度可以分為兩類，一類是 buffer，一類是 image。buffer 是一維的，image 可以是二維、三維的 texture、frame-buffer 或 image。本例僅僅使用 buffer，可以通過 clCreateBuffer 函數來建立。
cl_mem outputBuffer = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR, 4 * 4 * 4, NULL, &status);
if (status != CL_SUCCESS) {
printf("Error: Create Buffer, outputBuffer. (clCreateBuffer)\n");
return EXIT_FAILURE;
}
為 kernel 設置參數
使用 clSetKernelArg 函數為 kernel 設置參數。傳遞的參數既可以是常數，變數，也可以是記憶體物件。本例傳遞的就是記憶體物件。
status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&outputBuffer);
if (status != CL_SUCCESS) {
printf("Error: Setting kernel argument. (clSetKernelArg)\n");
return EXIT_FAILURE;
}
該函數每次只能設置一個參數，如有多個參數，需多次呼叫。而且 kernel 程式中所有的參數都必須被設置，否則在啟動 kernel 程式是會報錯。指定位置的參數的類型最好和對應 kernel 函數內參數類型一致，以免產生各種未知的錯誤。在設置好指定參數後，每次運行該 kernel 程式都會使用設置值，直到用戶使用次 API 重新設置參數。
在指定的 device 上建立 command queue
command queue 用於光裡將要執行的各種命令。可以通過 clCreateCommandQueue函數建立。其中的 device 必須為 context 的關聯設備，所有該 command queue 中的命令都會在這個指定的 device 上運行。
cl_command_queue commandQueue = clCreateCommandQueue(context,
devices[0],
0,
&status);
if (status != CL_SUCCESS) {
printf("Error: Create Command Queue. (clCreateCommandQueue)\n");
return EXIT_FAILURE;
}
將要執行的 kernel 放入 command queue
建立好 command queue 後，用戶可以建立相應的命令並放入 command queue 中執行。OpenCL 提供了三種方案來建立 kernel 執行命令。最常用的即為本例所示的運行在指定工作空間上的 kernel 程式，使用了 clEnqueueNDRangeKernel 函數。
size_t globalThreads[] = {4, 4};
size_t localThreads[] = {2, 2};
status = clEnqueueNDRangeKernel(commandQueue, kernel,
2, NULL, globalThreads,
localThreads, 0,
NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Enqueueing kernel\n");
return EXIT_FAILURE;
}
clEnqueueNDRangeKernel 函數每次只能將一個 kernel 物件放入 command queue 中，用戶可以多次呼叫該 API 將多個 kernel 物件放置到一個 command queue 中，command queue 中的不同 kernel 物件的工作區域完全不相關。其餘兩個 API clEnqueueTask 和 clEnqueueNativeKernel 的用法就不多講了，詳情請參閱 OpenCL Specificarions。
最後可以用 clFinish 函數來確認一個 command queue 中所有的命令都執行完畢。函數會在 command queue 中所有 kernel 執行完畢後返回。
// 確認 command queue 中所有命令都執行完畢
status = clFinish(commandQueue);
if (status != CL_SUCCESS) {
printf("Error: Finish command queue\n");
return EXIT_FAILURE;
}
將結果讀回 host
計算完畢，將結果讀回 host 端。使用 clEnqueueReadBuffer 函數將 OpenCL buffer 物件中的內容讀取到 host 可以訪問的記憶體空間。
// 將記憶體物件中的結果讀回Host
status = clEnqueueReadBuffer(commandQueue,
outputBuffer, CL_TRUE, 0,
4 * 4 * 4, outbuffer, 0, NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Read buffer queue\n");
return EXIT_FAILURE;
}
當然，為了看下程式的運行效果，咱們當然得看看運行結果啦。列印一下吧：
// Host端列印結果
printf("out:\n");
for (int i = 0; i < 16; ++i) {
printf("%x ", outbuffer[i]);
if ((i + 1) % 4 == 0)
printf("\n");
}
資源回收
程式的最後是對所有建立的物件進行釋放回收，與C/C++的記憶體回收同理。
// 資源回收
status = clReleaseKernel(kernel);
status = clReleaseProgram(program);
status = clReleaseMemObject(outputBuffer);
status = clReleaseCommandQueue(commandQueue);
status = clReleaseContext(context);
free(devices);
delete outbuffer;
總結
這次使用一個小例子來詳細說明瞭 OpenCL 程式設計的一般步驟。其實這些步驟一般都是固定的。真正需要我們注意的是 OpenCL Kernel 程式的編寫。當然，合理高效的利用 API 也是一門技術活。
最後給出本實例的全部程式碼：
/* OpenCL_01.cpp
* (c) by keyring <keyrings@163.com>
* 2013.10.26
*/
#include <iostream>
#if defined(__APPLE__) || defined(__MACOSX)
#include <OpenCL/cl.hpp>
#else
#include <CL/cl.h>
#endif
#define KERNEL(...)#__VA_ARGS__
const char *kernelSourceCode = KERNEL(
__kernel void hellocl(__global uint *buffer)
{
size_t gidx = get_global_id(0);
size_t gidy = get_global_id(1);
size_t lidx = get_local_id(0);
buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);
}
);
int main(int argc, char const *argv[])
{
printf("hello OpenCL\n");
cl_int status = 0;
size_t deviceListSize;
// 得到並選擇可用平臺
cl_uint numPlatforms;
cl_platform_id platform = NULL;
status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS) {
printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");
return EXIT_FAILURE;
}
if (numPlatforms > 0) {
cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");
return -1;
}
for (unsigned int i = 0; i < numPlatforms; ++i) {
char pbuff[100];
status = clGetPlatformInfo(
platforms[i],
CL_PLATFORM_VENDOR,
sizeof(pbuff),
pbuff,
NULL);
platform = platforms[i];
if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {
break;
}
}
delete platforms;
}
// 如果我們能找到相應平臺，就使用它，否則返回NULL
cl_context_properties cps[3] = {
CL_CONTEXT_PLATFORM,
(cl_context_properties)platform,
0
};
cl_context_properties *cprops = (NULL == platform) ? NULL : cps;
// 生成 context
cl_context context = clCreateContextFromType(
cprops,
CL_DEVICE_TYPE_GPU,
NULL,
NULL,
&status);
if (status != CL_SUCCESS) {
printf("Error: Creating Context.(clCreateContexFromType)\n");
return EXIT_FAILURE;
}
// 尋找OpenCL設備
// 首先得到設備清單的長度
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
0,
NULL,
&deviceListSize);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info device list size, clGetContextInfo)\n");
return EXIT_FAILURE;
}
cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);
if (devices == 0) {
printf("Error: No devices found.\n");
return EXIT_FAILURE;
}
// 現在得到設備清單
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
deviceListSize,
devices,
NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info (device list, clGetContextInfo)\n");
return EXIT_FAILURE;
}
// 裝載核心程式，編譯CL program ,生成CL核心實例
size_t sourceSize[] = {strlen(kernelSourceCode)};
cl_program program = clCreateProgramWithSource(context,
1,
&kernelSourceCode,
sourceSize,
&status);
if (status != CL_SUCCESS) {
printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");
return EXIT_FAILURE;
}
// 為指定的設備編譯CL program.
status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Building Program (clBuildingProgram)\n");
return EXIT_FAILURE;
}
// 得到指定名字的核心實例的控制碼
cl_kernel kernel = clCreateKernel(program, "hellocl", &status);
if (status != CL_SUCCESS) {
printf("Error: Creating Kernel from program.(clCreateKernel)\n");
return EXIT_FAILURE;
}
// 建立 OpenCL buffer 對象
unsigned int *outbuffer = new unsigned int [4 * 4];
memset(outbuffer, 0, 4 * 4 * 4);
cl_mem outputBuffer = clCreateBuffer(
context,
CL_MEM_ALLOC_HOST_PTR,
4 * 4 * 4,
NULL,
&status);
if (status != CL_SUCCESS) {
printf("Error: Create Buffer, outputBuffer. (clCreateBuffer)\n");
return EXIT_FAILURE;
}
// 為核心程式設置參數
status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&outputBuffer);
if (status != CL_SUCCESS) {
printf("Error: Setting kernel argument. (clSetKernelArg)\n");
return EXIT_FAILURE;
}
// 建立一個OpenCL command queue
cl_command_queue commandQueue = clCreateCommandQueue(context,
devices[0],
0,
&status);
if (status != CL_SUCCESS) {
printf("Error: Create Command Queue. (clCreateCommandQueue)\n");
return EXIT_FAILURE;
}
// 將一個kernel 放入 command queue
size_t globalThreads[] = {4, 4};
size_t localThreads[] = {2, 2};
status = clEnqueueNDRangeKernel(commandQueue, kernel,
2, NULL, globalThreads,
localThreads, 0,
NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Enqueueing kernel\n");
return EXIT_FAILURE;
}
// 確認 command queue 中所有命令都執行完畢
status = clFinish(commandQueue);
if (status != CL_SUCCESS) {
printf("Error: Finish command queue\n");
return EXIT_FAILURE;
}
// 將記憶體物件中的結果讀回Host
status = clEnqueueReadBuffer(commandQueue,
outputBuffer, CL_TRUE, 0,
4 * 4 * 4, outbuffer, 0, NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Read buffer queue\n");
return EXIT_FAILURE;
}
// Host端列印結果
printf("out:\n");
for (int i = 0; i < 16; ++i) {
printf("%x ", outbuffer[i]);
if ((i + 1) % 4 == 0)
printf("\n");
}
// 資源回收
status = clReleaseKernel(kernel);
status = clReleaseProgram(program);
status = clReleaseMemObject(outputBuffer);
status = clReleaseCommandQueue(commandQueue);
status = clReleaseContext(context);
free(devices);
delete outbuffer;
system("pause");
return 0;
}
platform
OpenCL 全稱 Open Computing Language，作為開放語言自然有其開放的特質。其為異構平臺提供了一個編寫程式，尤其是並行程式的開放式框架標準。它支援的異構平臺可由多核 CPU、GPU 或其他類型的處理器（DSP）組成。
這個框架標準具體來說就是 platform。OpenCL 使用的是 Installable Client Driver,ICD 模型，意味著一個系統上可以有多個 OpenCL 實現並存。簡單地說，你的系統上有多少個 OpenCL 框架/實現/ platform，取決於你有什麼硬體和與之配對的驅動程式。比如，我的系統是 Windows+Intel處理器+AMD顯卡，所以裝好顯卡驅動後，我的PC上就有兩個 OpenCL 框架/實現/platform，如果你再裝個 NVIDIA 顯卡和驅動，那就又多了一個 platform。例外的是蘋果電腦，由於它的封閉統一性，便只有一個 OpenCL platform。
到此，我們可以明白，platfrom == implement，你的系統上有多少個 OpenCL 實現，就有多少個 platform。一般具體的實現都在驅動程式裡面，所以千萬記得裝好對應的驅動程式！
我們可以通過 clGetPlatformIDs 來查詢 platform 集。
cl_int clGetPlatformIDs( cl_uint num_entries,
cl_platform_id *platforms,
cl_uint *num_platforms )
這類查詢函數在 OpenCL 裡用法相似。兩次呼叫，第一次獲取結果的個數/大小以便申請記憶體空間，第二次呼叫來獲取具體的結果集。
舉個例子來說明一下：
cl_int status; // 獲取函數返回值，一般是狀態碼
cl_uint numPlatforms; // platfrom 數量
cl_platform_id *platormIds; // 存儲所有 platform 的空間
// 第一次呼叫，獲取 platform 個數
status = clGetPlatformIDs( 0, NULL, &numPlatforms );
// 利用獲取的個數來申請記憶體空間
platformIds = (cl_platform_id *)alloca(sizeof(cl_platform_id) * numPlatforms);
// 第二次呼叫，獲取 platform ID 列表
status = clGetPlatformIDs( numPlatforms, platformIds, NULL );
注意：對於蘋果系統，這一步沒有必要，因為 MAC 上只有一個 platform。
在已知了 platform 之後，可以通過 clGetPlatformInfo 函數來獲取 platform 的各個屬性。
cl_int clGetPlatformInfo( cl_platform_id platform,
cl_platform_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret )
其中 param_name 的取值如下：
-CL_PLATFORM_PROFILE // OpenCL 簡檔
-CL_PLATFORM_VERSION // OpenCL 版本
-CL_PLATFORM_NAME // platform 名字
-CL_PLATFORM_VENDOR // platform 開發商
-CL_PLATFORM_EXTENSIONS // platform 支持的擴展列表
// 它們的返回值都是 字串。
該 API 的具體使用請參考 OpenCL Specification。
device
如果說 platform 只是一個框架實現，看得見摸不著。那麼 device 就是一個非常具體的存在了。簡單地說，device 就是硬體設備，比如 CPU,GPU。你的應用程式將利用這些計算設備來執行程式碼。
各個 platform 可能會分別關聯一組計算設備。在給定的 platform 上，可以使用 clGetDeviceIDs 來獲取與之關聯的受支援的設備清單。
cl_int clGetDeviceIDs( cl_platform_id platform,
cl_device_type device_type,
cl_uint num_entries,
cl_device_id *devices,
cl_uint *num_devices )
其中 device_type 代表你獲取的設備的類型，其取值如下：
-CL_DEVICE_TYPE_CPU // host 端的 CPU 設備
-CL_DEVICE_TYPE_GPU // GPU 設備
-CL_DEVICE_TYPE_ACCELERATOR // OpenCL 加速器
-CL_DEVICE_TYPE_DEFAULT // 預設設備
-CL_DEVICE_TYPE_ALL // 與相應 platform 關聯的所有設備
CPU是一個同構設備，通常可以利用大緩存進行優化來減少延遲。比如 AMD 的 Opteron 系列和 Intel 的 Core 系列。
GPU是面向圖形和通用計算的大輸送量設備。比如 AMD 的 Radeon 系列和 NVIDIA 的 GTX 系列。
ACCELERATOR涵蓋了從 IBM 的 Cell Broadband 體系結構到不太著名的 DSP 型等大量的設備。
DEFAULT允許 OpenCL Running time library 指定一個“首選”設備。
ALL允許 OpenCL Running time library 指定所有可用設備。
對於 CPU,GPU 和 ACCELERATOR 設備，對 platform 提供的設備數量沒有任何限制，由應用程式負責查詢來確定具體的數目。當然，使用的 API 還是 clGetDeviceIDs。
對於獲取的某個 device，可以利用 clGetDeviceInfo 函數查詢其各項屬性：
cl_int clGetDeviceInfo( cl_device_id device,
cl_device_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret )
這個函數能夠返回的資訊非常多,大概好幾十種吧，從設備類型到設備的記憶體狀態均能獲得。具體的請查閱 OpenCL Specificaion。至於函數的使用，咱們以後在具體範例中來展現。
context
context 是所有 OpenCL 應用的核心。它為與之關聯的 device,memory,command queue 提供一個容器。如果說 platform 等於 實現，device 就是 設備，那麼 context 就是一個 管理員。它管理著應用程式與設備，設備與設備之間的通信。
正如 platform,device 有很多個一樣，context 也可以有很多個。不同之處在於，前面兩個的個數我們沒法控制，是硬體/軟體提供商提供的，而 context 是由我們自己決定的。
我們可以由不同的 platform 建立多個 context，並把工作分佈到這些 context和與之關聯的 device 上。下面這幅圖可以直觀的展示它們之間的關係，簡單的說，都是 一對多 的關係。
context 不再由別人提供，需要我們自己來建立。有兩種方式來建立：已知一個 platform 和一組關聯的 device，可以使用 clCreateContext 函數來建立；已知一個 platform 和 device 的類型，可以使用 clCreateContextFromType 函數來建立。
cl_context clCreateContext( const cl_context_properties *properties,
cl_uint num_devices,
const cl_device_id *devices,
void (CL_CALLBACK *pfn_notify)( const char *errorinfo,
void *user_data,
cl_int *errcode_ret )
cl_context clCreateContextFromType( const cl_context_properties *properties,
cl_device_type device_type,
void (CALLBACK *pfn_notify) (const char *errinfo,
void *user_data,
cl_int *errcode_ret )
參數 pfn_notify 和 user_data 用來共同定義一個回檔，用於報告 context 生命期中所出現錯誤的有關資訊，記得把 user_data 最為最後一個參數傳給回呼函數。
同樣的，我們也可以通過 API 來查詢 context 的屬性，clGetContextInfo 函數的具體用法將在以後的範例程式裡展現。
cl_int clGetContextInfo( cl_context context,
cl_context_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret )
所有的 OpenCL 物件都是引用計數的，對於 context 這類有我們自己建立的物件，需要我們自己來遞增和遞減引用數。
cl_int clRetainContext (cl_context context) // 遞增引用數
cl_int clReleaseContext(cl_context context) // 遞減引用數
program
program 的意思相信大家都懂的，畢竟我們自己就是一名苦逼的 programmer。顧名思義，program 代表的是一個程式物件，裡面包含著我們所寫的 cl 程式。
做 OpenCL 程式設計，要寫兩種程式碼，一種是 host 程式，一種是 cl 程式，類比於圖形程式設計，cl 程式就等同於 shader 程式。host 端的程式碼幹的是為 cl 程式設置環境，打掃記憶體等保姆型工作， cl 程式才是真正 解決問題 的地方。
要讓 cl 程式發揮作用，得先讓 host 端把它送進 device 是吧。送之前得先載入進 host 吧。載入後存哪呢？存在 program 裡面！！
建立一個 program 物件有兩種方法，一種是利用 clCreateProgramWithSource函數從 cl 原始程式碼建立，另一種是利用 clCreateProgramWithBinary 函數從二進位中建立。兩個函數的原型如下：
cl_program clCreateProgramWithSource( cl_context context,
cl_uint count,
const char **strings,
const size_t *lengths,
cl_int *errcode_ret )
cl_program clCreateProgramWithBinary( cl_context context,
cl_uint num_devices,
const cl_device_id *device_list,
const size_t *lengths,
cosnt unsigned char **binaries,
cl_int *binary_status,
cl_int *errcode_ret )
關於兩種建立方式實際使用，我後面會寫個小範例。當然，OpenCL Specification 一直都是最佳最標準的參考讀物。
一般建立好一個 program 物件後，接著就是對其進行編譯。從上面兩個建立函數可以看出，現在 cl 程式存在 program 物件中，其形式莫過於原始程式碼和二進位。而我們的 device （無論是CPU、GPU還是DSP）能處理的僅僅是機器碼和指令。所以與一般的程式設計相同，我們需要將 cl 程式進行編譯。不同之處在於編譯方法不在是點一下IDE上的按鈕或者make一下，而是在 host 程式內部以 API 呼叫的方式來控制 編譯。控制編譯的 API 是 clBuildProgram，其原型如下：
cl_int clBuildProgram(cl_program program,
cl_uint num,
const cl_device_id *device_list,
const char *options,
void(CL_CALLBACK *pfn_notify)(cl_program program, void *user_data),
void *user_data )
這個函數會把裝有 cl 程式的 program 物件傳入指定的 device（第三個參數）進行編譯。而且，你也可以設定多種編譯 options（第四個參數）。options 包括前置處理器定義和各種優化以及程式碼生成（比如，-DUSE_FEATURE = 1 -cl -mad -enable）。
編譯完成並最終生成的可執行程式碼還是存在了指定 device（第三個參數指定的）的 program 中。如果在所有指定 device 上編譯成功，則該函數返回 CL_SUCCESS；否則會返回一個錯誤碼。如果存在錯誤，可以呼叫 clGetProgramBuildInfo，並指定 param_name（第三個參數） 為 CL_PROGRAM_BUILD_LOG來查看某個 device（第二個參數）上的詳細構建日誌。
cl_int clGetProgramBuildInfo( cl_program program,
cl_device_id device,
cl_program_buld_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret )
這類由我們自己 create 的物件，可以手動控制其引用計數，物件使用完了一定要記得 release。不然記憶體洩漏是很慘的。
cl_int clRetainProgram( cl_program program ) // 遞減引用計數
cl_int clReleaseProgram(cl_program program ) // 遞增引用計數
kernel
如果說上述 program 物件是一個容器，裝著一個 cl 程式，那麼一個 kernel 物件也是一個容器，裝的是 cl 程式中的一個 kernel 函數。
kernel 函數指的是 cl 程式中那些以 kernel 或 __kernel 限定的函數。一段 cl 程式碼可以包含多個 kernel 函數，自然就對應著多個 kernel` 物件。
kernel 物件有什麼用呢？想想我們要執行某個 kernel 函數，要傳遞一堆的參數，在平時的C程式設計裡面，直接呼叫函數即可是吧。可問題在於現在 host 端與 cl端是分離的，不光檔分離，甚至連執行的地方都是分離的（比如，host 在CPU上，cl 在GPU上）。這種情況下怎麼向 kernel 函數 傳遞參數呢？
我們的 kernel 物件的用處就在於此：參數傳遞（資料傳遞）。我們一直說 OpenCL 提供了一個異構計算的平臺，這些各種物件其實就是對異構程式設計的一個抽象。
瞭解了 what 和 why，接下來就是 how了。
建立 kernel 物件的方法是將 kernel 函數名（第二個參數） 傳入 clCreateKernel：
cl_kernel clCreateKernel( cl_program program,
const char *kernel_name,
cl_int *errcode_ret )
kernel_name 指的就是 cl 程式碼裡面 kernel 或 __kernel 關鍵字後面的函數名。
clCreateKernel 一次只能為一個 kernel 函數 建立一個 kernel 物件。如果你的 cl 程式碼裡面有很多很多 kernel 函數 怎麼辦？一個個來？那還不煩死！所以這種情況請使用 clCreateKernelsInProgram 為 program 中的所有 kernel 函數 建立物件。
cl_int clCreateKernelsInProgram(cl_program program,
cl_uint num_kernels,
cl_kernel *kernels,
cl_uint *num_kernels_ret)
使用 clCreateKernelsInProgram 時需要呼叫兩次：第一次確定 kernel 數量；然後申請空間；第二次具體建立 kernel 對象。比如下面的範例程式碼：
cl_uint numKernels;
clCreateKernelsInProgram( program, 0, NULL, &numKernels ); // 第一次呼叫確定數量
cl_kernel *kernels = new cl_kernel[numKernels]; // 根據數量申請空間
clCreateKernelsInProgram( program, numKernels, kernels, &numKernels ); // 具體建立對象
建立好之後，就可以為相應的 kernel 函數 傳遞參數了。使用的 API 是 clSetKernelArg
cl_int clSetKernelArg( cl_kernel kernel,
cl_uint arg_index,
size_t *arg_size,
const void *arg_value )
其中，arg_index 是從 0 開始的，即第一個參數索引為 0，第二個為 1，依次類推。
最後當然得記得增減計數器：
cl_int RetainKernel( cl_kernel kernel )
cl_int ReleaseKernel(cl_kernel kernel )
注意：有時我們會考慮對一個 program 物件重新編譯。記住，重新編譯前一定要把與該 program 物件相關的 kernel 物件全部釋放掉，否則重新編譯會出錯！！！
存儲 opencl 程式為二進位
我們的第一個問題是：二進位版的 opencl 程式從哪裡來？前文說過，所有的 cl 程式碼都要經過載入並建立 program 物件，然後由 program 物件在 device 上面編譯並執行。難道還有其他方式編譯 opencl 程式碼？答案是：NO!
意味著我們還是需要將程式碼送到 device 裡面編譯。你會說，這不是多此一舉嗎？看起來確實有點，不過一般的做法是在軟體安裝的時候就進行編譯保存二進位形式，然後真正運行時才載入二進位。這樣分成兩個步驟的話，倒也說的過去。
省去前面那些與 platform、device 和 context的程式碼，我們直接進入建立 program 的地方。首先還是利用 clCreateProgramWithSource 函數讀取原始程式碼檔並用 clBuildProgram 函數編譯。範例程式碼如下：
cl_int status;
cl_program program;
ifstream kernelFile("binary_kernel.ocl", ios::in);
if(!kernelFile.is_open())
return;
ostringstream oss;
oss << kernelFile.rdbuf();
string srcStdStr = oss.str();
const char *srcStr = srcStdStr.c_str();
program = clCreateProgramWithSource(context, 1, (const char **)&srcStr, NULL, NULL);
if(program ==NULL)
return;
status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if(status != CL_SUCCESS)
return;
程式碼可能不完整，完整範例請看文末。
現在我們已經將 opencl 程式碼在 device 編譯完成了。接下來要做的就是將編譯好的二進位取出來存在磁片上。使用的 API 就是 clGetProgramInfo:
cl_int clGetProgramInfo(cl_program program,
cl_program_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
使用方法見如下程式碼片段（為使邏輯清晰，省略了錯誤檢測，實際開發可不要省啊）：
cl_uint numDevices = 0;
// 獲取 program 綁定過的 device 數量
clGetProgramInfo(program,
CL_PROGRAM_NUM_DEVICES,
sizeof(cl_uint),
&numDevices,
NULL);
// 獲取所有的 device ID
cl_device_id *devices = new cl_device_id[numDevices];
clGetProgramInfo(program,
CL_PROGRAM_DEVICES,
sizeof(cl_device_id) * numDevices,
devices,
NULL);
// 決定每個 program 二進位的大小
size_t *programBinarySizes = new size_t[numDevices];
clGetProgramInfo(program,
CL_PROGRAM_BINARY_SIZES,
sizeof(size_t) * numDevices,
programBinarySizes,
NULL);
unsigned char **programBinaries = new unsigned char *[numDevices];
for(cl_uint i = 0; i < numDevices; ++i)
programBinaries[i] = new unsigned char[programBinarySizes[i]];
// 獲取所有的 program 二進位
clGetProgramInfo(program,
CL_PROGRAM_BINARIES,
sizeof(unsigned char *) * numDevices,
programBinaries,
NULL);
// 存儲 device 所需要的二進位
for(cl_uint i = 0; i < numDevices; ++i){
// 只存儲 device 需要的二進位，多個 device 需要存儲多個二進位
if(devices[i] == device){
FILE *fp = fopen("kernel_binary_ocl.bin", "wb");
fwrite(programBinaries[i], 1， programBinarySizes[i], fp);
fclose(fp);
break;
}
}
// 清理
delete[] devices;
delete [] programBinarySizes;
for(cl_uint i = 0; i < numDevices; ++i)
delete [] programBinaries[i];
delete[] programBinaries;
要注意的是，可能有很多個 device 都編譯了 program，所以將二進位提取出來時，我們是搜尋了所有編譯了 program 的 device。
讀取二進位版opencl程式
經過上面一系列的操作，我們的磁片上應該存在一個二進位版的 opencl 程式了。裡面的內容可能是可讀的，也可能是不可讀的。這個視不同廠商實現而不同。
相對於存儲，讀取看起來就清爽的多，無非是打開二進位檔案，然後呼叫clCreateProgramWithBinary函數。範例如下：
FILE *fp= fopen("kernel_binary_ocl.bin", "rb");
// 獲取二進位的大小
size_t binarySize;
fseek(fp, 0, SEEK_END);
binarySize = ftell(fp);
rewind(fp);
// 載入二進位檔案
unsigned char *programBinary = new unsigned char[binarySize];
fread(programBinary, 1, binarySize, fp);
fclose(fp);
cl_program program;
program = clCreateProgramWithBinary(context,
1, &device, &binarySize,
(const unsigned char**)&programBinary,
NULL，
NULL）；
delete [] programBinary;
clBuildProgram(program 0, NULL, NULL, NULL, NULL);
這裡要注意，即使載入是二進位版，我們在之後還是要對其進行 clBuildProgram。原因在於，我們無法保證所謂的二進位一定是可執行碼。因為每個廠商的實現不一，有的可能就是最終執行碼，而有的卻是中間碼。所以無論是從原始程式碼還是二進位建立 program，之後都需要 clBuildProgram。
這樣兜了一圈，發現要使用二進位版還是用了一遍原始程式碼方式，感覺程式碼複雜好多，有點多餘。其實換個角度來看，我們完全可以寫成兩個程式，一個專門用來讀取原始程式碼並編譯生成二進位，另一個才是讀取二進位運行軟體。前者開發人員使用，後者才是給用戶使用的。只有這樣才能體現二進位版的優勢。
範例程式碼
OpenCLCompileToBin.cpp
/* OpenCLCompileToBin.cpp
* (c) by keyring <keyrings@163.com>
* 2016.01.13
*/
#include <iostream>
#if defined(__APPLE__) || defined(__MACOSX)
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
#define KERNEL(...)#__VA_ARGS__
const char *kernelSourceCode = KERNEL(
__kernel void hellocl(__global uint *buffer)
{
size_t gidx = get_global_id(0);
size_t gidy = get_global_id(1);
size_t lidx = get_local_id(0);
buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);
}
);
int main(int argc, char const *argv[])
{
printf("hello OpenCL\n");
cl_int status = 0;
size_t deviceListSize;
// 得到並選擇可用平臺
cl_uint numPlatforms;
cl_platform_id platform = NULL;
status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS) {
printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");
return EXIT_FAILURE;
}
if (numPlatforms > 0) {
cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");
return -1;
}
for (unsigned int i = 0; i < numPlatforms; ++i) {
char pbuff[100];
status = clGetPlatformInfo(
platforms[i],
CL_PLATFORM_VENDOR,
sizeof(pbuff),
pbuff,
NULL);
platform = platforms[i];
if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {
break;
}
}
delete platforms;
}
// 如果我們能找到相應平臺，就使用它，否則返回NULL
cl_context_properties cps[3] = {
CL_CONTEXT_PLATFORM,
(cl_context_properties)platform,
0
};
cl_context_properties *cprops = (NULL == platform) ? NULL : cps;
// 生成 context
cl_context context = clCreateContextFromType(
cprops,
CL_DEVICE_TYPE_GPU,
NULL,
NULL,
&status);
if (status != CL_SUCCESS) {
printf("Error: Creating Context.(clCreateContexFromType)\n");
return EXIT_FAILURE;
}
// 尋找OpenCL設備
// 首先得到設備清單的長度
status = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceListSize);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info device list size, clGetContextInfo)\n");
return EXIT_FAILURE;
}
cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);
if (devices == 0) {
printf("Error: No devices found.\n");
return EXIT_FAILURE;
}
// 現在得到設備清單
status = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceListSize, devices, NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info (device list, clGetContextInfo)\n");
return EXIT_FAILURE;
}
// 裝載核心程式，編譯CL program ,生成CL核心實例
size_t sourceSize[] = {strlen(kernelSourceCode)};
cl_program program = clCreateProgramWithSource(context, 1, &kernelSourceCode, sourceSize,
&status);
if (status != CL_SUCCESS) {
printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");
return EXIT_FAILURE;
}
// 為指定的設備編譯CL program.
status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Building Program (clBuildingProgram)\n");
return EXIT_FAILURE;
}
deviceListSize = 1;
// 決定每個 program 二進位的大小
size_t *programBinarySizes = new size_t[deviceListSize];
clGetProgramInfo(program,
CL_PROGRAM_BINARY_SIZES,
sizeof(size_t) * deviceListSize,
programBinarySizes,
NULL);
printf("%lu\n", deviceListSize);
unsigned char **programBinaries = new unsigned char *[deviceListSize];
for(cl_uint i = 0; i < deviceListSize; ++i)
programBinaries[i] = new unsigned char[programBinarySizes[i]];
// 獲取所有的 program 二進位
clGetProgramInfo(program,
CL_PROGRAM_BINARIES,
sizeof(unsigned char *) * deviceListSize,
programBinaries,
NULL);
printf("ready write to file\n");
// 寫入文件
FILE *fp = fopen("./kernel_binary_ocl.bin", "wb");
fwrite(programBinaries[0], 1, programBinarySizes[0], fp);
fclose(fp);
// 資源回收
status = clReleaseProgram(program);
status = clReleaseContext(context);
free(devices);
delete [] programBinarySizes;
for(cl_uint i = 0; i < deviceListSize; ++i)
delete [] programBinaries[i];
delete programBinaries;
return 0;
}
OpenCLRunWithBin.cpp
/* OpenCLRunWithBin.cpp
* (c) by keyring <keyrings@163.com>
* 2016.01.13
*/
#include <iostream>
#if defined(__APPLE__) || defined(__MACOSX)
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
// #define KERNEL(...)#__VA_ARGS__
// const char *kernelSourceCode = KERNEL(
// __kernel void hellocl(__global uint *buffer)
// {
// size_t gidx = get_global_id(0);
// size_t gidy = get_global_id(1);
// size_t lidx = get_local_id(0);
// buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);
// }
// );
int main(int argc, char const *argv[])
{
printf("hello OpenCL\n");
cl_int status = 0;
size_t deviceListSize;
// 得到並選擇可用平臺
cl_uint numPlatforms;
cl_platform_id platform = NULL;
status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS) {
printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");
return EXIT_FAILURE;
}
if (numPlatforms > 0) {
cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");
return -1;
}
for (unsigned int i = 0; i < numPlatforms; ++i) {
char pbuff[100];
status = clGetPlatformInfo(
platforms[i],
CL_PLATFORM_VENDOR,
sizeof(pbuff),
pbuff,
NULL);
platform = platforms[i];
if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {
break;
}
}
delete platforms;
}
// 如果我們能找到相應平臺，就使用它，否則返回NULL
cl_context_properties cps[3] = {
CL_CONTEXT_PLATFORM,
(cl_context_properties)platform,
0
};
cl_context_properties *cprops = (NULL == platform) ? NULL : cps;
// 生成 context
cl_context context = clCreateContextFromType(
cprops,
CL_DEVICE_TYPE_GPU,
NULL,
NULL,
&status);
if (status != CL_SUCCESS) {
printf("Error: Creating Context.(clCreateContexFromType)\n");
return EXIT_FAILURE;
}
// 尋找OpenCL設備
// 首先得到設備清單的長度
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
0,
NULL,
&deviceListSize);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info device list size, clGetContextInfo)\n");
return EXIT_FAILURE;
}
cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);
if (devices == 0) {
printf("Error: No devices found.\n");
return EXIT_FAILURE;
}
// 現在得到設備清單
status = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
deviceListSize,
devices,
NULL);
if (status != CL_SUCCESS) {
printf("Error: Getting Context Info (device list, clGetContextInfo)\n");
return EXIT_FAILURE;
}
FILE *fp= fopen("./kernel_binary_ocl.bin", "rb");
// 獲取二進位的大小
size_t binarySize;
fseek(fp, 0, SEEK_END);
binarySize = ftell(fp);
rewind(fp);
// 載入二進位檔案
unsigned char *programBinary = new unsigned char[binarySize];
fread(programBinary, 1, binarySize, fp);
fclose(fp);
cl_program program;
program = clCreateProgramWithBinary(context,
1,
&devices[0],
&binarySize,
(const unsigned char**)&programBinary,
NULL,
NULL);
delete [] programBinary;
// 裝載核心程式，編譯CL program ,生成CL核心實例
// size_t sourceSize[] = {strlen(kernelSourceCode)};
// cl_program program = clCreateProgramWithSource(context,
// 1,
// &kernelSourceCode,
// sourceSize,
// &status);
// if (status != CL_SUCCESS) {
// printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");
// return EXIT_FAILURE;
// }
// 為指定的設備編譯CL program.
status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Building Program (clBuildingProgram)\n");
return EXIT_FAILURE;
}
// 得到指定名字的核心實例的控制碼
cl_kernel kernel = clCreateKernel(program, "hellocl", &status);
if (status != CL_SUCCESS) {
printf("Error: Creating Kernel from program.(clCreateKernel)\n");
return EXIT_FAILURE;
}
// 建立 OpenCL buffer 對象
unsigned int *outbuffer = new unsigned int [4 * 4];
memset(outbuffer, 0, 4 * 4 * 4);
cl_mem outputBuffer = clCreateBuffer(
context,
CL_MEM_ALLOC_HOST_PTR,
4 * 4 * 4,
NULL,
&status);
if (status != CL_SUCCESS) {
printf("Error: Create Buffer, outputBuffer. (clCreateBuffer)\n");
return EXIT_FAILURE;
}
// 為核心程式設置參數
status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&outputBuffer);
if (status != CL_SUCCESS) {
printf("Error: Setting kernel argument. (clSetKernelArg)\n");
return EXIT_FAILURE;
}
// 建立一個OpenCL command queue
cl_command_queue commandQueue = clCreateCommandQueue(context,
devices[0],
0,
&status);
if (status != CL_SUCCESS) {
printf("Error: Create Command Queue. (clCreateCommandQueue)\n");
return EXIT_FAILURE;
}
// 將一個kernel 放入 command queue
size_t globalThreads[] = {4, 4};
size_t localThreads[] = {2, 2};
status = clEnqueueNDRangeKernel(commandQueue, kernel,
2, NULL, globalThreads,
localThreads, 0,
NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Enqueueing kernel\n");
return EXIT_FAILURE;
}
// 確認 command queue 中所有命令都執行完畢
status = clFinish(commandQueue);
if (status != CL_SUCCESS) {
printf("Error: Finish command queue\n");
return EXIT_FAILURE;
}
// 將記憶體物件中的結果讀回Host
status = clEnqueueReadBuffer(commandQueue,
outputBuffer, CL_TRUE, 0,
4 * 4 * 4, outbuffer, 0, NULL, NULL);
if (status != CL_SUCCESS) {
printf("Error: Read buffer queue\n");
return EXIT_FAILURE;
}
// Host端列印結果
printf("out:\n");
for (int i = 0; i < 16; ++i) {
printf("%x ", outbuffer[i]);
if ((i + 1) % 4 == 0)
printf("\n");
}
// 資源回收
status = clReleaseKernel(kernel);
status = clReleaseProgram(program);
status = clReleaseMemObject(outputBuffer);
status = clReleaseCommandQueue(commandQueue);
status = clReleaseContext(context);
free(devices);
delete outbuffer;
system("pause");
return 0;
}
先使用compile檔編譯一個bin，然後使用run檔載入bin運行。
http://www.photoneray.com/opencl_03/
留言列表
程式語言 (17)
