[[414533]]

本文轉載自微信公眾號「LoyenWang」，作者LoyenWang。轉載本文請聯系LoyenWang公眾號。

背景

• 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