[[414533]]

本文轉(zhuǎn)載自微信公眾號「LoyenWang」，作者LoyenWang。轉(zhuǎn)載本文請聯(lián)系LoyenWang公眾號。

背景

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

說明：

• 對不起，我竟然用了一個奪人眼球的標(biāo)題;
• 我會盡量從一個程序員的角度來闡述OpenCL，目標(biāo)是淺顯易懂，如果沒有達(dá)到這個效果，就當(dāng)我沒說這話;
• 子曾經(jīng)曰過：不懂Middleware的系統(tǒng)軟件工程師，不是一個好碼農(nóng);

1. 介紹

• OpenCL(Open Computing Language，開放計算語言)：從軟件視角看，它是用于異構(gòu)平臺編程的框架;從規(guī)范視角看，它是異構(gòu)并行計算的行業(yè)標(biāo)準(zhǔn)，由Khronos Group來維護(hù);
• 異構(gòu)平臺包括了CPU、GPU、FPGA、DSP，以及最近幾年流行的各類AI加速器等;
• OpenCL包含兩部分：

1)用于編寫運行在OpenCL device上的kernels的語言(基于C99);

2)OpenCL API，至于Runtime的實現(xiàn)交由各個廠家，比如Intel發(fā)布的opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz

以人工智能場景為例來理解一下，假如在某個AI芯片上跑人臉識別應(yīng)用，CPU擅長控制，AI processor擅長計算，軟件的flow就可以進(jìn)行拆分，用CPU來負(fù)責(zé)控制視頻流輸入輸出前后處理，AI processor來完成深度學(xué)習(xí)模型運算完成識別，這就是一個典型的異構(gòu)處理場景，如果該AI芯片的SDK支持OpenCL，那么上層的軟件就可以基于OpenCL進(jìn)行開發(fā)了。

話不多說，看看OpenCL的架構(gòu)吧。

2. OpenCL架構(gòu)

OpenCL架構(gòu)，可以從平臺模型、內(nèi)存模型、執(zhí)行模型、編程模型四個角度來展開。

2.1 Platform Model

平臺模型：硬件拓?fù)潢P(guān)系的抽象描述

• 平臺模型由一個Host連接一個或多個OpenCL Devices組成;
• OpenCL Device，可以劃分成一個或多個計算單元Compute Unit(CU);
• CU可以進(jìn)一步劃分成一個或多個處理單元Processing Unit(PE)，最終的計算由PE來完成;
• OpenCL應(yīng)用程序分成兩部分：host代碼和device kernel代碼，其中Host運行host代碼，并將kernel代碼以命令的方式提交到OpenCL devices，由OpenCL device來運行kernel代碼;

2.2 Execution Model

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

Context

OpenCL的Execution Model由兩個不同的執(zhí)行單元定義：1)運行在OpenCL設(shè)備上的kernel;2)運行在Host上的Host program;其中，OpenCL使用Context代表kernel的執(zhí)行環(huán)境：

Context包含以下資源：

• Devices：一個或多個OpenCL設(shè)備;
• Kernel Objects：OpenCL Device的執(zhí)行函數(shù)及相關(guān)的參數(shù)值，通常定義在cl文件中;
• Program Objects：實現(xiàn)kernel的源代碼和可執(zhí)行程序，每個program可以包含多個kernel;
• Memory Objects：Host和OpenCL設(shè)備可見的變量，kernel執(zhí)行時對其進(jìn)行操作;

NDrange

• kernel是Execution Model的核心，放置在設(shè)備上執(zhí)行，當(dāng)kernel執(zhí)行前，需要創(chuàng)建一個索引空間NDRange(一維/二維/三維);
• 執(zhí)行kernel實例的稱為work-item，work-item組織成work-group，work-group組織成NDRange，最終將NDRange映射到OpenCL Device的計算單元上;

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

1. 通過work-item的全局索引;
2. 先查找到所在work-group的索引號，再根據(jù)局部索引號確定;

以一維為例：

• 上圖中總共有四個work-group，每個work-group包含四個work-item，所以local_size的大小為4，而local_id都是從0開始重新計數(shù);
• global_size代表總體的大小，也就是16個work-item，而global_id則是從0開始計數(shù);

以二維為例：

• 二維的計算方式與一維類似，也是結(jié)合global和local的size，可以得出global_id和local_id的大小，細(xì)節(jié)不表了;

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

2.3 Memory Model

內(nèi)存模型：Host和OpenCL Device怎么來看待數(shù)據(jù)

OpenCL的內(nèi)存模型中，包含以下幾類類型的內(nèi)存：

• Host memory：Host端的內(nèi)存，只能由Host直接訪問;
• Global Memory：設(shè)備內(nèi)存，可以由Host和OpenCL Device訪問，允許Host的讀寫操作，也允許OpenCL Device中PE讀寫，Host負(fù)責(zé)該內(nèi)存中Buffer的分配和釋放;
• Constant Global Memory：設(shè)備內(nèi)存，允許Host進(jìn)行讀寫操作，而設(shè)備只能進(jìn)行讀操作，用于傳輸常量數(shù)據(jù);
• Local Memory：單個CU中的本地內(nèi)存，Host看不到該區(qū)域并無法對其操作，該區(qū)域允許內(nèi)部的PE進(jìn)行讀寫操作，也可以用于PE之間的共享，需要注意同步和并發(fā)問題;
• Private Memory：PE的私有內(nèi)存，Host與PE之間都無法看到該區(qū)域;

2.4 Programming Model

• 在編程模型中，有兩部分代碼需要編寫：一部分是Host端，一部分是OpenCL Device端;
• 編程過程中，核心是要維護(hù)一個Context，代表了整個Kernel執(zhí)行的環(huán)境;
• 從cl源代碼中創(chuàng)建Program對象并編譯，在運行時創(chuàng)建Kernel對象以及內(nèi)存對象，設(shè)置好相關(guān)的參數(shù)和輸入之后，就可以將Kernel送入到隊列中執(zhí)行，也就是Launch kernel的流程;
• 最終等待運算結(jié)束，獲取計算結(jié)果即可;

3. 編程流程

• 上圖為一個OpenCL應(yīng)用開發(fā)涉及的基本過程;

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

4. 示例代碼

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

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函數(shù)

• 在Host程序中，創(chuàng)建program對象時會去讀取kernel的源代碼，本示例源代碼位于：vector_add.cl文件中

內(nèi)容如下：

__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