這是第一篇真正的OpenCL教程。這篇文章不會(huì)從GPU結(jié)構(gòu)的技術(shù)概念和性能指標(biāo)入手。我們將會(huì)從OpenCL的基礎(chǔ)API開(kāi)始，使用一個(gè)小的kernel作為例子來(lái)講解基本的計(jì)算管理。
首先我們需要明白的是，OpenCL程序是分成兩部分的：一部分是在設(shè)備上執(zhí)行的（對(duì)于我們，是GPU），另一部分是在主機(jī)上運(yùn)行的（對(duì)于我們，是CPU）。在設(shè)備上執(zhí)行的程序或許是你比較關(guān)注的。它是OpenCL產(chǎn)生神奇力量的地方。為了能在設(shè)備上執(zhí)行代碼，程序員需要寫(xiě)一個(gè)特殊的函數(shù)（kernel函數(shù)）。這個(gè)函數(shù)需要使用OpenCL語(yǔ)言編寫(xiě)。OpenCL語(yǔ)言采用了C語(yǔ)言的一部分加上一些約束、關(guān)鍵字和數(shù)據(jù)類(lèi)型。在主機(jī)上運(yùn)行的程序提供了API，所以i可以管理你在設(shè)備上運(yùn)行的程序。主機(jī)程序可以用C或者C++編寫(xiě)，它控制OpenCL的環(huán)境（上下文，指令隊(duì)列…）。
設(shè)備（Device）
我們來(lái)簡(jiǎn)單的說(shuō)一下設(shè)備。設(shè)備，像上文介紹的一樣，OpenCL編程最給力的地方。
我們必須了解一些基本概念：
Kernel：你可以把它想像成一個(gè)可以在設(shè)備上執(zhí)行的函數(shù)。當(dāng)然也會(huì)有其他可以在設(shè)備上執(zhí)行的函數(shù)，但是他們之間是有一些區(qū)別的。Kernel是設(shè)備程序執(zhí)行的入口點(diǎn)。換言之，Kernel是唯一可以從主機(jī)上調(diào)用執(zhí)行的函數(shù)。
現(xiàn)在的問(wèn)題是：我們?nèi)绾蝸?lái)編寫(xiě)一個(gè)Kernel？在Kernel中如何表達(dá)并行性？它的執(zhí)行模型是怎樣的？解決這些問(wèn)題，我們需要引入下面的概念：
SIMT：?jiǎn)沃噶疃嗑�程（SINGLE INSTRUCTION MULTI THREAD）的簡(jiǎn)寫(xiě)。就像這名字一樣，相同的代碼在不同線程中并行執(zhí)行，每個(gè)線程使用不同的數(shù)據(jù)來(lái)執(zhí)行同一段代碼。
Work-item（工作項(xiàng)）：Work-item與CUDA Threads是一樣的，是最小的執(zhí)行單元。每次一個(gè)Kernel開(kāi)始執(zhí)行，很多（程序員定義數(shù)量）的Work-item就開(kāi)始運(yùn)行，每個(gè)都執(zhí)行同樣的代碼。每個(gè)work-item有一個(gè)ID，這個(gè)ID在kernel中是可以訪問(wèn)的，每個(gè)運(yùn)行在work-item上的kernel通過(guò)這個(gè)ID來(lái)找出work-item需要處理的數(shù)據(jù)。
Work-group（工作組）：work-group的存在是為了允許work-item之間的通信和協(xié)作。它反映出work-item的組織形式（work-group是以N維網(wǎng)格形式組織的，N=1，2或3）。
Work-group等價(jià)于CUDA thread blocks。像work-items一樣，work-groups也有一個(gè)kernel可以讀取的唯一的ID。
ND-Range:ND-Range是下一個(gè)組織級(jí)別，定義了work-group的組織形式（ND-Rang以N維網(wǎng)格形式組織的，N=1，2或3）；
這是ND-Range組織形式的例子
Kernel
現(xiàn)在該寫(xiě)我們的第一個(gè)kernel了。我們寫(xiě)一個(gè)小的kernel將兩個(gè)向量相加。這個(gè)kernel需要四個(gè)參數(shù)：兩個(gè)要相加的向量，一個(gè)存儲(chǔ)結(jié)果的向量，和向量個(gè)數(shù)。如果你寫(xiě)一個(gè)程序在cpu上解決這個(gè)問(wèn)題，將會(huì)是下面這個(gè)樣子：
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上，邏輯就會(huì)有一些不同。我們使每個(gè)線程計(jì)算一個(gè)元素的方法來(lái)代替cpu程序中的循環(huán)計(jì)算。每個(gè)線程的index與要計(jì)算的向量的index相同。我們來(lái)看一下代碼實(shí)現(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í)行的這個(gè)線程的ID。
許多線程會(huì)在同一時(shí)間開(kāi)始執(zhí)行同一個(gè)kernel，
每個(gè)線程都會(huì)收到一個(gè)不同的ID，所以必然會(huì)執(zhí)行一個(gè)不同的計(jì)算。*/
const int idx = get_global_id(0);
/* 每個(gè)work-item都會(huì)檢查自己的id是否在向量數(shù)組的區(qū)間內(nèi)。
如果在，work-item就會(huì)執(zhí)行相應(yīng)的計(jì)算。*/
if (idx < num)
res[idx] = src_a[idx] + src_b[idx];
}
有一些需要注意的地方：
1. Kernel關(guān)鍵字定義了一個(gè)函數(shù)是kernel函數(shù)。Kernel函數(shù)必須返回void。
2. Global關(guān)鍵字位于參數(shù)前面。它定義了參數(shù)內(nèi)存的存放位置。
另外，所有kernel都必須寫(xiě)在“.cl”文件中，“.cl”文件必須只包含OpenCL代碼。
主機(jī)（Host）
我們的kernel已經(jīng)寫(xiě)好了，現(xiàn)在我們來(lái)寫(xiě)host程序。
建立基本OpenCL運(yùn)行環(huán)境
有一些東西我們必須要弄清楚：
Plantform（平臺(tái)）：主機(jī)加上OpenCL框架管理下的若干設(shè)備構(gòu)成了這個(gè)平臺(tái)，通過(guò)這個(gè)平臺(tái)，應(yīng)用程序可以與設(shè)備共享資源并在設(shè)備上執(zhí)行kernel。平臺(tái)通過(guò)cl_plantform來(lái)展現(xiàn)，可以使用下面的代碼來(lái)初始化平臺(tái)：
// Returns the error code cl_int oclGetPlatformID (cl_platform_id *platforms) // Pointer to the platform object
Device（設(shè)備）：通過(guò)cl_device來(lái)表現(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（上下文）：定義了整個(gè)OpenCL化境，包括OpenCL kernel、設(shè)備、內(nèi)存管理、命令隊(duì)列等。上下文使用cl_context來(lái)表現(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（指令隊(duì)列）：就像它的名字一樣，他是一個(gè)存儲(chǔ)需要在設(shè)備上執(zhí)行的OpenCL指令的隊(duì)列�！爸噶铌�(duì)列建立在一個(gè)上下文中的指定設(shè)備上。多個(gè)指令隊(duì)列允許應(yīng)用程序在不需要同步的情況下執(zhí)行多條無(wú)關(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)存
主機(jī)的基本環(huán)境已經(jīng)配置好了，為了可以執(zhí)行我們的寫(xiě)的小kernel，我們需要分配3個(gè)向量的內(nèi)存空間，然后至少初始化它們其中的兩個(gè)。
在主機(jī)環(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類(lèi)型，像下面這樣：
// 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是逐位的，選項(xiàng)如下：
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ù)
我們通過(guò)下面的代碼使用這個(gè)函數(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)在為止，你可能會(huì)問(wèn)自己一些問(wèn)題，比如：我們?cè)趺凑{(diào)用kernel？編譯器怎么知道如何將代碼放到設(shè)備上？我們?cè)趺淳幾gkernel？
下面是我們?cè)趯?duì)比OpenCL程序和OpenCL kernel時(shí)的一些容易混亂的概念：
Kernel：你應(yīng)該已經(jīng)知道了，像在上文中描述的一樣，kernel本質(zhì)上是一個(gè)我們可以從主機(jī)上調(diào)用的，運(yùn)行在設(shè)備上的函數(shù)。你或許不知道kernel是在運(yùn)行的時(shí)候編譯的！更一般的講，所有運(yùn)行在設(shè)備上的代碼，包括kernel和kernel調(diào)用的其他的函數(shù)，都是在運(yùn)行的時(shí)候編譯的。這涉及到下一個(gè)概念，Program。
Program：OpenCL Program由kernel函數(shù)、其他函數(shù)和聲明組成。它通過(guò)cl_program表示。當(dāng)創(chuàng)建一個(gè)program時(shí)，你必須指定它是由哪些文件組成的，然后編譯它。
你需要用到下面的函數(shù)來(lái)建立一個(gè)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的入口點(diǎn)。使用cl_kernel：
cl_kernel clCreateKernel (cl_program program, // The program where the kernel is const char *kernel_name, // The name of the kernel, i.e. the name of the kernel function as it's declared in the code cl_int *errcode_ret)
注意我們可以創(chuàng)建多個(gè)OpenCL program，每個(gè)program可以擁有多個(gè)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);
運(yùn)行kernel
一旦我們的kernel建立好，我們就可以運(yùn)行它。
首先，我們必須設(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
每個(gè)參數(shù)都需要調(diào)用一次這個(gè)函數(shù)。
當(dāng)所有參數(shù)設(shè)置完畢，我們就可以調(diào)用這個(gè)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é)果非常簡(jiǎn)單。與之前講到的寫(xiě)入內(nèi)存（設(shè)備內(nèi)存）的操作相似，現(xiàn)在我們需要存入隊(duì)列一個(gè)讀取緩沖區(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ǐng)的（緩沖區(qū)、kernel、隊(duì)列）必須使用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);
這是文章的全部?jī)?nèi)容了，碼農(nóng)們，作者最后說(shuō)，如果你有任何問(wèn)題，都可以馬上聯(lián)系他。
下載 
下載 
下載 
下載 
下載 
下載 
喜歡 
頂
難過(guò)
囧
圍觀
無(wú)聊
