以下文章來源于淘系技術，作者恬步

MNN(Mobile Neural Network)是一個高性能、通用的深度學習框架，支持在移動端、PC端、服務端、嵌入式等各種設備上高效運行。MNN利用設備的GPU能力，全面充分“榨干”設備的GPU資源，來進行深度學習的高性能部署與訓練。

概述

MNN自開源以來，一直以高性能、通用性、易用性等特性聞名于業界。近一年來，MNN GPU再發力，OpenCL后端針對移動端(Adreno/Mali GPU)、PC端性能總體提升超過100%，部分機型性能翻幾番。請看下圖：

移動端GPU受限于芯片面積、能耗以及成本等因素，通常需要在IO帶寬和運算資源上嚴格受控壓縮。GPU硬件多樣性、OpenCL本身支持力度取決于設備生產商，這導致了不同設備GPU軟硬件層面的差異化。那么，如何在有限且碎片化的資源下充分發揮出硬件的性能優勢，加快深度學習模型推理速度，給MNN GPU平臺通用性與高性能兼備的定位帶來很高的挑戰。

是什么優化手段能在高端手機、低端千元機、服務端顯卡上，模型推理都能加速如此之多呢？MNN OpenCL在最近一年內做了什么優化呢？且看下文~

內存訪問效率角度

? 內存對象多元化調優

OpenCL提供兩種內存對象接口，Buffer-object和Image-object。有的GPU廠商對Image-object具有更快速的訪問和支持，高通Adreno GPU用戶手冊明確建議使用Image-object進行存儲更有利于訪問效率；但是ARM Mali GPU系列眾多，有不少機型對Image-object內存支持沒有那么好，適合更通用的Buffer-object。面對不同GPU內存存儲格式引起的不同訪問效率，MNN支持多元的內存存儲格式，以應對差異化的GPU型號。

下表為OpenCL Imgae-object與Buffer-object的區別對比。

高通平臺明確給出下圖數據流路結構圖，當使用Image-object存儲時，在進去讀數據時候可以使用Texture處理器和L1 cache進行快速讀取。如果使用Buffer-object的話將無法使用該硬件緩存優勢。所以對于高通驍龍系列手機，使用其GPU資源，選擇OpenCL Image-object性能更優。

Buffer-object也有其使用范圍和優勢。首先，Image-object的2D/3D內存每個維度申請尺寸是有上限的(不同硬件不同)，當要申請的內存某個維度的尺寸超過硬件支持的上限時，Image-object內存會申請失敗，這個時候只能使用Buffer-object。其次，Buffer-object內存排布是線性的，排布緊密，cache miss較友好，Image-object如果第一維度尺寸太小容易造成很嚴重的cache miss。ARM Mali GPU并沒有很友好的Texture處理器和L1 Cache(官方未提及)，官方開源項目采用Buffer-object存儲模式，可見目前Buffer-object對ARM Mali GPU更有優勢。

MNN OpenCL2020年之前已經支持Image-object，今年新增了對Buffer-object內存對象格式的支持，針對不同硬件平臺和算法模型，MNN OpenCL目前框架內支持根據實驗經驗化，在ARM-Mali GPU上采用Buffer內存，其他GPU型號采用Image內存。目前這個策略只是比較粗糙的經驗化手段。為了精確的內存選擇，用戶可以通過提前試跑兩種內存設置，來得到兩種內存模式推理性能的更優者，在實際推理過程中設置此內存模式即可。

? 內存對齊優化

下圖是CPU和GPU擁有的硬件資源(運算單元ALU/控制單元/緩存等)的示意圖。

CPU配置強大的控制單元和緩存機制，具有厲害的分支預測能力，擅長處理復雜邏輯。GPU硬件資源大量運算單元，控制單元和緩存相對來說較薄弱，使得分支預測能力較弱，所以在編程過程中需要盡量避免邏輯分支。

通常某個維度的并行量不會太大，在MNN OpenCL實現中，寬方向最高并行量是4，在申請內存時可以將該維度4對齊向上取整，這樣可以避免在讀取數據時為了防止讀越界而帶來的邊界判斷需求，從而減少GPU kernel內部分支。同時在遇到部分實在無法避免的分支判斷時，盡可能選擇三目運算符替代if分支。

? local memory并行歸約優化

歸約是一種并行算法，對于傳入進來N個輸入，通過二元操作，得到一個輸出值。典型的就是求最值、取平均值、求和等操作。以取最大值為例，傳統的串行算法，實現簡單，需要N次迭代運算操作。通常如下：

float maxValue = -MAXFLOAT;
for (int i = 0; i < N; i++) {
    maxValue = max(maxValue, array[i]);
}

下圖二分法并行歸約算法示意圖，每個步驟可以并行去求最大值，在opencl實現中，將放在同一個工作組的線程采用局部內存(local memory)進行存儲，因為同一工作組中的線程訪問共享數據時，local memory由于其得天獨厚的物理設計，效率遠高于global memory。二分歸約只需要logN次迭代操作。相較于傳統串行方法時間復雜度有降維優勢。對于歸約數目N較時，該方法性能提升明顯；但是N較少時使用串行方法即可，因為線程間barrier開銷會明顯蓋過local memory和操作次數帶來的優勢。

如下是opencl歸約算法kernel代碼示例：

const int idx = get_local_id(0);
FLOAT local sum[256];
sum[idx] = -MAXFLOAT;
const int reduce_num = get_local_size(0);//獲取工作組中的線程數量


for (int h = idx; h < total_num; h+=reduce_num) {//將多個工作組的值映射到當前工作組
    FLOAT in = read_input_data(input);
    sum[idx] = max(sum[idx], in);
}


barrier(CLK_LOCAL_MEM_FENCE);
for(int i = reduce_num/2; i > 0; i /= 2){//對當前工作組進行二分歸約運算
    if (idx < i)
        sum[idx] = max(sum[idx], sum[idx + i]);
    barrier(CLK_LOCAL_MEM_FENCE);
}
if (idx == 0) {
    write_output_data(output, sum[0]);//將sum[0]的值寫入輸出地址處
}

GPU計算分塊角度

? 工作組大小選擇

下圖示意的是一個GPU任務分塊執行示意圖。左側NDRange size表示的是任務擁有的總子任務(子線程)數量。子線程會被組織成一系列work-group分塊線程，每個work-group分塊會被分配到一個SP上面執行。這個執行過程是GPU SIMT架構的必然映射。

work-group大小的劃分會影響到整個GPU硬件資源的利用情況。針對一個特定的任務，最合適的work-group大小受單個線程需要完成的任務量、機型GPU硬件資源的強弱以及總線程數目等很多綜合因素影響，選擇不合適的work-group會對計算效率性產生不好的影響。

OpenCL框架允許編程者不去設定這個尺寸大小，會根據情況自行調度決定。但這往往不能帶來較好的性能效果。通常較合適的work-group size是NDRange size的因子或者是2的冪次方，這種極簡的設置往往可以帶來“還不錯”的性能。為了極佳的性能，MNN在預推理階段會選取多組work-group分塊大小進行Auto-Tuning試跑，選出性能最佳的work-group size，以此配置在實際推理中應用。

? 數據分塊復用方案

對于CNN網絡中常見的二維卷積運算，實際上是三維張量到三維張量的映射操作。拿kernel為3x3，stride為1x1，pad為1x1，dilate為1x1的卷積運算(暫不考慮偏置的情況)舉例，對于

輸入維度是Cin*H*W

權重維度是Cout*Cin*Kh*Kw,（3x3卷積核，Kw=Kh=3）

輸出維度則是Cout*H*W

下表給出了，單個線程不同粒度計算量對應的總的計算復雜度和內存訪問量大小。可以看出總的計算復雜度都是恒定的，但是隨著單個線程計算量的則增加，數據可復用力度越高，總的內存訪問量將會越少。拿單個像素輸出粒度為基準，單個線程輸出連續4個通道和連續4個寬方向的16個像素時，內存訪問量將減少75%。

增大數據分塊可以在計算復雜度不變的前提下，有效的降低數據內存訪問次數，對于提升性能有很重要的作用。但是，隨著單個線程GPU kernel計算量越大，需要使用的寄存器資源越多，全局工作項數目也將對應減少。單個線程計算量過大時必然會導致寄存器等資源不夠用，也可能導致全局工作項數目過小起不到較佳的GPU任務發射并行度。

最適宜的數據分塊量，會隨著總線程數量/硬件平臺寄存器/硬件ALU數目/IO帶寬等資源以及單個線程的計算量大小的不同，也會有較大的差別。譬如，當總線程數目較少時，此時為了足夠的并發量，數據分塊大小需要相對應盡可能減小；當總線程數/寄存器資源足夠的時候，可以考慮加大數據分塊大小，來增大數據復用的優勢。

當前MNN OpenCL針對核心算子支持多種數據復用分塊量，支持在預推理階段提前Auto-Tuning試跑找到最適宜當前設備/計算規模的數據分塊量，以獲得最佳的優化性能。

異構調度角度

異構系統的調度，相比于同構系統調度會復雜一些，因為會涉及到主機端與設備端交互部分。下圖是一個典型的OpenCL異構系統調度圖。主要包含了三個部分。其一：主機CPU端，負責整個異構系統的主控調度，包括資源的申請調配、任務的發射等；其二：各種異構設備端（比如GPU/DSP/FPGA等），是異構系統的核心處理器；其三：OpenCL kernel代碼，負責對異構處理器進行操縱。

OpenCL kernel從CPU提交到任務隊列后會經歷Queued/Submitted/Ready/Running等整個執行狀態。和CPU算子執行存在明顯的差別，GPU算子執行需要統一入隊，然后GPU會對同一個任務隊列上的任務進行隊列式排隊。對于某個算子從入隊等待到真正執行是有延遲的。具體延遲情況依賴廠商平臺調度。

從Queued到Submitted狀態之間的軟件開銷/CPU cache開銷，調度好的系統能盡量最優化處理調度。但是，由于不同機型，對OpenCL的任務調度不盡如人意，往往這段時延較大。OpenCL提供flush接口，可以在一定kernel量的時候手動加速提交任務，在調度上加入人為的動態隊列刷新機制。

在華為系列手機(Mali GPU)，需要在kernel累計量較少的時候就需要人為加入刷新機制，整體性能提升可觀，對人為刷新機制依賴比較嚴重。高通系列手機，在kernel累計量較多的時候加入刷新機制即可，不依賴人為刷新機制，整體系統調度情況較好。MNN根據實驗經驗公式，針對不同機型調優出動態的命令隊列刷新機制。

預推理Auto-tuning調優

GPU種類眾多，不同廠家設計差異大，相同廠家GPU設計更新變遷復雜。這一系列導致的GPU碎片化，使得不同機型最優的算子實現都不一樣。很可能會出現某些機型算子實現了最優化，其他部分機型上出現負優化的情況。這給MNN GPU平臺通用且高性能兼備的定位帶來很大的挑戰。要想使得全機型、多模型性能都能達到很優，Auto-tuning試跑是一個很有效的方式。

但是，Auto-tuning必然會增加不少額外的試跑耗時。MNN之前已經支持了“預推理”機制，核心目的是：在真正推理之前，將推理過程中需要的內存/任務準備與分發等提前推理出來，從而優化降低實際推理過程中的耗時。功能主要包括：

進行內存管理：申請每個算子的輸入輸出Tensor內存與運算時所需的緩存。
任務準備與分發：對CPU來說，可以在這個環節生成Lambda函數。對GPU來說，可以制作相關算子的命令緩沖（Command Buffer），填充參數等等。

MNN的OpenCL后端擴充新增“預推理”的功能——任務Auto-tuning試跑，找出最優的計算配置方案。增加此功能主要基于以下考慮：

每個模型的算子固定，每次推理只是算子的輸入數據不同，計算方式和計算量完全一致。最優的實現方式一致，可以在“預推理”階段提前Auto-tuning出每個算子的最優配置。
“預推理”機制可以有效降低推理階段的耗時。

MNN OpenCL增加此“預推理”功能后，在推理階段，可以直接使用“預推理”出的Auto-tuning計算配置，獲得優化的性能。目前支持了工作組大小選擇和數據分塊復用兩種優化策略的Auto-tuning試跑，具體內容在GPU分塊角度優化模塊里已經闡述過。

GPU業務落地設計與建議

? 用戶可配置的Gpu-Mode

上述通過經驗實驗方法與Auto-Tuning試跑的方式來提升推理速度。由于經驗實驗公式不可能覆蓋所有情況，支持增加接口給用戶自行選擇。目前GPU內存對象選擇提供開放選項給用戶自行配置。Auto-Tuning試跑會增加“預推理”的耗時，MNN OpenCL提供不同的Auto-tuning力度選項可供用戶選擇。用戶可選取性能滿足要求的前提下盡可能縮減Auto-tuning的力度。

MNN OpenCL提供用戶可自行配置的MNNGpuMode，具體選項如下圖所示。

指定需要使用的Tuning-mode和Gpu-Memory類型，在代碼中設置config的mode設置即可。代碼示例如下：

MNN::ScheduleConfig config;
config.mode = MNN_GPU_TUNING_NORMAL | MNN_GPU_MEMORY_IMAGE;

通常如果介意“預推理”耗時較長可以選取較低level的Tuning-mode(下面也會介紹Cache機制解決初始化耗時長的問題)。Gpu-Memory用戶可以Buffer模式和Image模式都自行設置一次，選擇推理速度較優的模式，當然如果不設置的話框架會根據機型進行自動選擇(不可能保證所有情況下都最優)。

? MNN Cache技術設計

由于OpenCL kernel需要根據不同機型在線編譯源碼program，以及加入精細化調優Auto-tuning試跑機制后，獲得極佳性能的同時會帶來啟動時間較慢的代價。很多情況下，用戶對于初始化時間不太能接受，導致很多業務難以真正落地。

為了優化GPU初始化時間，MNN將當前機型編譯好的program轉成二進制、Auto-tuning出的最佳配置進行記錄，并存儲成Cache文件。之后初始化的時候加載Cache文件讀取二進制版program(無需編譯源碼)和tune好的配置信息(無需再次Auto-tuning)，從而大大提高初始化速度。

如果應用僅限某種或某幾種特定機型，可以事先生成好該機型的cache文件。這樣實際啟動的時候就可以直接加載cache文件，享受快速的啟動速度。如果應用機型太多，不能接受每種機型都事先提前生成好Cache文件，可以考慮在調度上可以事先提前初始化，在跑其他應用的時候就提前初始化生成Cache文件，在調度上“隱藏”掉生成Cache的時間。

MNN Cache使用上極其方便簡潔，用戶只需要調用一行代碼，接口如下：

setCacheFile(const std::string& fileName, size_t keySize = 128)//keySize: 使用模型Buffer的前 keySize 個 byte 作為校驗

下表給出小米6(Adreno-540 GPU)設備上，在使用Cache前后不同tuning-mode時OpenCL總初始化耗時。可以看到使用Cache能極大優化啟動速度，助力業務落地。

? 用戶透明的性能分析

用戶在設計了一個模型后，使用MNN benchmark工具測試GPU性能，當遇到不太符合預期的性能時，這個時候用戶就像是用了個盲盒似的，無從分析，無從下手。在此需求下，MNN GPU提供了性能熱點分析工具，幫助用戶定位性能熱點。

性能熱點分析對于提升總體推理性能的重要性，就好比生病了去醫院需要抽血分析各項指標報告一樣，只有知道各處的指標詳情才好對癥下藥。MNN GPU提供統計OpenCL kernel耗時方法，使用opencl event進行GPU端計時統計，可以精確地進行單個kernel耗時分析，準確性能明顯高于CPU端計時器。

MNN OpenCL用戶可以在編譯庫的時候打開MNN_OPENCL_PROFILE宏，運行程序可以看到每個部分的耗時，去進行性能熱點定位分析。下圖給出的是部分算子耗時圖，可以看到各個算子的耗時情況，可以看出第一個Conv2D算子耗時是絕對的熱點。用戶可以為了提升性能，對這個卷積算子考慮采用減小通道數或者使用多個小卷積核代替一個大卷積核方式。提供這樣的Profile功能，對模型設計帶來更多指導性靈感和參考，提供了結合框架去設計模型的可能性。

? 適合GPU加速模型設計建議

經常會有用戶反饋，為什么使用GPU加速反而性能不如運行CPU上呢？通常在用戶潛意識里，GPU總比CPU快。其實這是個誤解。GPU的硬件結構設計特性，決定了GPU對具有大量可并行的運算才更有優勢。對于運算量過小或者并行度較低的模型，通常GPU上運行效率不如CPU。

對于CPU耗時本身較少的小模型(如幾個ms)，不建議使用GPU加速。因為GPU運行啟動調度本身需要一定耗時，其次CPU/GPU數據拷貝耗時，加之不符合GPU適合大量運算的特性。因此，模型太小選擇CPU就好。在移動端和PC端，要用GPU加速，模型設計方面要盡可能設計一些并行量大的高速模型。具體給出以下幾點建議：