OpenCL OpenCL快速入門教程 OpenCL快速入門教程

阿新 • • 發佈：2018-12-27

OpenCL快速入門教程

原文地址：http://opencl.codeplex.com/wikipage?title=OpenCL%20Tutorials%20-%201

翻譯日期：2012年6月4日星期一

這是第一篇真正的OpenCL教程。這篇文章不會從GPU結構的技術概念和效能指標入手。我們將會從OpenCL的基礎API開始，使用一個小的kernel作為例子來講解基本的計算管理。

首先我們需要明白的是，OpenCL程式是分成兩部分的：一部分是在裝置上執行的（對於我們，是GPU），另一部分是在主機上執行的（對於我們，是CPU）。在裝置上執行的程式或許是你比較關注的。它是OpenCL產生神奇力量的地方。為了能在裝置上執行程式碼，程式設計師需要寫一個特殊的函式（kernel函式）。這個函式需要使用OpenCL語言編寫。OpenCL語言採用了C語言的一部分加上一些約束、關鍵字和資料型別。在主機上執行的程式提供了API，所以i可以管理你在裝置上執行的程式。主機程式可以用C或者C++編寫，它控制OpenCL的環境（上下文，指令佇列…）。

裝置（Device）

我們來簡單的說一下裝置。裝置，像上文介紹的一樣，OpenCL程式設計最給力的地方。

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

Kernel：你可以把它想像成一個可以在裝置上執行的函式。當然也會有其他可以在裝置上執行的函式，但是他們之間是有一些區別的。Kernel是裝置程式執行的入口點。換言之，Kernel是唯一可以從主機上呼叫執行的函式。

現在的問題是：我們如何來編寫一個Kernel？在Kernel中如何表達並行性？它的執行模型是怎樣的？解決這些問題，我們需要引入下面的概念：

SIMT：單指令多執行緒（SINGLE INSTRUCTION MULTI THREAD）的簡寫。就像這名字一樣，相同的程式碼在不同執行緒中並行執行，每個執行緒使用不同的資料來執行同一段程式碼。

Work-item（工作項）：Work-item與CUDA Threads是一樣的，是最小的執行單元。每次一個Kernel開始執行，很多（程式設計師定義數量）的Work-item就開始執行，每個都執行同樣的程式碼。每個work-item有一個ID，這個ID在kernel中是可以訪問的，每個執行在work-item上的kernel通過這個ID來找出work-item需要處理的資料。

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

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

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

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

Kernel

現在該寫我們的第一個kernel了。我們寫一個小的kernel將兩個向量相加。這個kernel需要四個引數：兩個要相加的向量，一個儲存結果的向量，和向量個數。如果你寫一個程式在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程式中的迴圈計算。每個執行緒的index與要計算的向量的index相同。我們來看一下程式碼實現：

__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) 返回正在執行的這個執行緒的ID。
   許多執行緒會在同一時間開始執行同一個kernel，
   每個執行緒都會收到一個不同的ID，所以必然會執行一個不同的計算。*/
   const int idx = get_global_id(0);

   /* 每個work-item都會檢查自己的id是否在向量陣列的區間內。
   如果在，work-item就會執行相應的計算。*/
   if (idx < num)
      res[idx] = src_a[idx] + src_b[idx];
}

有一些需要注意的地方：

1. Kernel關鍵字定義了一個函式是kernel函式。Kernel函式必須返回void。

2. Global關鍵字位於引數前面。它定義了引數記憶體的存放位置。

另外，所有kernel都必須寫在“.cl”檔案中，“.cl”檔案必須只包含OpenCL程式碼。

主機（Host）

我們的kernel已經寫好了，現在我們來寫host程式。

建立基本OpenCL執行環境

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

Plantform（平臺）：主機加上OpenCL框架管理下的若干裝置構成了這個平臺，通過這個平臺，應用程式可以與裝置共享資源並在裝置上執行kernel。平臺通過cl_plantform來展現，可以使用下面的程式碼來初始化平臺：

// Returns the error code

cl_int oclGetPlatformID (cl_platform_id *platforms) // Pointer to the platform object

Device（裝置）：通過cl_device來表現，使用下面的程式碼：

// 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、裝置、記憶體管理、命令佇列等。上下文使用cl_context來表現。使用以下程式碼初始化：

// 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（指令佇列）：就像它的名字一樣，他是一個儲存需要在裝置上執行的OpenCL指令的佇列。“指令佇列建立在一個上下文中的指定裝置上。多個指令佇列允許應用程式在不需要同步的情況下執行多條無關聯的指令。”

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);
}

分配記憶體

主機的基本環境已經配置好了，為了可以執行我們的寫的小kernel，我們需要分配3個向量的記憶體空間，然後至少初始化它們其中的兩個。

在主機環境下執行這些操作，我們需要像下面的程式碼這樣去做：

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;
}

在裝置上分配記憶體，我們需要使用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處拷貝資料

我們通過下面的程式碼使用這個函式：

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

到現在為止，你可能會問自己一些問題，比如：我們怎麼呼叫kernel？編譯器怎麼知道如何將程式碼放到裝置上？我們怎麼編譯kernel？

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

Kernel：你應該已經知道了，像在上文中描述的一樣，kernel本質上是一個我們可以從主機上呼叫的，執行在裝置上的函式。你或許不知道kernel是在執行的時候編譯的！更一般的講，所有執行在裝置上的程式碼，包括kernel和kernel呼叫的其他的函式，都是在執行的時候編譯的。這涉及到下一個概念，Program。

Program：OpenCL Program由kernel函式、其他函式和宣告組成。它通過cl_program表示。當建立一個program時，你必須指定它是由哪些檔案組成的，然後編譯它。

你需要用到下面的函式來建立一個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

當我們建立了Program我們可以用下面的函式執行編譯操作：

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，必須使用下面的函式：

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 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)

注意我們可以建立多個OpenCL program，每個program可以擁有多個kernel。

以下是這一章節的程式碼：

// 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建立好，我們就可以執行它。

首先，我們必須設定kernel的引數：

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

每個引數都需要呼叫一次這個函式。

當所有引數設定完畢，我們就可以呼叫這個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)

下面是這一章節的程式碼：

// 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);

讀取結果

讀取結果非常簡單。與之前講到的寫入記憶體（裝置記憶體）的操作相似，現在我們需要存入佇列一個讀取緩衝區的操作：

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的程式設計師我們肯定要考慮如何清理記憶體！

你需要知道最基本東西：使用clCreate申請的（緩衝區、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);

這是文章的全部內容了，碼農們，作者最後說，如果你有任何問題，都可以馬上聯絡他。

作者： Leo_wl 　　　　出處： http://www.cnblogs.com/Leo_wl/ 　　　　本文版權歸作者和部落格園共有，歡迎轉載，但未經作者同意必須保留此段宣告，且在文章頁面明顯位置給出原文連線，否則保留追究法律責任的權利。版權資訊

再分享一下我老師大神的人工智慧教程吧。零基礎！通俗易懂！風趣幽默！還帶黃段子！希望你也加入到我們人工智慧的隊伍中來！https://www.cnblogs.com/captainbed