移動端異構運算技術-GPU OpenCL 程式設計（基礎篇）

全文3829字，預計閱讀時間10分鐘。

一、前言

随着移動端晶片性能的不斷提升，在移動端上實時進行計算機圖形學、深度學習模型推理等計算密集型任務不再是一個奢望。在移動端裝置上，GPU 憑借其優秀的浮點運算性能，以及良好的 API 相容性，成為移動端異構計算中非常重要的計算單元。現階段，在 Android 裝置市場，高通 Adreno 和華為Mali已經占據了手機 GPU 晶片的主要份額，二者均提供了強勁的 GPU 運算能力。OpenCL，作為 Android 的系統庫，在兩個晶片上均得到良好的支援。

目前，百度APP 已經将 GPU 計算加速手段，應用在深度模型推理及一些計算密集型業務上，本文将介紹 OpenCL 基礎概念與簡單的 OpenCL 程式設計。

（注：Apple 對于 GPU 推薦的使用方式是 Metal，此處暫不做展開）

二、基礎概念

2.1 異構計算

異構計算（Heterogeneous Computing），主要是指使用不同類型指令集和體系架構的計算單元組成系統的計算方式。常見的計算單元類别包括 CPU、GPU 等協處理器、DSP、ASIC、FPGA 等。

2.2 GPU

GPU（Graphics Processing Unit），圖形處理器，又稱顯示核心、顯示卡、視覺處理器、顯示晶片或繪圖晶片，是一種專門在個人電腦、工作站、遊戲機和一些移動裝置（如平闆電腦、智能手機等）上執行繪圖運算工作的微處理器。傳統方式中提升 CPU 時鐘頻率和核心數量而提高計算能力的方式已經遇到了散熱以及能耗的瓶頸。雖然 GPU 單個計算單元的工作頻率較低，卻具備更多的核心數及并行計算能力。相比于 CPU，GPU 的總體性能-晶片面積比，性能-功耗比都更高。

三、OpenCL

OpenCL（Open Computing Language）是一個由非盈利性技術組織 Khronos Group 掌管的異構平台程式設計架構，支援的異構平台涵蓋 CPU、GPU、DSP、FPGA 以及其他類型的處理器與硬體加速器。OpenCL 主要包含兩部分，一部分是一種基于 C99 标準用于編寫核心的語言，另一部分是定義并控制平台的API。

OpenCL 類似于另外兩個開放的工業标準 OpenGL 和 OpenAL ，二者分别用于三維圖形和計算機音頻方面。OpenCL 主要擴充了 GPU 圖形生成之外的計算能力。

3.1 OpenCL 程式設計模型

使用 OpenCL 程式設計需要了解 OpenCL 程式設計的三個核心模型，OpenCL 平台、執行和記憶體模型。

平台模型（Platform Model）

Platform 代表 OpenCL 視角上的系統中各計算資源之間的拓撲聯系。對于Android 裝置，Host 即是 CPU。每個 GPU計算裝置（Compute Device）均包含了多個計算單元（Compute Unit），每個計算單元包含多個處理元素（Processing Element）。對于 GPU 而言，計算單元和處理元素就是 GPU 内的流式多處理器。

執行模型 （Execution Model）

通過 OpenCL 的 clEnqueueNDRangeKernel 指令，可以啟動預編譯好的 OpenCL 核心，OpenCL 架構上可以支援N維的資料并行處理。以二維圖檔為例，如果将圖檔的寬高作為 NDRange，在 OpenCL 的核心中可以把圖檔的每個像素放在一個處理元素上執行，借此可以達到并行化執行的目地。

從上面平台模型部分可以知道，為了提高執行效率，處理器通常會将處理元素配置設定到執行單元中。我們可以在 clEnqueueNDRangeKernel 中指定工作組大小。同一個工作組中的工作項可以共享本地記憶體，可以使用屏障（Barriers）去進行同步，也可以通過特定的工作組函數（比如async_work_group_copy）來進行協作。

記憶體模型 （Memory Model）

下圖中描述了 OpenCL 的記憶體結構：

• 宿主記憶體（Host Memory）：宿主 CPU 可直接通路的記憶體。
• 全局/常量記憶體 （Global/Constant Memory）：可以用于計算裝置中的所有計算單元。
• 本地記憶體（Local Memory）：對計算單元中的所有處理元素可用。
• 私有記憶體（Private Memory）：用于單個處理元素。

3.2 OpenCL 程式設計

OpenCL 的程式設計實際應用中需要一些工程化的封裝，本文僅以兩個數組相加作為舉例，并提供一個簡單的示例代碼作為參考 ARRAY_ADD_SAMPLE (https://github.com/xiebaiyuan/opencl_cook/blob/master/array_add/array_add.cpp)。

本文将用此作為示例，來闡述 OpenCL 的工作流程。

OpenCL 整體流程主要分為以下幾個步驟：

初始化 OpenCL 相關環境，如 cl_device、cl_context、cl_command_queue 等

cl_int status;
// init device
    runtime.device = init_device();
// create context
    runtime.context = clCreateContext(nullptr, 1, &runtime.device, nullptr, nullptr, &status);
// create queue
    runtime.queue = clCreateCommandQueue(runtime.context, runtime.device, 0, &status);
           

初始化程式要執行的 program、kernel

cl_int status;
    // init program
    runtime.program = build_program(runtime.context, runtime.device, PROGRAM_FILE);
    // create kernel
    runtime.kernel = clCreateKernel(runtime.program, KERNEL_FUNC, &status);
           

準備輸入輸出，設定到 CLKernel

// init datas 
    float input_data[ARRAY_SIZE];
    float bias_data[ARRAY_SIZE];
    float output_data[ARRAY_SIZE];
    for (int i = 0; i < ARRAY_SIZE; i++) {
        input_data[i] = 1.f * (float) i;
        bias_data[i] = 10000.f;
    }
    // create buffers
    runtime.input_buffer = clCreateBuffer(runtime.context, CL_MEM_READ_ONLY |
        CL_MEM_COPY_HOST_PTR, ARRAY_SIZE * sizeof(float), input_data, &status);
    runtime.bias_buffer = clCreateBuffer(runtime.context, CL_MEM_READ_ONLY |
        CL_MEM_COPY_HOST_PTR, ARRAY_SIZE * sizeof(float), bias_data, &status);
    runtime.output_buffer = clCreateBuffer(runtime.context, CL_MEM_READ_ONLY |
        CL_MEM_COPY_HOST_PTR, ARRAY_SIZE * sizeof(float), output_data, &status);
    // config cl args
    status = clSetKernelArg(runtime.kernel, 0, sizeof(cl_mem), &runtime.input_buffer);
    status |= clSetKernelArg(runtime.kernel, 1, sizeof(cl_mem), &runtime.bias_buffer);
    status |= clSetKernelArg(runtime.kernel, 2, sizeof(cl_mem), &runtime.output_buffer);
           

執行擷取結果

// clEnqueueNDRangeKernel
    status = clEnqueueNDRangeKernel(runtime.queue, runtime.kernel, 1, nullptr, &ARRAY_SIZE,
                                    nullptr, 0, nullptr, nullptr);
    // read from output
    status = clEnqueueReadBuffer(runtime.queue, runtime.output_buffer, CL_TRUE, 0,
                                 sizeof(output_data), output_data, 0, nullptr, nullptr);
    // do with output_data
    ...
           

‍

四、總結

随着 CPU 瓶頸的到來，GPU 或者其他專用計算裝置的程式設計将是未來的一個重要的技術方向，本文僅介紹 OpenCL 程式設計基礎。可以關注 百度APP技術公衆号，獲得後續《移動端異構運算技術》文章更新。

五、參考資料

1.  OpenCL-Guide

https://github.com/KhronosGroup/OpenCL-Guide/blob/main/chapters/opencl_programming_model.md

2.  OpenCL-Examples

https://github.com/rsnemmen/OpenCL-examples

推薦閱讀：

雲原生賦能開發測試

基于Saga的分布式事務排程落地