一文說清OpenCL框架

背景

• Read the fucking official documents! --By 魯迅
• A picture is worth a thousand words. --By 高爾基

說明：

• 對不起，我居然用了一個奪人眼球的標題；
• 我會盡可能從一個程序員的角度來闡述OpenCL，目標是淺顯易懂，若是沒有達到這個效果，就當我沒說這話；
• 子曾經曰過：不懂Middleware的系統軟件工程師，不是一個好碼農；

1. 介紹

• OpenCL(Open Computing Language，開放計算語言）：
  從軟件視角看，它是用於異構平臺編程的框架；
  從規範視角看，它是異構並行計算的行業標準，由Khronos Group來維護；
• 異構平臺包括了CPU、GPU、FPGA、DSP，以及最近幾年流行的各種AI加速器等；
• OpenCL包含兩部分：
  1）用於編寫運行在OpenCL device上的kernels的語言（基於C99）；
  2）OpenCL API，至於Runtime的實現交由各個廠家，好比Intel發佈的opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz

 以人工智能場景爲例來理解一下，假如在某個AI芯片上跑人臉識別應用，CPU擅長控制，AI processor擅長計算，軟件的flow就能夠進行拆分，用CPU來負責控制視頻流輸入輸出先後處理，AI processor來完成深度學習模型運算完成識別，這就是一個典型的異構處理場景，若是該AI芯片的SDK支持OpenCL，那麼上層的軟件就能夠基於OpenCL進行開發了。

話很少說，看看OpenCL的架構吧。

2. OpenCL架構

OpenCL架構，能夠從平臺模型、內存模型、執行模型、編程模型四個角度來展開。

2.1 Platform Model

平臺模型：硬件拓撲關係的抽象描述

• 平臺模型由一個Host鏈接一個或多個OpenCL Devices組成；
• OpenCL Device，能夠劃分紅一個或多個計算單元Compute Unit（CU）；
• CU能夠進一步劃分紅一個或多個處理單元Processing Unit（PE），最終的計算由PE來完成；
• OpenCL應用程序分紅兩部分：host代碼和device kernel代碼，其中Host運行host代碼，並將kernel代碼以命令的方式提交到OpenCL devices，由OpenCL device來運行kernel代碼；

2.2 Execution Model

執行模型：Host如何利用OpenCL Device的計算資源完成高效的計算處理過程

Context

OpenCL的Execution Model由兩個不一樣的執行單元定義：1）運行在OpenCL設備上的kernel；2）運行在Host上的Host program；
其中，OpenCL使用Context表明kernel的執行環境：

Context包含如下資源：

• Devices：一個或多個OpenCL設備；
• Kernel Objects：OpenCL Device的執行函數及相關的參數值，一般定義在cl文件中；
• Program Objects：實現kernel的源代碼和可執行程序，每一個program能夠包含多個kernel；
• Memory Objects：Host和OpenCL設備可見的變量，kernel執行時對其進行操做；

NDrange

• kernel是Execution Model的核心，放置在設備上執行，當kernel執行前，須要建立一個索引空間NDRange（一維/二維/三維）；
• 執行kernel實例的稱爲work-item，work-item組織成work-group，work-group組織成NDRange，最終將NDRange映射到OpenCL Device的計算單元上；

有兩種方式來找到work-item：

1. 經過work-item的全局索引；
2. 先查找到所在work-group的索引號，再根據局部索引號肯定；

以一維爲例：

• 上圖中總共有四個work-group，每一個work-group包含四個work-item，因此local_size的大小爲4，而local_id都是從0開始從新計數；
• global_size表明整體的大小，也就是16個work-item，而global_id則是從0開始計數；

以二維爲例：

• 二維的計算方式與一維相似，也是結合global和local的size，能夠得出global_id和local_id的大小，細節不表了；

三維的方式也相似，略去。

2.3 Memory Model

內存模型：Host和OpenCL Device怎麼來看待數據

OpenCL的內存模型中，包含如下幾類類型的內存：

• Host memory：Host端的內存，只能由Host直接訪問；
• Global Memory：設備內存，能夠由Host和OpenCL Device訪問，容許Host的讀寫操做，也容許OpenCL Device中PE讀寫，Host負責該內存中Buffer的分配和釋放；
• Constant Global Memory：設備內存，容許Host進行讀寫操做，而設備只能進行讀操做，用於傳輸常量數據；
• Local Memory：單個CU中的本地內存，Host看不到該區域並沒有法對其操做，該區域容許內部的PE進行讀寫操做，也能夠用於PE之間的共享，須要注意同步和併發問題；
• Private Memory：PE的私有內存，Host與PE之間都沒法看到該區域；

2.4 Programming Model

• 在編程模型中，有兩部分代碼須要編寫：一部分是Host端，一部分是OpenCL Device端；
• 編程過程當中，核心是要維護一個Context，表明了整個Kernel執行的環境；
• 從cl源代碼中建立Program對象並編譯，在運行時建立Kernel對象以及內存對象，設置好相關的參數和輸入以後，就能夠將Kernel送入到隊列中執行，也就是Launch kernel的流程；
• 最終等待運算結束，獲取計算結果便可；

3. 編程流程

• 上圖爲一個OpenCL應用開發涉及的基本過程；

下邊來一個實際的代碼測試跑跑，Talk is cheap, show me the code!

4. 示例代碼

• 測試環境：Ubuntu16.04，安裝Intel CPU OpenCL SDK（opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz）；
• 爲了簡化流程，示例代碼都不作容錯處理，僅保留關鍵的操做；
• 整個代碼的功能是完成向量的加法操做；

4.1 Host端程序

#include <iostream>
#include <fstream>
#include <sstream>

#include <CL/cl.h>

const int DATA_SIZE = 10;

int main(void)
{
    /* 1. get platform & device information */
    cl_uint num_platforms;
    cl_platform_id first_platform_id;
    clGetPlatformIDs(1, &first_platform_id, &num_platforms);


    /* 2. create context */
    cl_int err_num;
    cl_context context = nullptr;
    cl_context_properties context_prop[] = {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)first_platform_id,
        0
    };
    context = clCreateContextFromType(context_prop, CL_DEVICE_TYPE_CPU, nullptr, nullptr, &err_num);


    /* 3. create command queue */
    cl_command_queue command_queue;
    cl_device_id *devices;
    size_t device_buffer_size = -1;

    clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, nullptr, &device_buffer_size);
    devices = new cl_device_id[device_buffer_size / sizeof(cl_device_id)];
    clGetContextInfo(context, CL_CONTEXT_DEVICES, device_buffer_size, devices, nullptr);
    command_queue = clCreateCommandQueueWithProperties(context, devices[0], nullptr, nullptr);
    delete [] devices;


    /* 4. create program */
    std::ifstream kernel_file("vector_add.cl", std::ios::in);
    std::ostringstream oss;

    oss << kernel_file.rdbuf();
    std::string srcStdStr = oss.str();
    const char *srcStr = srcStdStr.c_str();
    cl_program program;
    program = clCreateProgramWithSource(context, 1, (const char **)&srcStr, nullptr, nullptr);


    /* 5. build program */
    clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr);


    /* 6. create kernel */
    cl_kernel kernel;
    kernel = clCreateKernel(program, "vector_add", nullptr);


    /* 7. set input data && create memory object */
    float output[DATA_SIZE];
    float input_x[DATA_SIZE];
    float input_y[DATA_SIZE];
    for (int i = 0; i < DATA_SIZE; i++) {
        input_x[i] = (float)i;
        input_y[i] = (float)(2 * i);
    }

    cl_mem mem_object_x;
    cl_mem mem_object_y;
    cl_mem mem_object_output;
    mem_object_x = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_x, nullptr);
    mem_object_y = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_y, nullptr);
    mem_object_output = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * DATA_SIZE, nullptr, nullptr);


    /* 8. set kernel argument */
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &mem_object_x);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &mem_object_y);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &mem_object_output);


    /* 9. send kernel to execute */
    size_t globalWorkSize[1] = {DATA_SIZE};
    size_t localWorkSize[1] = {1};
    clEnqueueNDRangeKernel(command_queue, kernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);


    /* 10. read data from output */
    clEnqueueReadBuffer(command_queue, mem_object_output, CL_TRUE, 0, DATA_SIZE * sizeof(float), output, 0, nullptr, nullptr);
    for (int i = 0; i < DATA_SIZE; i++) {
        std::cout << output[i] << " ";
    }
    std::cout << std::endl;


    /* 11. clean up */
    clRetainMemObject(mem_object_x);
    clRetainMemObject(mem_object_y);
    clRetainMemObject(mem_object_output);
    clReleaseCommandQueue(command_queue);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseContext(context);

    return 0;
}

4.2 OpenCL Kernel函數

• 在Host程序中，建立program對象時會去讀取kernel的源代碼，本示例源代碼位於：vector_add.cl文件中

內容以下：

__kernel void vector_add(__global const float *input_x,
	__global const float *input_y,
	__global float *output)
{
	int gid = get_global_id(0);
 
	output[gid] = input_x[gid] + input_y[gid];
}

4.3 輸出

參考

The OpenCL Specification

歡迎關注公衆號，不按期分享技術文章