Optimizing MNN Mobile Neural Network Inference on GPU with OpenCL: Memory Objects, Work‑Group Tuning, and Auto‑Tuning

MNN (Mobile Neural Network) is a high‑performance, general‑purpose deep‑learning framework that runs efficiently on mobile, PC, server and embedded devices. It fully exploits GPU resources to accelerate model inference.

Memory Object Diversification : OpenCL provides Buffer‑object and Image‑object. Qualcomm Adreno GPUs favor Image‑object for faster texture access, while ARM Mali GPUs perform better with Buffer‑object. MNN supports both and selects the appropriate type per device, allowing users to manually override via MNN::ScheduleConfig config; config.mode = MNN_GPU_TUNING_NORMAL | MNN_GPU_MEMORY_IMAGE; .

Memory Alignment Optimization : Aligning dimensions to multiples of 4 reduces boundary checks and branch instructions in kernels. Image‑object size limits and cache‑miss behavior are discussed, with Buffer‑object preferred when dimensions exceed hardware limits.

Local Memory Parallel Reduction : The article presents a parallel reduction algorithm using local memory to compute maximum values, showing both the conceptual description and the OpenCL kernel code:

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]);
}

GPU Compute Chunking : Work‑group size directly impacts hardware utilization. MNN auto‑tunes multiple candidate sizes during a pre‑inference stage and selects the best configuration for the target device.

Data Chunk Reuse : For 2‑D convolutions, increasing per‑thread compute granularity reduces total memory traffic while keeping computational complexity constant. The article provides tables and diagrams illustrating the trade‑off.

Heterogeneous Scheduling : Describes the three‑part OpenCL execution pipeline (host CPU, heterogeneous devices, kernel code) and the state transitions (Queued → Submitted → Ready → Running). It notes that different GPUs (Qualcomm vs. Mali) require different command‑queue flushing strategies, which MNN abstracts via dynamic tuning.

Pre‑Inference Auto‑Tuning : MNN performs memory allocation, task preparation, and command‑buffer generation before actual inference. Auto‑tuning identifies optimal work‑group sizes and data‑chunk parameters, at the cost of extra one‑time overhead.

Cache Mechanism : To avoid costly runtime compilation, MNN stores compiled OpenCL binaries and tuned configurations as cache files. Loading the cache eliminates both compilation and auto‑tuning delays, dramatically reducing startup time on devices such as Xiaomi 6 (Adreno‑540).

Performance Analysis Tools : By enabling the MNN_OPENCL_PROFILE macro, developers can obtain per‑kernel timing via OpenCL events, allowing hotspot identification and model‑level optimizations (e.g., replacing large kernels with stacked smaller ones).

GPU‑Friendly Model Design Recommendations :

Prefer 1×1 or 3×3 convolutions; replace 5×5 with 5×1 + 1×5 or two 3×3 layers.

Keep channel counts 4‑aligned.

Use depthwise convolutions for large feature maps.

Minimize low‑compute, high‑memory‑traffic ops (squeeze, transpose, reshape, concat, slice, global pooling).

Avoid excessive branching; align data to reduce cache misses.

The article concludes with a reference list covering OpenCL specifications, Qualcomm and ARM GPU guides, and related research papers.