Mastering Mobile OpenCL on Qualcomm Adreno: Architecture & Performance Tips

Introduction

This article extends a previous overview of mobile GPU computing and OpenCL basics, focusing on practical optimization techniques for Qualcomm's Adreno GPU, a representative mobile processor.

OpenCL Basics

OpenCL, maintained by the Khronos Group, is an open, royalty‑free standard that enables cross‑platform parallel programming on heterogeneous systems (CPU, GPU, DSP, FPGA, etc.).

OpenCL Architecture Overview

An abstract OpenCL application consists of three components:

CPU host – manages the OpenCL context, command queues, and kernel compilation.

OpenCL devices – hardware such as GPUs, DSPs, or FPGAs that execute kernels.

OpenCL kernels – device‑side code compiled by the host and run on the device.

Mobile OpenCL Usage

On Android devices OpenCL typically runs on the GPU. Qualcomm recommends explicitly selecting the GPU as the OpenCL accelerator; most devices expose a single GPU device.

Compatibility

Program Portability

OpenCL source code is largely portable across vendors, with differences limited to optional extensions.

Performance Portability

Performance varies between vendors (e.g., Adreno vs. Mali) and even between driver versions. Device‑specific tuning is required for optimal speed.

Backward Compatibility

The core specification strives for backward compatibility, but vendor extensions may not be supported across hardware generations.

Qualcomm Adreno OpenCL Architecture

Shader (Streaming Processor, SP)

Core modules include ALU, load/store units, control‑flow units, and register files.

Executes graphics shaders and OpenCL kernels; each SP maps to one or more OpenCL compute units.

SPs can load Buffer objects or __read_write images via the unified L2 cache.

Read‑only images may use the L1 cache or the texture processor.

Texture Processor (TP)

Handles texture reads, filtering, and other image operations scheduled by the kernel.

Works with L1 cache to reduce L2 cache miss latency.

Unified L2 Cache (UCHE)

Serves SP buffer reads/writes and TP image loads, providing a shared high‑bandwidth memory layer.

Writing High‑Performance OpenCL Code

Optimization Categories

Program/algorithm‑level optimizations – portable across devices.

API‑level optimizations – e.g., command‑queue properties.

Kernel‑level optimizations – device‑specific tuning.

When to Use OpenCL

Large input data sets.

Compute‑intensive workloads.

High degree of data parallelism.

Minimal control‑flow complexity.

CPU‑to‑GPU Conversion Tips

Combine multiple CPU operations into a single kernel to reduce host‑to‑device transfers.

Split complex CPU code into several simple kernels to increase parallelism.

Redesign data structures to minimize transfer overhead.

Parallel CPU‑GPU Workflow

Assign branch‑heavy or serial tasks to the CPU.

Avoid GPU idle time by overlapping CPU work.

Move suitable CPU tasks to the GPU to lower data‑transfer cost.

Performance Analysis

CPU Profiling

Standard C/C++ timing APIs such as gettimeofday can measure host‑side execution time. Example:

#include <time.h>
#include <sys/time.h>
int main() {
    struct timeval start, end;
    gettimeofday(&start, NULL);
    // enqueue and execute OpenCL kernel
    clFinish(command_queue);
    gettimeofday(&end, NULL);
    double elapsed = (end.tv_sec - start.tv_sec) * 1000.0 +
                     (end.tv_usec - start.tv_usec) / 1000.0;
    printf("CPU total time: %f ms
", elapsed);
    return 0;
}

GPU Profiling

OpenCL provides event‑based profiling to capture timestamps for each kernel phase.

// Example of kernel profiling
cl_event prof_event;
cl_ulong t_queued, t_submit, t_start, t_end;
int status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
                                    &global_size, NULL,
                                    0, NULL, &prof_event);
clWaitForEvents(1, &prof_event);
clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_QUEUED,
                        sizeof(cl_ulong), &t_queued, NULL);
clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_SUBMIT,
                        sizeof(cl_ulong), &t_submit, NULL);
clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START,
                        sizeof(cl_ulong), &t_start, NULL);
clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END,
                        sizeof(cl_ulong), &t_end, NULL);
printf("Kernel time: %f ns
", (t_end - t_start) * 1e-0);

Identifying Bottlenecks

OpenCL kernels are typically limited by either memory bandwidth or compute capacity. Simple tests can reveal the dominant factor:

Add extra arithmetic; if runtime does not change, the kernel is not compute‑bound.

Add extra memory loads; if runtime does not change, the kernel is not memory‑bound.

Resolving Bottlenecks

Compute‑bound: Reduce algorithmic complexity, enable fast‑math flags (e.g., -cl-fast-relaxed-math or native_math), or switch to lower‑precision types such as half (fp16) when acceptable.

Memory‑bound: Use vectorized loads/stores, exploit local or texture memory, and employ narrower data types to lower bandwidth demand.

References

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

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

Mali GPU Wikipedia: https://zh.wikipedia.org/wiki/Mali_(GPU)

Adreno GPU Wikipedia: https://zh.wikipedia.org/wiki/Adreno