On-Device GPU Image Processing via OpenCL and the Android NDK

An 8.4x Gaussian blur speedup from JNI-bridged OpenCL kernels on a mobile GPU

01Problem

A Gaussian blur computes a weighted average over each pixel's neighbourhood. On a CPU that is a serial pass over every pixel; on a GPU it is one work-item per pixel and the arithmetic is unchanged. This project tested that difference on an actual Android board, which meant writing the OpenCL host code by hand against the NDK: at the time almost nothing was documented about running OpenCL on Android, so the driver was built up call by call from the Khronos guide.

02What I built

Wrote OpenCLDriver.c from scratch — platform and device enumeration, context and command queue, runtime kernel build, buffer I/O, NDRange dispatch, and explicit release of every OpenCL object.
Exposed four native methods to the Java activity through JNI (GaussianBlurBitmap on the CPU, GaussianBlurGPU, GaussianGreyGPU, GaussianRotateGPU), all sharing one driver path that differs only in the .cl file and the kernel name.
Timed each path from Java with System.nanoTime(): the same blur takes 2059 ms on the CPU and 246 ms on the GPU, an 8.4x speedup with no change to the algorithm.

03How it works

The Java activity holds an ARGB_8888 bitmap and calls a native method. The driver locks the bitmap with AndroidBitmap_lockPixels, copies the pixels into a staging buffer so the original survives, initialises OpenCL, reads and builds the .cl file at run time, writes the pixels into a device buffer, dispatches a one-dimensional NDRange over all pixels with a local work size of 64, reads the result back, releases every OpenCL object, unlocks the bitmap, and hands it back to Java for display.

One call path: Activity (Java) → JNI → OpenCLDriver.c → clBuildProgram and clEnqueueNDRangeKernel → Mali GPU. Pixels cross the same boundary twice, between AndroidBitmap_lockPixels and clEnqueueWriteBuffer on the way in and clEnqueueReadBuffer on the way out.

1BuildCMake 3.22 compiles OpenCLDriver.c into a shared library and links libGLES_mali.so together with jnigraphics and the Android log library.

2Read the bitmapThe driver reads the bitmap's width, height, stride, and format, checks it is RGBA, and copies the pixels into a native staging buffer.

3Initialise OpenCLclGetPlatformIDs, then clGetDeviceIDs with CL_DEVICE_TYPE_GPU, clCreateContext, and clCreateCommandQueue, with every call wrapped in an error-check macro that logs the OpenCL error code.

4Compile and dispatchThe .cl file for the requested operation is read from the device, built with clCreateProgramWithSource and clBuildProgram, given the source buffer, destination buffer, width, and height as arguments, and enqueued as a 1-D NDRange over all pixels with a local size of 64.

5Read backclEnqueueReadBuffer copies the result into the locked bitmap, every OpenCL object is released, the bitmap is unlocked, and the processed image returns to the ImageView.

04Results

8.4x

Gaussian blur, CPU to GPU

246 ms

Gaussian blur on the GPU

81 ms

Grayscale on the GPU

opencl-result.png — output of the three GPU kernels on the test image, photographed from the device for the project slides: original, Gaussian blur, grayscale, and rotation.

Key takeaway Measured on the device with System.nanoTime(): Gaussian blur 2059 ms on the CPU against 246 ms on the GPU, grayscale 81 ms, rotation 72 ms. Blur costs more than the other two on the GPU because each work-item reads a neighbourhood instead of a single pixel; grayscale and rotation touch one pixel each and finish under 100 ms.

05What went wrong

Symptom

On the target board the application crashed on launch whenever the camera preview was part of the build, even though the permission was declared in AndroidManifest.xml.

Hypotheses

The same crash had appeared in an earlier lab when the camera permission was handled incorrectly, so that was the first hypothesis; the second was an incompatibility between the board's Android version and the Camera API it was being called with.

Tests

Corrected and re-checked the permission both in the manifest and at run time, and the crash stayed. Rebuilt the same preview and capture flow in a separate project on the same board, where it ran without a problem, which removed the permission hypothesis without identifying a replacement.

Fix

The root cause was never found. I removed the camera preview, capture, and hold path and loaded a stored image with BitmapFactory.decodeFile instead, which left the OpenCL pipeline — the actual subject of the project — intact and measurable.

Takeaway

Isolating the camera proved where the fault was not, never where it was. Cutting an unexplained subsystem out to protect the deliverable was the right call against the deadline; closing the project without an explanation is the part I would not repeat.

06Engineering decisions

Copy the pixels into a staging buffer instead of processing the bitmap in place. The original image has to survive so the user can return to it and run a different filter on the same input.
Keep one driver path for all three GPU kernels and vary only the .cl file and the kernel name. The OpenCL host code is identical for all of them, so three copies would have meant fixing the same bug three times.
Compile the kernel at run time from a file on the device rather than embedding it in the library. A kernel can then be edited and rerun without rebuilding and reinstalling the APK.