以下文章來(lái)源于淘系技術(shù) ，作者恬步

MNN(Mobile Neural Network)是一個(gè)高性能、通用的深度學(xué)習(xí)框架，支持在移動(dòng)端、PC端、服務(wù)端、嵌入式等各種設(shè)備上高效運(yùn)行。MNN利用設(shè)備的GPU能力，全面充分“榨干”設(shè)備的GPU資源，來(lái)進(jìn)行深度學(xué)習(xí)的高性能部署與訓(xùn)練。

概述

MNN自開(kāi)源以來(lái)，一直以高性能、通用性、易用性等特性聞名于業(yè)界。近一年來(lái)，MNN GPU再發(fā)力，OpenCL后端針對(duì)移動(dòng)端(Adreno/Mali GPU)、PC端性能總體提升超過(guò)100%，部分機(jī)型性能翻幾番。請(qǐng)看下圖：

移動(dòng)端GPU受限于芯片面積、能耗以及成本等因素，通常需要在IO帶寬和運(yùn)算資源上嚴(yán)格受控壓縮。GPU硬件多樣性、OpenCL本身支持力度取決于設(shè)備生產(chǎn)商，這導(dǎo)致了不同設(shè)備GPU軟硬件層面的差異化。那么，如何在有限且碎片化的資源下充分發(fā)揮出硬件的性能優(yōu)勢(shì)，加快深度學(xué)習(xí)模型推理速度，給MNN GPU平臺(tái)通用性與高性能兼?zhèn)涞亩ㄎ粠?lái)很高的挑戰(zhàn)。

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

內(nèi)存訪(fǎng)問(wèn)效率角度

? 內(nèi)存對(duì)象多元化調(diào)優(yōu)

OpenCL提供兩種內(nèi)存對(duì)象接口，Buffer-object和Image-object。有的GPU廠(chǎng)商對(duì)Image-object具有更快速的訪(fǎng)問(wèn)和支持，高通Adreno GPU用戶(hù)手冊(cè)明確建議使用Image-object進(jìn)行存儲(chǔ)更有利于訪(fǎng)問(wèn)效率；但是ARM Mali GPU系列眾多，有不少機(jī)型對(duì)Image-object內(nèi)存支持沒(méi)有那么好，適合更通用的Buffer-object。面對(duì)不同GPU內(nèi)存存儲(chǔ)格式引起的不同訪(fǎng)問(wèn)效率，MNN支持多元的內(nèi)存存儲(chǔ)格式，以應(yīng)對(duì)差異化的GPU型號(hào)。

下表為OpenCL Imgae-object與Buffer-object的區(qū)別對(duì)比。

高通平臺(tái)明確給出下圖數(shù)據(jù)流路結(jié)構(gòu)圖，當(dāng)使用Image-object存儲(chǔ)時(shí)，在進(jìn)去讀數(shù)據(jù)時(shí)候可以使用Texture處理器和L1 cache進(jìn)行快速讀取。如果使用Buffer-object的話(huà)將無(wú)法使用該硬件緩存優(yōu)勢(shì)。所以對(duì)于高通驍龍系列手機(jī)，使用其GPU資源，選擇OpenCL Image-object性能更優(yōu)。

Buffer-object也有其使用范圍和優(yōu)勢(shì)。首先，Image-object的2D/3D內(nèi)存每個(gè)維度申請(qǐng)尺寸是有上限的(不同硬件不同)，當(dāng)要申請(qǐng)的內(nèi)存某個(gè)維度的尺寸超過(guò)硬件支持的上限時(shí)，Image-object內(nèi)存會(huì)申請(qǐng)失敗，這個(gè)時(shí)候只能使用Buffer-object。其次，Buffer-object內(nèi)存排布是線(xiàn)性的，排布緊密，cache miss較友好，Image-object如果第一維度尺寸太小容易造成很?chē)?yán)重的cache miss。ARM Mali GPU并沒(méi)有很友好的Texture處理器和L1 Cache(官方未提及)，官方開(kāi)源項(xiàng)目采用Buffer-object存儲(chǔ)模式，可見(jiàn)目前Buffer-object對(duì)ARM Mali GPU更有優(yōu)勢(shì)。

MNN OpenCL2020年之前已經(jīng)支持Image-object，今年新增了對(duì)Buffer-object內(nèi)存對(duì)象格式的支持，針對(duì)不同硬件平臺(tái)和算法模型，MNN OpenCL目前框架內(nèi)支持根據(jù)實(shí)驗(yàn)經(jīng)驗(yàn)化，在A(yíng)RM-Mali GPU上采用Buffer內(nèi)存，其他GPU型號(hào)采用Image內(nèi)存。目前這個(gè)策略只是比較粗糙的經(jīng)驗(yàn)化手段。為了精確的內(nèi)存選擇，用戶(hù)可以通過(guò)提前試跑兩種內(nèi)存設(shè)置，來(lái)得到兩種內(nèi)存模式推理性能的更優(yōu)者，在實(shí)際推理過(guò)程中設(shè)置此內(nèi)存模式即可。

? 內(nèi)存對(duì)齊優(yōu)化

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

CPU配置強(qiáng)大的控制單元和緩存機(jī)制，具有厲害的分支預(yù)測(cè)能力，擅長(zhǎng)處理復(fù)雜邏輯。GPU硬件資源大量運(yùn)算單元，控制單元和緩存相對(duì)來(lái)說(shuō)較薄弱，使得分支預(yù)測(cè)能力較弱，所以在編程過(guò)程中需要盡量避免邏輯分支。

通常某個(gè)維度的并行量不會(huì)太大，在MNN OpenCL實(shí)現(xiàn)中，寬方向最高并行量是4，在申請(qǐng)內(nèi)存時(shí)可以將該維度4對(duì)齊向上取整，這樣可以避免在讀取數(shù)據(jù)時(shí)為了防止讀越界而帶來(lái)的邊界判斷需求，從而減少GPU kernel內(nèi)部分支。同時(shí)在遇到部分實(shí)在無(wú)法避免的分支判斷時(shí)，盡可能選擇三目運(yùn)算符替代if分支。

? local memory并行歸約優(yōu)化

歸約是一種并行算法，對(duì)于傳入進(jìn)來(lái)N個(gè)輸入，通過(guò)二元操作，得到一個(gè)輸出值。典型的就是求最值、取平均值、求和等操作。以取最大值為例，傳統(tǒng)的串行算法，實(shí)現(xiàn)簡(jiǎn)單，需要N次迭代運(yùn)算操作。通常如下：

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

下圖二分法并行歸約算法示意圖，每個(gè)步驟可以并行去求最大值，在opencl實(shí)現(xiàn)中，將放在同一個(gè)工作組的線(xiàn)程采用局部?jī)?nèi)存(local memory)進(jìn)行存儲(chǔ)，因?yàn)橥还ぷ鹘M中的線(xiàn)程訪(fǎng)問(wèn)共享數(shù)據(jù)時(shí)，local memory由于其得天獨(dú)厚的物理設(shè)計(jì)，效率遠(yuǎn)高于global memory。二分歸約只需要logN次迭代操作。相較于傳統(tǒng)串行方法時(shí)間復(fù)雜度有降維優(yōu)勢(shì)。對(duì)于歸約數(shù)目N較時(shí)，該方法性能提升明顯；但是N較少時(shí)使用串行方法即可，因?yàn)榫€(xiàn)程間barrier開(kāi)銷(xiāo)會(huì)明顯蓋過(guò)local memory和操作次數(shù)帶來(lái)的優(yōu)勢(shì)。

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

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


for (int h = idx; h < total_num; h+=reduce_num) {//將多個(gè)工作組的值映射到當(dāng)前工作組
    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){//對(duì)當(dāng)前工作組進(jìn)行二分歸約運(yùn)算
    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]的值寫(xiě)入輸出地址處
}

GPU計(jì)算分塊角度

? 工作組大小選擇

下圖示意的是一個(gè)GPU任務(wù)分塊執(zhí)行示意圖。左側(cè)NDRange size表示的是任務(wù)擁有的總子任務(wù)(子線(xiàn)程)數(shù)量。子線(xiàn)程會(huì)被組織成一系列work-group分塊線(xiàn)程，每個(gè)work-group分塊會(huì)被分配到一個(gè)SP上面執(zhí)行。這個(gè)執(zhí)行過(guò)程是GPU SIMT架構(gòu)的必然映射。

work-group大小的劃分會(huì)影響到整個(gè)GPU硬件資源的利用情況。針對(duì)一個(gè)特定的任務(wù)，最合適的work-group大小受單個(gè)線(xiàn)程需要完成的任務(wù)量、機(jī)型GPU硬件資源的強(qiáng)弱以及總線(xiàn)程數(shù)目等很多綜合因素影響，選擇不合適的work-group會(huì)對(duì)計(jì)算效率性產(chǎn)生不好的影響。

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

? 數(shù)據(jù)分塊復(fù)用方案

對(duì)于CNN網(wǎng)絡(luò)中常見(jiàn)的二維卷積運(yùn)算，實(shí)際上是三維張量到三維張量的映射操作。拿kernel為3x3，stride為1x1，pad為1x1，dilate為1x1的卷積運(yùn)算(暫不考慮偏置的情況)舉例，對(duì)于

輸入維度是Cin*H*W

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

輸出維度則是Cout*H*W

下表給出了，單個(gè)線(xiàn)程不同粒度計(jì)算量對(duì)應(yīng)的總的計(jì)算復(fù)雜度和內(nèi)存訪(fǎng)問(wèn)量大小。可以看出總的計(jì)算復(fù)雜度都是恒定的，但是隨著單個(gè)線(xiàn)程計(jì)算量的則增加，數(shù)據(jù)可復(fù)用力度越高，總的內(nèi)存訪(fǎng)問(wèn)量將會(huì)越少。拿單個(gè)像素輸出粒度為基準(zhǔn)，單個(gè)線(xiàn)程輸出連續(xù)4個(gè)通道和連續(xù)4個(gè)寬方向的16個(gè)像素時(shí)，內(nèi)存訪(fǎng)問(wèn)量將減少75%。

增大數(shù)據(jù)分塊可以在計(jì)算復(fù)雜度不變的前提下，有效的降低數(shù)據(jù)內(nèi)存訪(fǎng)問(wèn)次數(shù)，對(duì)于提升性能有很重要的作用。但是，隨著單個(gè)線(xiàn)程GPU kernel計(jì)算量越大，需要使用的寄存器資源越多，全局工作項(xiàng)數(shù)目也將對(duì)應(yīng)減少。單個(gè)線(xiàn)程計(jì)算量過(guò)大時(shí)必然會(huì)導(dǎo)致寄存器等資源不夠用，也可能導(dǎo)致全局工作項(xiàng)數(shù)目過(guò)小起不到較佳的GPU任務(wù)發(fā)射并行度。

最適宜的數(shù)據(jù)分塊量，會(huì)隨著總線(xiàn)程數(shù)量/硬件平臺(tái)寄存器/硬件ALU數(shù)目/IO帶寬等資源以及單個(gè)線(xiàn)程的計(jì)算量大小的不同，也會(huì)有較大的差別。譬如，當(dāng)總線(xiàn)程數(shù)目較少時(shí)，此時(shí)為了足夠的并發(fā)量，數(shù)據(jù)分塊大小需要相對(duì)應(yīng)盡可能減??；當(dāng)總線(xiàn)程數(shù)/寄存器資源足夠的時(shí)候，可以考慮加大數(shù)據(jù)分塊大小，來(lái)增大數(shù)據(jù)復(fù)用的優(yōu)勢(shì)。

當(dāng)前MNN OpenCL針對(duì)核心算子支持多種數(shù)據(jù)復(fù)用分塊量，支持在預(yù)推理階段提前Auto-Tuning試跑找到最適宜當(dāng)前設(shè)備/計(jì)算規(guī)模的數(shù)據(jù)分塊量，以獲得最佳的優(yōu)化性能。

異構(gòu)調(diào)度角度

異構(gòu)系統(tǒng)的調(diào)度，相比于同構(gòu)系統(tǒng)調(diào)度會(huì)復(fù)雜一些，因?yàn)闀?huì)涉及到主機(jī)端與設(shè)備端交互部分。下圖是一個(gè)典型的OpenCL異構(gòu)系統(tǒng)調(diào)度圖。主要包含了三個(gè)部分。其一：主機(jī)CPU端，負(fù)責(zé)整個(gè)異構(gòu)系統(tǒng)的主控調(diào)度，包括資源的申請(qǐng)調(diào)配、任務(wù)的發(fā)射等；其二：各種異構(gòu)設(shè)備端（比如GPU/DSP/FPGA等），是異構(gòu)系統(tǒng)的核心處理器；其三：OpenCL kernel代碼，負(fù)責(zé)對(duì)異構(gòu)處理器進(jìn)行操縱。

OpenCL kernel從CPU提交到任務(wù)隊(duì)列后會(huì)經(jīng)歷Queued/Submitted/Ready/Running等整個(gè)執(zhí)行狀態(tài)。和CPU算子執(zhí)行存在明顯的差別，GPU算子執(zhí)行需要統(tǒng)一入隊(duì)，然后GPU會(huì)對(duì)同一個(gè)任務(wù)隊(duì)列上的任務(wù)進(jìn)行隊(duì)列式排隊(duì)。對(duì)于某個(gè)算子從入隊(duì)等待到真正執(zhí)行是有延遲的。具體延遲情況依賴(lài)廠(chǎng)商平臺(tái)調(diào)度。

從Queued到Submitted狀態(tài)之間的軟件開(kāi)銷(xiāo)/CPU cache開(kāi)銷(xiāo)，調(diào)度好的系統(tǒng)能盡量最優(yōu)化處理調(diào)度。但是，由于不同機(jī)型，對(duì)OpenCL的任務(wù)調(diào)度不盡如人意，往往這段時(shí)延較大。OpenCL提供flush接口，可以在一定kernel量的時(shí)候手動(dòng)加速提交任務(wù)，在調(diào)度上加入人為的動(dòng)態(tài)隊(duì)列刷新機(jī)制。

在華為系列手機(jī)(Mali GPU)，需要在kernel累計(jì)量較少的時(shí)候就需要人為加入刷新機(jī)制，整體性能提升可觀(guān)，對(duì)人為刷新機(jī)制依賴(lài)比較嚴(yán)重。高通系列手機(jī)，在kernel累計(jì)量較多的時(shí)候加入刷新機(jī)制即可，不依賴(lài)人為刷新機(jī)制，整體系統(tǒng)調(diào)度情況較好。MNN根據(jù)實(shí)驗(yàn)經(jīng)驗(yàn)公式，針對(duì)不同機(jī)型調(diào)優(yōu)出動(dòng)態(tài)的命令隊(duì)列刷新機(jī)制。

預(yù)推理Auto-tuning調(diào)優(yōu)

GPU種類(lèi)眾多，不同廠(chǎng)家設(shè)計(jì)差異大，相同廠(chǎng)家GPU設(shè)計(jì)更新變遷復(fù)雜。這一系列導(dǎo)致的GPU碎片化，使得不同機(jī)型最優(yōu)的算子實(shí)現(xiàn)都不一樣。很可能會(huì)出現(xiàn)某些機(jī)型算子實(shí)現(xiàn)了最優(yōu)化，其他部分機(jī)型上出現(xiàn)負(fù)優(yōu)化的情況。這給MNN GPU平臺(tái)通用且高性能兼?zhèn)涞亩ㄎ粠?lái)很大的挑戰(zhàn)。要想使得全機(jī)型、多模型性能都能達(dá)到很優(yōu)，Auto-tuning試跑是一個(gè)很有效的方式。

但是，Auto-tuning必然會(huì)增加不少額外的試跑耗時(shí)。MNN之前已經(jīng)支持了“預(yù)推理”機(jī)制，核心目的是：在真正推理之前，將推理過(guò)程中需要的內(nèi)存/任務(wù)準(zhǔn)備與分發(fā)等提前推理出來(lái)，從而優(yōu)化降低實(shí)際推理過(guò)程中的耗時(shí)。功能主要包括：

進(jìn)行內(nèi)存管理：申請(qǐng)每個(gè)算子的輸入輸出Tensor內(nèi)存與運(yùn)算時(shí)所需的緩存。
任務(wù)準(zhǔn)備與分發(fā)：對(duì)CPU來(lái)說(shuō)，可以在這個(gè)環(huán)節(jié)生成Lambda函數(shù)。對(duì)GPU來(lái)說(shuō)，可以制作相關(guān)算子的命令緩沖（Command Buffer），填充參數(shù)等等。

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

每個(gè)模型的算子固定，每次推理只是算子的輸入數(shù)據(jù)不同，計(jì)算方式和計(jì)算量完全一致。最優(yōu)的實(shí)現(xiàn)方式一致，可以在“預(yù)推理”階段提前Auto-tuning出每個(gè)算子的最優(yōu)配置。
“預(yù)推理”機(jī)制可以有效降低推理階段的耗時(shí)。

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

GPU業(yè)務(wù)落地設(shè)計(jì)與建議

? 用戶(hù)可配置的Gpu-Mode

上述通過(guò)經(jīng)驗(yàn)實(shí)驗(yàn)方法與Auto-Tuning試跑的方式來(lái)提升推理速度。由于經(jīng)驗(yàn)實(shí)驗(yàn)公式不可能覆蓋所有情況，支持增加接口給用戶(hù)自行選擇。目前GPU內(nèi)存對(duì)象選擇提供開(kāi)放選項(xiàng)給用戶(hù)自行配置。Auto-Tuning試跑會(huì)增加“預(yù)推理”的耗時(shí)，MNN OpenCL提供不同的Auto-tuning力度選項(xiàng)可供用戶(hù)選擇。用戶(hù)可選取性能滿(mǎn)足要求的前提下盡可能縮減Auto-tuning的力度。

MNN OpenCL提供用戶(hù)可自行配置的MNNGpuMode，具體選項(xiàng)如下圖所示。

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

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

通常如果介意“預(yù)推理”耗時(shí)較長(zhǎng)可以選取較低level的Tuning-mode(下面也會(huì)介紹Cache機(jī)制解決初始化耗時(shí)長(zhǎng)的問(wèn)題)。Gpu-Memory用戶(hù)可以Buffer模式和Image模式都自行設(shè)置一次，選擇推理速度較優(yōu)的模式，當(dāng)然如果不設(shè)置的話(huà)框架會(huì)根據(jù)機(jī)型進(jìn)行自動(dòng)選擇(不可能保證所有情況下都最優(yōu))。

? MNN Cache技術(shù)設(shè)計(jì)

由于OpenCL kernel需要根據(jù)不同機(jī)型在線(xiàn)編譯源碼program，以及加入精細(xì)化調(diào)優(yōu)Auto-tuning試跑機(jī)制后，獲得極佳性能的同時(shí)會(huì)帶來(lái)啟動(dòng)時(shí)間較慢的代價(jià)。很多情況下，用戶(hù)對(duì)于初始化時(shí)間不太能接受，導(dǎo)致很多業(yè)務(wù)難以真正落地。

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

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

MNN Cache使用上極其方便簡(jiǎn)潔，用戶(hù)只需要調(diào)用一行代碼，接口如下：

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

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

? 用戶(hù)透明的性能分析

用戶(hù)在設(shè)計(jì)了一個(gè)模型后，使用MNN benchmark工具測(cè)試GPU性能，當(dāng)遇到不太符合預(yù)期的性能時(shí)，這個(gè)時(shí)候用戶(hù)就像是用了個(gè)盲盒似的，無(wú)從分析，無(wú)從下手。在此需求下，MNN GPU提供了性能熱點(diǎn)分析工具，幫助用戶(hù)定位性能熱點(diǎn)。

性能熱點(diǎn)分析對(duì)于提升總體推理性能的重要性，就好比生病了去醫(yī)院需要抽血分析各項(xiàng)指標(biāo)報(bào)告一樣，只有知道各處的指標(biāo)詳情才好對(duì)癥下藥。MNN GPU提供統(tǒng)計(jì)OpenCL kernel耗時(shí)方法，使用opencl event進(jìn)行GPU端計(jì)時(shí)統(tǒng)計(jì)，可以精確地進(jìn)行單個(gè)kernel耗時(shí)分析，準(zhǔn)確性能明顯高于CPU端計(jì)時(shí)器。

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

? 適合GPU加速模型設(shè)計(jì)建議

經(jīng)常會(huì)有用戶(hù)反饋，為什么使用GPU加速反而性能不如運(yùn)行CPU上呢？通常在用戶(hù)潛意識(shí)里，GPU總比CPU快。其實(shí)這是個(gè)誤解。GPU的硬件結(jié)構(gòu)設(shè)計(jì)特性，決定了GPU對(duì)具有大量可并行的運(yùn)算才更有優(yōu)勢(shì)。對(duì)于運(yùn)算量過(guò)小或者并行度較低的模型，通常GPU上運(yùn)行效率不如CPU。

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

卷積核不宜太大，常用的1x1和3x3較好。如果模型需要更大卷積核（如5x5）可以考慮使用5x1和1x5來(lái)代替，或者采用兩個(gè)3x3卷積去替代5x5卷積。
通道數(shù)設(shè)計(jì)盡量保持4對(duì)齊
對(duì)于feature map和通道數(shù)都較大的卷積，可以考慮使用depthwise卷積。
加減乘除乘方這類(lèi)binary/unary運(yùn)算量較低的算子，可以有，但是不要過(guò)多。
盡量減少只改變形狀沒(méi)有計(jì)算量的算子，如squeeze、transpose、permute、reshape等。
盡量減少concat/slice這類(lèi)純?cè)L存類(lèi)算子，無(wú)計(jì)算量。
盡量避免使用global pooling。
盡量減少使用除reduce軸之外的維度尺寸較大的reduction操作。

總之，適合GPU計(jì)算的模型，就是模型中的算子，盡可能多的滿(mǎn)足具有大量可并行的特點(diǎn)；減少低計(jì)算量、高訪(fǎng)存算子的使用，避免不好并行運(yùn)算的算子。

? 參考文獻(xiàn)

[1] https://en.wikipedia.org/wiki/OpenCL

[2] "Qualcomm Snapdragon Mobile Platform OpenCL General Programming and Optimization Guide"

[3] "The OpenCL Specification，Version: 1.2，Document Revision: 15"

[4] "ARM Mali GPU OpenCL，Version 3.3 Developer Guide"

[5] "Arm Mali GPU Datasheet 2020"

[6] "Building Heterogeneous Systems with PowerVROpenCL Programmer’s Reference"

[7] "顏深根等. "基于OpenCL 的歸約算法優(yōu)化." 軟件學(xué)報(bào) (2011)."

[8] https://www.cnblogs.com/xudong-bupt/p/3586518.html

[9] https://zhuanlan.zhihu.com/p/273657259

[10] https://developer.arm.com/solutions/graphics-and-gaming/developer-guides/advanced-guides/mali-gpu-best-practices