一文說清OpenCL框架
作者：LoyenWang
OpenCL(Open Computing Language，開放計(jì)算語言）：從軟件視角看，它是用于異構(gòu)平臺(tái)編程的框架；從規(guī)范視角看，它是異構(gòu)并行計(jì)算的行業(yè)標(biāo)準(zhǔn)，由Khronos Group來維護(hù)。
本文轉(zhuǎn)載自微信公眾號(hào)「LoyenWang」，作者LoyenWang。轉(zhuǎn)載本文請(qǐng)聯(lián)系LoyenWang公眾號(hào)。
背景
- Read the fucking official documents! --By 魯迅
- A picture is worth a thousand words. --By 高爾基
說明：
- 對(duì)不起，我竟然用了一個(gè)奪人眼球的標(biāo)題;
- 我會(huì)盡量從一個(gè)程序員的角度來闡述OpenCL，目標(biāo)是淺顯易懂，如果沒有達(dá)到這個(gè)效果，就當(dāng)我沒說這話;
- 子曾經(jīng)曰過：不懂Middleware的系統(tǒng)軟件工程師，不是一個(gè)好碼農(nóng);
1. 介紹
- OpenCL(Open Computing Language，開放計(jì)算語言)：從軟件視角看，它是用于異構(gòu)平臺(tái)編程的框架;從規(guī)范視角看，它是異構(gòu)并行計(jì)算的行業(yè)標(biāo)準(zhǔn)，由Khronos Group來維護(hù);
- 異構(gòu)平臺(tái)包括了CPU、GPU、FPGA、DSP，以及最近幾年流行的各類AI加速器等;
- OpenCL包含兩部分：
1)用于編寫運(yùn)行在OpenCL device上的kernels的語言(基于C99);
2)OpenCL API，至于Runtime的實(shí)現(xiàn)交由各個(gè)廠家，比如Intel發(fā)布的opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz
以人工智能場景為例來理解一下，假如在某個(gè)AI芯片上跑人臉識(shí)別應(yīng)用，CPU擅長控制，AI processor擅長計(jì)算，軟件的flow就可以進(jìn)行拆分，用CPU來負(fù)責(zé)控制視頻流輸入輸出前后處理，AI processor來完成深度學(xué)習(xí)模型運(yùn)算完成識(shí)別，這就是一個(gè)典型的異構(gòu)處理場景，如果該AI芯片的SDK支持OpenCL，那么上層的軟件就可以基于OpenCL進(jìn)行開發(fā)了。
話不多說，看看OpenCL的架構(gòu)吧。
2. OpenCL架構(gòu)
OpenCL架構(gòu)，可以從平臺(tái)模型、內(nèi)存模型、執(zhí)行模型、編程模型四個(gè)角度來展開。
2.1 Platform Model
平臺(tái)模型：硬件拓?fù)潢P(guān)系的抽象描述
- 平臺(tái)模型由一個(gè)Host連接一個(gè)或多個(gè)OpenCL Devices組成;
- OpenCL Device，可以劃分成一個(gè)或多個(gè)計(jì)算單元Compute Unit(CU);
- CU可以進(jìn)一步劃分成一個(gè)或多個(gè)處理單元Processing Unit(PE)，最終的計(jì)算由PE來完成;
- OpenCL應(yīng)用程序分成兩部分：host代碼和device kernel代碼，其中Host運(yùn)行host代碼，并將kernel代碼以命令的方式提交到OpenCL devices，由OpenCL device來運(yùn)行kernel代碼;
2.2 Execution Model
執(zhí)行模型：Host如何利用OpenCL Device的計(jì)算資源完成高效的計(jì)算處理過程
Context
OpenCL的Execution Model由兩個(gè)不同的執(zhí)行單元定義：1)運(yùn)行在OpenCL設(shè)備上的kernel;2)運(yùn)行在Host上的Host program;其中，OpenCL使用Context代表kernel的執(zhí)行環(huán)境：
Context包含以下資源：
- Devices：一個(gè)或多個(gè)OpenCL設(shè)備;
- Kernel Objects：OpenCL Device的執(zhí)行函數(shù)及相關(guān)的參數(shù)值，通常定義在cl文件中;
- Program Objects：實(shí)現(xiàn)kernel的源代碼和可執(zhí)行程序，每個(gè)program可以包含多個(gè)kernel;
- Memory Objects：Host和OpenCL設(shè)備可見的變量，kernel執(zhí)行時(shí)對(duì)其進(jìn)行操作;
NDrange
- kernel是Execution Model的核心，放置在設(shè)備上執(zhí)行，當(dāng)kernel執(zhí)行前，需要?jiǎng)?chuàng)建一個(gè)索引空間NDRange(一維/二維/三維);
- 執(zhí)行kernel實(shí)例的稱為work-item，work-item組織成work-group，work-group組織成NDRange，最終將NDRange映射到OpenCL Device的計(jì)算單元上;
有兩種方式來找到work-item：
- 通過work-item的全局索引;
- 先查找到所在work-group的索引號(hào)，再根據(jù)局部索引號(hào)確定;
以一維為例：
- 上圖中總共有四個(gè)work-group，每個(gè)work-group包含四個(gè)work-item，所以local_size的大小為4，而local_id都是從0開始重新計(jì)數(shù);
- global_size代表總體的大小，也就是16個(gè)work-item，而global_id則是從0開始計(jì)數(shù);
以二維為例：
- 二維的計(jì)算方式與一維類似，也是結(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：單個(gè)CU中的本地內(nèi)存，Host看不到該區(qū)域并無法對(duì)其操作，該區(qū)域允許內(nèi)部的PE進(jìn)行讀寫操作，也可以用于PE之間的共享，需要注意同步和并發(fā)問題;
- Private Memory：PE的私有內(nèi)存，Host與PE之間都無法看到該區(qū)域;
2.4 Programming Model
- 在編程模型中，有兩部分代碼需要編寫：一部分是Host端，一部分是OpenCL Device端;
- 編程過程中，核心是要維護(hù)一個(gè)Context，代表了整個(gè)Kernel執(zhí)行的環(huán)境;
- 從cl源代碼中創(chuàng)建Program對(duì)象并編譯，在運(yùn)行時(shí)創(chuàng)建Kernel對(duì)象以及內(nèi)存對(duì)象，設(shè)置好相關(guān)的參數(shù)和輸入之后，就可以將Kernel送入到隊(duì)列中執(zhí)行，也就是Launch kernel的流程;
- 最終等待運(yùn)算結(jié)束，獲取計(jì)算結(jié)果即可;
3. 編程流程
- 上圖為一個(gè)OpenCL應(yīng)用開發(fā)涉及的基本過程;
下邊來一個(gè)實(shí)際的代碼測(cè)試跑跑，Talk is cheap, show me the code!
4. 示例代碼
- 測(cè)試環(huán)境：Ubuntu16.04，安裝Intel CPU OpenCL SDK(opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz);
- 為了簡化流程，示例代碼都不做容錯(cuò)處理，僅保留關(guān)鍵的操作;
- 整個(gè)代碼的功能是完成向量的加法操作;
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對(duì)象時(shí)會(huì)去讀取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
責(zé)任編輯：武曉燕
來源：
LoyenWang
