背景

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：

通過work-item的全域索引；
先查找到所在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

歡迎關注公眾號，不定期分享技術文章

轉載請註明出處，本文鏈接：https://www.uj5u.com/qita/291162.html

標籤：其他

上一篇：IDEA工具匯入專案

下一篇：一文說清OpenCL框架