OpenCL快速入門教程

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

 

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

 

譯者注：對文章內容有任何疑問或建議可以加opencl cuda新手群 242337476 一起討論。