這是第一篇真正的OpenCL教程。這篇文章不會從GPU結(jié)構(gòu)的技術(shù)概念和性能指標(biāo)入手。我們將會從OpenCL的基礎(chǔ)API開始，使用一個小的kernel作為例子來講解基本的計算管理。

首先我們需要明白的是，OpenCL程序是分成兩部分的：一部分是在設(shè)備上執(zhí)行的（對于我們，是GPU），另一部分是在主機上運行的（對于我們，是CPU）。在設(shè)備上執(zhí)行的程序或許是你比較關(guān)注的。它是OpenCL產(chǎn)生神奇力量的地方。為了能在設(shè)備上執(zhí)行代碼，程序員需要寫一個特殊的函數(shù)（kernel函數(shù)）。這個函數(shù)需要使用OpenCL語言編寫。OpenCL語言采用了C語言的一部分加上一些約束、關(guān)鍵字和數(shù)據(jù)類型。在主機上運行的程序提供了API，所以i可以管理你在設(shè)備上運行的程序。主機程序可以用C或者C++編寫，它控制OpenCL的環(huán)境（上下文，指令隊列…）。

設(shè)備（Device）

我們來簡單的說一下設(shè)備。設(shè)備，像上文介紹的一樣，OpenCL編程最給力的地方。

我們必須了解一些基本概念：

Kernel：你可以把它想像成一個可以在設(shè)備上執(zhí)行的函數(shù)。當(dāng)然也會有其他可以在設(shè)備上執(zhí)行的函數(shù)，但是他們之間是有一些區(qū)別的。Kernel是設(shè)備程序執(zhí)行的入口點。換言之，Kernel是唯一可以從主機上調(diào)用執(zhí)行的函數(shù)。

現(xiàn)在的問題是：我們?nèi)绾蝸砭帉懸粋€Kernel？在Kernel中如何表達并行性？它的執(zhí)行模型是怎樣的？解決這些問題，我們需要引入下面的概念：

? ? SIMT：單指令多線程（SINGLE INSTRUCTION MULTI THREAD）的簡寫。就像這名字一樣，相同的代碼在不同線程中并行執(zhí)行，每個線程使用不同的數(shù)據(jù)來執(zhí)行同一段代碼。

? ? Work-item（工作項）：Work-item與CUDA Threads是一樣的，是最小的執(zhí)行單元。每次一個Kernel開始執(zhí)行，很多（程序員定義數(shù)量）的Work-item就開始運行，每個都執(zhí)行同樣的代碼。每個work-item有一個ID，這個ID在kernel中是可以訪問的，每個運行在work-item上的kernel通過這個ID來找出work-item需要處理的數(shù)據(jù)。

? ? Work-group（工作組）：work-group的存在是為了允許work-item之間的通信和協(xié)作。它反映出work-item的組織形式（work-group是以N維網(wǎng)格形式組織的，N=1，2或3）。

Work-group等價于CUDA thread blocks。像work-items一樣，work-groups也有一個kernel可以讀取的唯一的ID。

? ? ND-Range:ND-Range是下一個組織級別，定義了work-group的組織形式（ND-Rang以N維網(wǎng)格形式組織的，N=1，2或3）；

這是ND-Range組織形式的例子

Kernel

現(xiàn)在該寫我們的第一個kernel了。我們寫一個小的kernel將兩個向量相加。這個kernel需要四個參數(shù)：兩個要相加的向量，一個存儲結(jié)果的向量，和向量個數(shù)。如果你寫一個程序在cpu上解決這個問題，將會是下面這個樣子：

void vector_add_cpu (const float* src_a,
?????????????? const float* src_b,
?????????????? float*? res,
?????????????? const int num)
{
?? for (int i = 0; i < num; i++)
????? res[i] = src_a[i] + src_b[i];
}

在GPU上，邏輯就會有一些不同。我們使每個線程計算一個元素的方法來代替cpu程序中的循環(huán)計算。每個線程的index與要計算的向量的index相同。我們來看一下代碼實現(xiàn)：

__kernel void vector_add_gpu (__global const float* src_a,
???????????????????? __global const float* src_b,
???????????????????? __global float* res,
?????????? const int num)
{
?? /* get_global_id(0) 返回正在執(zhí)行的這個線程的ID。
?? 許多線程會在同一時間開始執(zhí)行同一個kernel，
?? 每個線程都會收到一個不同的ID，所以必然會執(zhí)行一個不同的計算。*/
?? const int idx = get_global_id(0);

?? /* 每個work-item都會檢查自己的id是否在向量數(shù)組的區(qū)間內(nèi)。
?? 如果在，work-item就會執(zhí)行相應(yīng)的計算。*/
?? if (idx < num)
????? res[idx] = src_a[idx] + src_b[idx];
}

有一些需要注意的地方：

1. Kernel關(guān)鍵字定義了一個函數(shù)是kernel函數(shù)。Kernel函數(shù)必須返回void。

2. Global關(guān)鍵字位于參數(shù)前面。它定義了參數(shù)內(nèi)存的存放位置。

另外，所有kernel都必須寫在“.cl”文件中，“.cl”文件必須只包含OpenCL代碼。

主機（Host）

我們的kernel已經(jīng)寫好了，現(xiàn)在我們來寫host程序。

建立基本OpenCL運行環(huán)境

有一些東西我們必須要弄清楚：

Plantform（平臺）：主機加上OpenCL框架管理下的若干設(shè)備構(gòu)成了這個平臺，通過這個平臺，應(yīng)用程序可以與設(shè)備共享資源并在設(shè)備上執(zhí)行kernel。平臺通過cl_plantform來展現(xiàn)，可以使用下面的代碼來初始化平臺：

// Returns the error codecl_int oclGetPlatformID (cl_platform_id *platforms) // Pointer to the platform object

Device（設(shè)備）：通過cl_device來表現(xiàn)，使用下面的代碼：

// Returns the error code

cl_int clGetDeviceIDs (cl_platform_id platform,

cl_device_type device_type, // Bitfield identifying the type. For the GPU we use CL_DEVICE_TYPE_GPU

cl_uint num_entries, // Number of devices, typically 1

cl_device_id *devices, // Pointer to the device object

cl_uint *num_devices) // Puts here the number of devices matching the device_type

Context（上下文）：定義了整個OpenCL化境，包括OpenCL kernel、設(shè)備、內(nèi)存管理、命令隊列等。上下文使用cl_context來表現(xiàn)。使用以下代碼初始化：

// Returs the context

cl_context clCreateContext (const cl_context_properties *properties, // Bitwise with the properties (see specification)

cl_uint num_devices, // Number of devices

const cl_device_id *devices, // Pointer to the devices object

void (*pfn_notify)(const char *errinfo, const void *private_info, size_t cb, void *user_data), // (don't worry about this)

void *user_data, // (don't worry about this)

cl_int *errcode_ret) // error code result

Command-Queue（指令隊列）：就像它的名字一樣，他是一個存儲需要在設(shè)備上執(zhí)行的OpenCL指令的隊列。“指令隊列建立在一個上下文中的指定設(shè)備上。多個指令隊列允許應(yīng)用程序在不需要同步的情況下執(zhí)行多條無關(guān)聯(lián)的指令。”

cl_command_queue clCreateCommandQueue (cl_context context,

cl_device_id device,

cl_command_queue_properties properties, // Bitwise with the properties

cl_int *errcode_ret) // error code result

下面的例子展示了這些元素的使用方法：

cl_int error = 0;?? // Used to handle error codes
cl_platform_id platform;
cl_context context;
cl_command_queue queue;
cl_device_id device;

// Platform
error = oclGetPlatformID(&platform);
if (error != CL_SUCCESS) {
?? cout << "Error getting platform id: " << errorMessage(error) << endl;
?? exit(error);
}
// Device
error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
if (err != CL_SUCCESS) {
?? cout << "Error getting device ids: " << errorMessage(error) << endl;
?? exit(error);
}
// Context
context = clCreateContext(0, 1, &device, NULL, NULL, &error);
if (error != CL_SUCCESS) {
?? cout << "Error creating context: " << errorMessage(error) << endl;
?? exit(error);
}
// Command-queue
queue = clCreateCommandQueue(context, device, 0, &error);
if (error != CL_SUCCESS) {
?? cout << "Error creating command queue: " << errorMessage(error) << endl;
?? exit(error);
}

分配內(nèi)存

主機的基本環(huán)境已經(jīng)配置好了，為了可以執(zhí)行我們的寫的小kernel，我們需要分配3個向量的內(nèi)存空間，然后至少初始化它們其中的兩個。

在主機環(huán)境下執(zhí)行這些操作，我們需要像下面的代碼這樣去做：

const int size = 1234567
float* src_a_h = new float[size];
float* src_b_h = new float[size];
float* res_h = new float[size];
// Initialize both vectors
for (int i = 0; i < size; i++) {
?? src_a_h = src_b_h = (float) i;
}

在設(shè)備上分配內(nèi)存，我們需要使用cl_mem類型，像下面這樣：

// Returns the cl_mem object referencing the memory allocated on the device

cl_mem clCreateBuffer (cl_context context, // The context where the memory will be allocated

cl_mem_flags flags,

size_t size, // The size in bytes

void *host_ptr,

cl_int *errcode_ret)

flags是逐位的，選項如下：

CL_MEM_READ_WRITE

CL_MEM_WRITE_ONLY

CL_MEM_READ_ONLY

CL_MEM_USE_HOST_PTR

CL_MEM_ALLOC_HOST_PTR

CL_MEM_COPY_HOST_PTR – 從 host_ptr處拷貝數(shù)據(jù)

我們通過下面的代碼使用這個函數(shù)：

const int mem_size = sizeof(float)*size;

// Allocates a buffer of size mem_size and copies mem_size bytes from src_a_h

cl_mem src_a_d = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size, src_a_h, &error);

cl_mem src_b_d = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size, src_b_h, &error);

cl_mem res_d = clCreateBuffer(context, CL_MEM_WRITE_ONLY, mem_size, NULL, &error);

程序和kernel

到現(xiàn)在為止，你可能會問自己一些問題，比如：我們怎么調(diào)用kernel？編譯器怎么知道如何將代碼放到設(shè)備上？我們怎么編譯kernel？

下面是我們在對比OpenCL程序和OpenCL kernel時的一些容易混亂的概念：

Kernel：你應(yīng)該已經(jīng)知道了，像在上文中描述的一樣，kernel本質(zhì)上是一個我們可以從主機上調(diào)用的，運行在設(shè)備上的函數(shù)。你或許不知道kernel是在運行的時候編譯的！更一般的講，所有運行在設(shè)備上的代碼，包括kernel和kernel調(diào)用的其他的函數(shù)，都是在運行的時候編譯的。這涉及到下一個概念，Program。

Program：OpenCL Program由kernel函數(shù)、其他函數(shù)和聲明組成。它通過cl_program表示。當(dāng)創(chuàng)建一個program時，你必須指定它是由哪些文件組成的，然后編譯它。

你需要用到下面的函數(shù)來建立一個Program：

// Returns the OpenCL program

cl_program clCreateProgramWithSource (cl_context context,

??? cl_uint count, // number of files

??? const char **strings, // array of strings, each one is a file

??? const size_t *lengths, // array specifying the file lengths

??? cl_int *errcode_ret) // error code to be returned

當(dāng)我們創(chuàng)建了Program我們可以用下面的函數(shù)執(zhí)行編譯操作：

cl_int clBuildProgram (cl_program program,

??? cl_uint num_devices,

??? const cl_device_id *device_list,

??? const char *options, // Compiler options, see the specifications for more details

??? void (*pfn_notify)(cl_program, void *user_data),

??? void *user_data)

查看編譯log，必須使用下面的函數(shù)：

cl_int clGetProgramBuildInfo (cl_program program,

??? cl_device_id device,

??? cl_program_build_info param_name, // The parameter we want to know

??? size_t param_value_size,

??? void *param_value, // The answer

??? size_t *param_value_size_ret)

最后，我們需要“提取”program的入口點。使用cl_kernel：

cl_kernel clCreateKernel (cl_program program, // The program where the kernel isconst char *kernel_name, // The name of the kernel, i.e. the name of the kernel function as it's declared in the codecl_int *errcode_ret)

注意我們可以創(chuàng)建多個OpenCL program，每個program可以擁有多個kernel。

以下是這一章節(jié)的代碼：

// Creates the program
// Uses NVIDIA helper functions to get the code string and it's size (in bytes)
size_t src_size = 0;
const char* path = shrFindFilePath("vector_add_gpu.cl", NULL);
const char* source = oclLoadProgSource(path, "", &src_size);
cl_program program = clCreateProgramWithSource(context, 1, &source, &src_size, &error);
assert(error == CL_SUCCESS);

// Builds the program
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
assert(error == CL_SUCCESS);

// Shows the log
char* build_log;
size_t log_size;
// First call to know the proper size
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
build_log = new char[log_size+1];
// Second call to get the log
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, build_log, NULL);
build_log[log_size] = '\0';
cout << build_log << endl;
delete[] build_log;

// Extracting the kernel
cl_kernel vector_add_kernel = clCreateKernel(program, "vector_add_gpu", &error);
assert(error == CL_SUCCESS);

運行kernel

一旦我們的kernel建立好，我們就可以運行它。

首先，我們必須設(shè)置kernel的參數(shù)：

cl_int clSetKernelArg (cl_kernel kernel, // Which kernel

??? cl_uint arg_index, // Which argument

??? size_t arg_size, // Size of the next argument (not of the value pointed by it!)

??? const void *arg_value) // Value

每個參數(shù)都需要調(diào)用一次這個函數(shù)。

當(dāng)所有參數(shù)設(shè)置完畢，我們就可以調(diào)用這個kernel：

cl_int? clEnqueueNDRangeKernel (cl_command_queue command_queue,?
??????????????????????????? cl_kernel kernel,?
??????????????????????????? cl_uint? work_dim,??? // Choose if we are using 1D, 2D or 3D work-items and work-groups
??????????????????????????? const size_t *global_work_offset,
??????????????????????????? const size_t *global_work_size,?? // The total number of work-items (must have work_dim dimensions)
??????????????????????????? const size_t *local_work_size,???? // The number of work-items per work-group (must have work_dim dimensions)
??????????????????????????? cl_uint num_events_in_wait_list,?
??????????????????????????? const cl_event *event_wait_list,?
??????????????????????????? cl_event *event)

下面是這一章節(jié)的代碼：

// Enqueuing parameters
// Note that we inform the size of the cl_mem object, not the size of the memory pointed by it
error = clSetKernelArg(vector_add_k, 0, sizeof(cl_mem), &src_a_d);
error |= clSetKernelArg(vector_add_k, 1, sizeof(cl_mem), &src_b_d);
error |= clSetKernelArg(vector_add_k, 2, sizeof(cl_mem), &res_d);
error |= clSetKernelArg(vector_add_k, 3, sizeof(size_t), &size);
assert(error == CL_SUCCESS);

// Launching kernel
const size_t local_ws = 512;??? // Number of work-items per work-group
// shrRoundUp returns the smallest multiple of local_ws bigger than size
const size_t global_ws = shrRoundUp(local_ws, size);??? // Total number of work-items
error = clEnqueueNDRangeKernel(queue, vector_add_k, 1, NULL, &global_ws, &local_ws, 0, NULL, NULL);
assert(error == CL_SUCCESS);

讀取結(jié)果

讀取結(jié)果非常簡單。與之前講到的寫入內(nèi)存（設(shè)備內(nèi)存）的操作相似，現(xiàn)在我們需要存入隊列一個讀取緩沖區(qū)的操作：

cl_int? clEnqueueReadBuffer (cl_command_queue command_queue,?
????????????????????? cl_mem buffer,?? // from which buffer
????????????????????? cl_bool blocking_read,?? // whether is a blocking or non-blocking read
????????????????????? size_t offset,?? // offset from the beginning
????????????????????? size_t cb,?? // size to be read (in bytes)
????????????????????? void *ptr,?? // pointer to the host memory
????????????????????? cl_uint num_events_in_wait_list,
????????????????????? const cl_event *event_wait_list,?
????????????????????? cl_event *event)

使用方法如下：

// Reading back float* check = new float[size]; clEnqueueReadBuffer(queue, res_d, CL_TRUE, 0, mem_size, check, 0, NULL, NULL);

清理

作為一名牛X的程序員我們肯定要考慮如何清理內(nèi)存！

你需要知道最基本東西：使用clCreate申請的（緩沖區(qū)、kernel、隊列）必須使用clRelease釋放。

代碼如下：

// Cleaning up

delete[] src_a_h;

delete[] src_b_h;

delete[] res_h;

delete[] check;

clReleaseKernel(vector_add_k);

clReleaseCommandQueue(queue);

clReleaseContext(context);

clReleaseMemObject(src_a_d);

clReleaseMemObject(src_b_d);

clReleaseMemObject(res_d);

總結(jié)

以上是生活随笔為你收集整理的OpenCL快速入门的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。