一文說清OpenCL框架

LoyenWang發表於2021-07-31
框架

背景

  • 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

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

相關文章