EdgeFirst Hardware Abstraction Layer

The EdgeFirst Hardware Abstraction Layer (HAL) is a Rust workspace providing hardware-accelerated tensor management, image processing, ML model output decoding, and multi-object tracking for edge AI inference pipelines. It ships as a Rust crate, a Python package, and a C library, all built from the same code. A single OpenGL ES engine runs on Linux (native EGL + DMA-BUF), macOS and iOS (ANGLE over Metal + IOSurface), and Android (native EGL + AHardwareBuffer), alongside NXP G2D on i.MX and a portable CPU fallback everywhere else.

• Zero-copy memory management — DMA-BUF, IOSurface, AHardwareBuffer, POSIX shared memory, OpenGL PBO, and heap, with automatic backend selection
• Zero-copy CUDA tensor mapping — convert() PBO output mapped directly to a CUDA device pointer for TensorRT and other CUDA consumers; no host round-trip on Jetson (Orin-series). See Zero-copy CUDA (TensorRT) input.
• Hardware-accelerated image processing — OpenGL → G2D → CPU dispatch with shared cache infrastructure
• Tiled inference (SAHI) — overlapping tile grid rendered in one GPU pass, with IoS-based merge of per-tile detections back to full-frame coordinates. See Tiled inference (SAHI).
• YOLO + ModelPack decoding — YOLOv5 / v8 / v11 / v26 (incl. end-to-end) and ModelPack post-processing
• Multi-object tracking — ByteTrack with Kalman filtering and stable per-track UUIDs
• Cross-platform — Linux (i.MX 8M Plus / i.MX 95 / RPi 5 / Jetson / desktop), macOS, iOS, and Android, over CPU / GPU / zero-copy-buffer tiers

Python:

pip install edgefirst-hal

Rust:

[dependencies]
edgefirst-hal = "0.28"

C: download a release archive from GitHub Releases and link against libedgefirst_hal.so (or .a); see crates/capi/README.md for full instructions.

Python:

import edgefirst_hal as ef

# Decode into a tensor you own. A real pipeline allocates once and reuses
# the tensor every frame; JPEG decodes to its native Nv12, PNG to Rgb/Rgba/Grey.
info = ef.Tensor.peek_image_info_file("image.jpg")
src = ef.Tensor.image(info.width, info.height, info.format)
src.decode_image_file("image.jpg")

processor = ef.ImageProcessor()
model_input = processor.create_image(640, 640, ef.PixelFormat.Rgb)

# convert() handles the colour conversion and the resize in one call.
# Omit letterbox= to stretch to fill instead of preserving aspect ratio.
processor.convert(src, model_input, letterbox=[114, 114, 114, 255])

# outputs is the list of Tensors your inference engine produced from model_input.
decoder = ef.Decoder(model_config, score_threshold=0.5, iou_threshold=0.45)
boxes, scores, classes, masks = decoder.decode(outputs)

# Fused decode + draw: masks never leave Rust.
processor.draw_masks(decoder, outputs, model_input)

Rust:

The umbrella edgefirst-hal crate re-exports its sub-crates as modules, so a single edgefirst-hal = "0.28" dependency is enough. There's no need to list edgefirst-image / edgefirst-tensor separately in Cargo.toml.

use edgefirst_hal::image::{ImageProcessor, ImageProcessorTrait, Rotation, Flip, Crop};
use edgefirst_hal::image::codec::{ImageDecoder, ImageLoad};
use edgefirst_hal::tensor::{PixelFormat, DType, CpuAccess};

let bytes = std::fs::read("image.jpg")?;
let mut processor = ImageProcessor::new()?;
let mut decoder = ImageDecoder::new();

// JPEG decodes to its native NV12 (colour); decode into an NV12 source tensor.
// load_image() reconfigures the tensor's shape and format to the decoded
// content, so allocate at or above the largest frame you expect.
let mut input =
    processor.create_image(1920, 1080, PixelFormat::Nv12, DType::U8, None, CpuAccess::Write)?;
let _info = input.load_image(&mut decoder, &bytes)?;

// convert() handles NV12 -> RGB and the letterbox resize in one call. The
// decode never rotates; pass the EXIF rotation here if you want it applied.
let mut output =
    processor.create_image(640, 640, PixelFormat::Rgb, DType::U8, None, CpuAccess::Read)?;
processor.convert(&input, &mut output, Rotation::None, Flip::None,
    Crop::letterbox([114, 114, 114, 255]))?;

If you prefer to depend on the sub-crates directly (e.g. to opt out of features or to track them at independent versions), add the relevant edgefirst-image, edgefirst-tensor, edgefirst-decoder, and edgefirst-tracker entries to your Cargo.toml and use the unprefixed edgefirst_image::* / edgefirst_tensor::* paths above.

C:

#include <edgefirst/hal.h>

struct hal_image_processor *proc = hal_image_processor_new();
/* `src` is decoded from disk with hal_tensor_decode_image_file(), or
 * imported from a DMA-BUF fd with hal_import_image() / hal_tensor_from_fd().
 * See the C API README for the full allocate-then-decode pattern. */
struct hal_tensor *src = /* ... */;
struct hal_tensor *dst = hal_image_processor_create_image(
    proc, 640, 640, HAL_PIXEL_FORMAT_RGB, HAL_DTYPE_U8, HAL_CPU_ACCESS_READ_WRITE);
hal_image_processor_convert(proc, src, dst, HAL_ROTATION_NONE, HAL_FLIP_NONE, NULL);

On CUDA-capable devices (e.g. Jetson Orin-series) the float PBO produced by convert() can be mapped directly to a CUDA device pointer with no host round-trip. The recommended pattern is to try cuda_map() first and fall back to the host map() when CUDA is unavailable:

Rust:

use edgefirst_hal::tensor::{Tensor, TensorTrait, is_cuda_available};

// At pipeline startup — check once
if is_cuda_available() {
    println!("CUDA present; will use zero-copy PBO→CUDA path");
}

// Per frame — try CUDA, fall back to host
if let Some(cuda) = dst.cuda_map() {
    // cuda.device_ptr() is a raw device pointer valid until `cuda` is dropped.
    // Drop `cuda` before the next convert() so the PBO is free to be reused.
    trt_enqueue(cuda.device_ptr(), cuda.len());
    // `cuda` drops here → PBO released
} else {
    let host = dst.map()?;
    trt_enqueue_host(host.as_slice());
}

Python:

import edgefirst_hal as ef

proc = ef.ImageProcessor()
dst = proc.create_image(640, 640, ef.PixelFormat.PlanarRgb, "float16")

for frame in camera_frames:
    proc.convert(frame, dst)
    cuda = dst.cuda_map()          # CudaMap | None
    if cuda is not None:
        with cuda:
            # cuda.device_ptr is a CUDA device pointer (int)
            trt_context.execute(cuda.device_ptr)
    else:
        host = dst.map()
        trt_context.execute_host(bytes(host))

cuda_map() fast-fails to None when libcudart is not present at runtime — no compile-time feature gate, no link-time dependency. CUDA register/map runs on the GL worker thread; the returned device pointer is usable from any thread. Drop the CudaMap guard before the next convert() call to release the PBO back to the GL pipeline.

For the full mechanism, aliasing rules, DMA-BUF import path, and per-language API reference, see crates/tensor/README.md § CUDA tensor mapping and crates/tensor/ARCHITECTURE.md § Zero-copy CUDA tensor mapping.

Small objects in a high-resolution frame disappear when the whole frame is squeezed down to a 640×640 model input. SAHI (Slicing Aided Hyper Inference) runs the same model at its native resolution over an overlapping grid of native-resolution crops instead, then stitches the per-tile detections back together. HAL covers both halves: edgefirst-image cuts and renders the grid, edgefirst-decoder lifts and merges the results.

The input side renders every tile into one tall packed batch tensor with a single GL import and a single flush, so N tiles cost roughly one GPU sync rather than N. The output side merges with GREEDYNMM using the IoS (intersection-over-smaller) metric, because an object split across a tile seam has low IoU with its own fragments but high IoS. A TilePlacement produced by plan_tiles / tile_into is the shared record of how each tile was cut, and it is what the merge uses to lift boxes back to full-frame coordinates.

Rust:

use edgefirst_hal::image::{ImageProcessor, ImageProcessorTrait, TilingConfig};
use edgefirst_hal::decoder::{DecoderBuilder, DetectBox, Nms, Segmentation};
use edgefirst_hal::decoder::tiling::{MergeConfig, TiledFrameAccumulator};
use edgefirst_hal::tensor::{CpuAccess, DType, PixelFormat};

// 640x640 tiles with at least 20% overlap. The realized overlap is
// redistributed evenly so every tile is full-size and the last one lands flush.
let cfg = TilingConfig::new(640, 640).with_overlap(0.2);

// plan_tiles is pure geometry (no GPU work), so its length sizes the batch.
let placements = processor.plan_tiles(src_w, src_h, &cfg)?;

// One tall [tile_w, N * tile_h] destination. Allocate once, reuse per frame.
let mut batch = processor.alloc_tile_batch(
    placements.len(), &cfg, PixelFormat::Rgb, DType::U8, None, CpuAccess::None)?;

// Render every tile: deferred convert per tile, one flush at the end.
let placements = processor.tile_into(&src, &mut batch, &cfg)?;

// Per-tile decoding is deliberately permissive. A fragment clipped at a seam
// scores low, and a high per-tile threshold discards it before the merge can
// rebuild the object. Gate the final scores in MergeConfig instead.
let decoder = DecoderBuilder::new()
    .with_config_yaml_str(model_config_yaml)
    .with_score_threshold(0.05)
    .with_nms(Some(Nms::ClassAware))
    .build()?;

let mut acc = TiledFrameAccumulator::new(
    (src_w as f32, src_h as f32),
    placements.len(),       // tiles_total — the fan-in fence
    MergeConfig::default(), // Ios metric, 0.5 threshold, max_det 300
    16,                     // estimated detections per tile (capacity hint)
);

// HAL does not run inference. `tile_results` pairs each placement with the
// output tensors your engine produced for that tile; tiles may arrive in any
// order, so pair them explicitly rather than relying on loop position.
for (tile_outputs, placement) in tile_results {
    let mut boxes: Vec<DetectBox> = Vec::new();
    let mut masks: Vec<Segmentation> = Vec::new();
    decoder.decode(&tile_outputs, &mut boxes, &mut masks)?;
    acc.push_tile(boxes, &placement);
}

// Merged, deduplicated, normalized to [0, 1] for the tracker.
let detections = acc.finalize_normalized();

Python:

import edgefirst_hal as ef

cfg = ef.TilingConfig(640, 640, overlap=0.2)
placements = processor.plan_tiles(src.width, src.height, cfg)

batch = processor.alloc_tile_batch(len(placements), cfg, ef.PixelFormat.Rgb)
placements = processor.tile_into(src, batch, cfg)

acc = ef.TiledFrameAccumulator(
    (float(src.width), float(src.height)), len(placements), ef.MergeConfig())

for tile_outputs, placement in tile_results:
    boxes, scores, classes, _masks = decoder.decode(tile_outputs)
    acc.push_tile(boxes, scores, classes, placement)

boxes, scores, classes = acc.finalize_normalized()

push_tile is idempotent per placement.index, so tiles can arrive in any order and an at-least-once delivery retry stays harmless. The merge runs once at finalize, never per push, which is what a pipelined runtime needs: plan_tiles sizes the ring up front, tile_one streams individual tiles through inference into a caller-owned slot, and is_complete() / remaining() fence the frame.

MergeConfig tunes the metric (Ios by default, or Iou), the match threshold (0.5), class_agnostic (false), max_det (300), and a final score_threshold (0.0). That last default is deliberate: per-tile decoding is the real flood control, and the merged score gate belongs after fragments have been joined.

The C API mirrors the same split (hal_image_processor_plan_tiles / _tile_into / _tile_one on the input side, hal_tiled_frame_accumulator_* on the output side). Full per-language detail lives in image/README.md § Tiled Preprocessing and decoder/README.md § Tiled Inference.

Per-language quick-starts and richer examples live in each crate's README: Rust (edgefirst-hal), C API, Python.

graph TB
    subgraph "EdgeFirst HAL Ecosystem"
        Python["Python Bindings (edgefirst-hal)<br/>PyO3"]
        CAPI["C API (edgefirst-hal-capi)<br/>cbindgen"]
        Main["Umbrella crate (edgefirst-hal)<br/>Re-exports"]

        Python --> Main
        CAPI --> Main

        Tensor["edgefirst-tensor<br/>Zero-copy buffers"]
        Codec["edgefirst-codec<br/>Image decode"]
        Image["edgefirst-image<br/>Format conv + draw"]
        Decoder["edgefirst-decoder<br/>Model output decode"]
        Tracker["edgefirst-tracker<br/>ByteTrack"]

        Main --> Tensor
        Main --> Codec
        Main --> Image
        Main --> Decoder
        Main -.->|tracker feature| Tracker
        CAPI --> Tracker

        Codec --> Tensor
        Image --> Tensor
        Image --> Decoder
        Image -.optional.-> G2D["g2d-sys<br/>NXP i.MX"]
    end

    Tensor -.-> DMA["Linux DMA-Heap<br/>Shared Memory"]
    Decoder -.-> PostProc["Model Output<br/>Post-Processing"]

    style Python fill:#e1f5ff
    style CAPI fill:#e1f5ff
    style Main fill:#fff4e1
    style Tensor fill:#e8f5e9
    style Codec fill:#e8f5e9
    style Image fill:#e8f5e9
    style Decoder fill:#e8f5e9
    style Tracker fill:#e8f5e9

The deep dive on each component (class diagrams, supported operations, backend dispatch, performance considerations) lives in the per-crate ARCHITECTURE.md. The cross-cutting story (DMA-BUF identity, performance tracing internals, design patterns) lives in the project ARCHITECTURE.md.

This section is the rules part of the cross-language performance contract. Each rule has a measurable cost when broken; see BENCHMARKS.md for empirical penalties per platform, ARCHITECTURE.md for why the rule exists, and TESTING.md for how to verify your integration follows it.

Allocate input and output tensors once at pipeline startup; reuse the same objects on every frame. The DMA memory backing a tensor is live: when an upstream producer (V4L2 DQBUF, codec output, ISP) writes new pixels into it, the existing tensor and its cached EGLImage remain valid. No re-import, no re-allocation.

let mut proc = ImageProcessor::new()?;
let mut dst = proc.create_image(640, 640, PixelFormat::Rgb, DType::U8, None, CpuAccess::ReadWrite)?;

for frame in camera_frames {
    proc.convert(&frame, &mut dst, Rotation::None, Flip::None, Crop::default())?;
    run_inference(&dst)?;
}

proc = ef.ImageProcessor()
dst = proc.create_image(640, 640, ef.PixelFormat.Rgb)
for frame in camera_frames:
    proc.convert(frame, dst)
    run_inference(dst)

create_image() selects the fastest memory backend for the active GPU at construction time:

The probe runs once at ImageProcessor::new() time. All subsequent create_image() calls reuse the same backend. Use create_image() for every destination passed to convert(); direct Tensor::new(memory=...) bypasses the probe.

Declare CPU access. Every image constructor takes a required CpuAccess parameter: hardware (GPU/NPU/ISP/codec) access is always implied, CPU access is the opt-in. Declare Write for decode targets, Read for buffers you verify/consume on the CPU, ReadWrite when both, and CpuAccess::None for pure hardware pipelines — on Android a hardware-only buffer is eligible for gralloc's vendor tile compression (UBWC/AFBC/PVRIC/DCC), and on every platform the declaration selects the cheapest mapping mode (write-combined for Write, read-only IOSurface locks / dma-buf sync direction for Read). Mapping beyond the declaration still works best-effort but warns once per buffer and counts in unplanned_cpu_access_count(); on Android hardware-only buffers refuse CPU maps deterministically.

For DMA-buf access, the process needs /dev/dma_heap/{linux,cma|system} and a DRM render/card node — the GL backend probes /dev/dri/renderD128, then /dev/dri/card0, then /dev/dri/card1 and uses the first one that opens. On embedded Linux, add the user to video and render groups, or set udev rules. If DMA-buf fails, create_image() transparently falls back to PBO or heap.

V4L2, libcamera, and codec output all surface frames as DMA-BUF file descriptors drawn from a small fixed pool (typically 4–16 buffers). The fd number is recycled: the same fd can refer to a different physical buffer between frames, and the same physical buffer can be exported with a different fd over time. A cache keyed by fd will produce false hits or false misses.

The kernel assigns each dma_buf object a unique inode in the anonymous inode filesystem. The inode is constant for the buffer's lifetime regardless of how many times it is exported. Cache imported HAL tensors by (inode, plane_offset):

#include <sys/stat.h>

typedef struct { ino_t inode; size_t offset; } BufferKey;

struct stat st;
if (fstat(fd, &st) != 0) continue;
BufferKey key = { .inode = st.st_ino, .offset = plane_offset };

struct hal_tensor *tensor = lookup_tensor(cache, &key);
if (!tensor) {
    struct hal_plane_descriptor *pd = hal_plane_descriptor_new(fd);
    if (!pd) { perror("hal_plane_descriptor_new"); continue; }
    tensor = hal_import_image(proc, pd, NULL, w, h,
                              HAL_PIXEL_FORMAT_NV12, HAL_DTYPE_U8,
                              NULL /* colorimetry: NULL = default */);
    // pd is consumed by hal_import_image (success or failure)
    if (!tensor) { perror("hal_import_image"); continue; }
    insert_tensor(cache, &key, tensor);
}
hal_image_processor_convert(proc, tensor, dst, /* ... */);

import os
buffer_cache: dict[tuple[int, int], ef.Tensor] = {}

def get_or_import(proc, fd, offset, width, height, fmt):
    key = (os.fstat(fd).st_ino, offset)
    t = buffer_cache.get(key)
    if t is None:
        t = proc.import_image(fd, width, height, fmt, "uint8", offset=offset)
        buffer_cache[key] = t
    return t

EdgeFirst's GStreamer elements implement this as a reference. For other pipelines (libcamera direct, custom V4L2, RTSP decoder) you are responsible for the equivalent layer above HAL. See ARCHITECTURE.md § Appendix C for the full identity-and-caching story.

Decoder parses the model output schema, resolves quantization, and allocates working buffers at construction time. Build it once outside the loop; the decoder clears its output vectors per call:

let decoder = DecoderBuilder::default()
    .with_config_yaml_str(config_yaml)
    .with_score_threshold(0.5)
    .with_iou_threshold(0.45)
    .build()?;

for frame in frames {
    let outputs = run_inference(frame)?;
    let refs: Vec<&TensorDyn> = outputs.iter().collect();
    decoder.decode(&refs, &mut boxes, &mut masks)?;
}

The same applies to ByteTrack: construct once, call update() per frame.

ImageProcessor owns its OpenGL context, dedicated GL thread, and EGL image cache. The EGL display itself is process-global (a shared SharedEglDisplay initialized once and never terminated), so additional processors don't pay the display-creation cost — but each one still creates a fresh context and per-instance caches. Whether GL operations across processors run in parallel is a per-driver policy: on Vivante (i.MX 8M Plus) and virtualized/paravirtual GPUs every command serializes on a global GL_MUTEX; on Mali, V3D, Tegra, llvmpipe, and real Apple GPUs they execute concurrently (override with EDGEFIRST_GL_SERIALIZE). Construct one per pipeline (or one per worker thread for parallel pipelines) and share it across all convert(), draw_*(), and create_image() calls.

ImageProcessor is Send + Sync, so it can be moved or shared across threads. On serializing drivers, concurrent use of a single shared instance funnels through GL_MUTEX; per-worker ownership runs in parallel wherever the driver allows and gives more predictable cache behaviour everywhere.

The default HAL binary is built to the target triple's guaranteed baseline ISA so a single distributed binary runs on every CPU within that triple. Richer ISAs (ARMv8.2-FP16, x86_64 F16C / FMA / AVX2) are not enabled by default; until HAL gains runtime CPU-feature detection with dynamic dispatch, baking them in would SIGILL on older CPUs.

For local benchmarking on supporting hosts, enable them via RUSTFLAGS:

# Orin Nano (Cortex-A78AE) — exclude the PyO3 binding (cross-Python toolchain not configured)
RUSTFLAGS="-C target-cpu=cortex-a78ae" cargo build --release \
  --target aarch64-unknown-linux-gnu --workspace --exclude edgefirst_hal

# Generic aarch64 with FEAT_FP16 (do NOT use on Cortex-A53 / imx8mp)
RUSTFLAGS="-C target-feature=+fp16" cargo build --release \
  --target aarch64-unknown-linux-gnu -p edgefirst-image

# x86_64 Haswell+ (F16C + FMA + AVX2)
RUSTFLAGS="-C target-feature=+f16c,+fma,+avx2" cargo build --release \
  -p edgefirst-image

When active, the f16 mask kernel at crates/image/src/cpu/masks.rs compiles to native widening (fcvt on aarch64, vcvtph2ps on x86_64), and on x86_64 with +f16c,+fma an explicit 8-lane _mm256_cvtph_ps + _mm256_fmadd_ps intrinsic path is enabled via cfg gate. Verify with scripts/audit_f16_codegen.sh.

Tensor.from_numpy() (and the implicit copy from numpy arrays passed to Decoder.decode_proto()) handles strided / non-contiguous sources internally. Do not maintain a manual np.ascontiguousarray() workaround — it wastes a copy.

The Python binding's copy_numpy_to_tensor_dyn selects one of three paths based on the source array's layout:

The fully-strided case is the one that bites users in practice: HailoRT's natural output is arr.transpose(0, 2, 1) over a (1, anchors, channels) buffer. PR #58 replaced the legacy element-wise loop with internal np.ascontiguousarray materialization (≈ 4× faster than the legacy loop, within ≈ 1.5× of the manual workaround).

# Wrong (post-PR #58): adds an extra copy above HAL.
tensor.from_numpy(np.ascontiguousarray(arr_strided))

# Right: HAL detects the strided layout and materializes internally.
tensor.from_numpy(arr_strided)

The regression tests in tests/test_tensor.py (test_from_numpy_hailort_shape, test_from_numpy_hailort_shape_perf_sanity) pin the behaviour and the ≤ 1.5× perf bound.

ImageProcessor.materialize_masks() accepts a MaskResolution parameter:

import edgefirst_hal as hal

# Wrong: threshold then upsample → blocky edges, mAP regression.
tiles = proc.materialize_masks(boxes, scores, classes, proto_data, letterbox=lb)
for tile, box in zip(tiles, boxes):
    binary = (tile[:, :, 0] > 127).astype(np.uint8)
    canvas[y:y+h, x:x+w] = cv2.resize(binary, (W, H), cv2.INTER_NEAREST)

# Right: HAL upsamples-then-thresholds inside its batched-GEMM kernel.
tiles = proc.materialize_masks(boxes, scores, classes, proto_data,
                               letterbox=lb,
                               resolution=hal.MaskResolution.Scaled(W, H))
for tile, box in zip(tiles, boxes):
    canvas[y:y+h, x:x+w] = (tile[:, :, 0] > 127).astype(np.uint8)

The Scaled path uses the batched-GEMM materializer (PR #54). At N ≥ 16 detections it amortizes a single GEMM at proto resolution and upsamples per-detection in rayon-parallel — both more accurate than threshold-then-resize and faster than per-detection scalar work in caller code.

¹ Android's bionic libc has no POSIX shm_open, so shared-memory tensor allocation reports NotImplemented; importing an existing segment received as a file descriptor (from_fd) works.

On macOS the OpenGL backend is enabled when ANGLE is installed — see macOS GPU Acceleration below for setup. If ANGLE is not present the HAL falls back to the CPU backend. On iOS the OpenGL backend uses the same ANGLE-over-Metal path — see iOS below. On Android the OpenGL backend uses the platform's native GLES driver directly (no translation layer) — see Android below.

The HAL uses Google's ANGLE to translate the same OpenGL ES 3.0 calls used on Linux to Metal, and Apple's IOSurface for zero-copy buffer interchange (the role DMA-BUF plays on Linux). ANGLE is not part of macOS and must be installed separately. If it is not present at runtime the HAL logs a warning and falls back to the CPU backend.

ANGLE access: ANGLE itself is an open-source Google project, and our pre-built, signed + notarized xcframework integration is published from the public repository (EdgeFirstAI/angle-package). Anyone can fetch it — no credentials or organization membership required. Two ways to get ANGLE:

• Recommended (macOS + iOS) — fetch the pre-built release with scripts/fetch-angle.sh (see Option A below). This is exactly what CI uses.
• macOS alternative — install ANGLE via the public Homebrew tap: brew install startergo/angle/angle (then re-sign the dylibs — see Option B — Homebrew tap below). The HAL finds it automatically.
• Build without macOS/iOS GL — the HAL's default features include opengl, but you can disable it (--no-default-features --features ndarray,tracing) to build the CPU-only path, which needs no ANGLE at all.

The HAL looks for libEGL.dylib / libGLESv2.dylib via the EDGEFIRST_ANGLE_PATH env var, then standard search paths (Homebrew, @loader_path, @executable_path). There are two ways to satisfy this:

Our pre-built, signed + notarized xcframeworks (built from a pinned ANGLE revision) are published in the EdgeFirstAI/angle-package releases. This repo is public — anyone can fetch the release with no credentials. A single helper downloads, sha256-verifies, and extracts them (into both the xcframework layout for iOS app embedding and a flat-lib layout for the macOS runtime dlopen path):

scripts/fetch-angle.sh                       # → target/angle/ (default tag v2.1.28252)
EDGEFIRST_ANGLE_PATH=target/angle/macos-flat-lib \
  cargo run --release --example pipeline_demo

Because the release is public, scripts/fetch-angle.sh needs no authentication — it works out of the box both locally and in CI. (It still honors gh auth login / GH_TOKEN / GITHUB_TOKEN if present, which raises GitHub's API rate limit, but none are required.)

Why a flat-lib dir for macOS? ANGLE's libEGL internally dlopens libGLESv2.dylib from its own directory (located via dladdr) to resolve GL entry points, so the two must be flat siblings. The signed framework bundles do not satisfy this, so the helper stages libEGL.dylib + libGLESv2.dylib siblings copied out of the framework binaries. Pulling a binary out of its framework invalidates the Developer-ID signature (it is scoped to the bundle's Info.plist, so dlopen then fails with "code signature invalid"), so scripts/fetch-angle.sh ad-hoc re-signs the two flat dylibs for you — you never re-sign manually (unlike the Homebrew path below).

ANGLE is also available via a public third-party Homebrew tap — an alternative to Option A if you prefer a package manager on macOS. Homebrew's install_name_tool step invalidates the bundled code signatures and macOS 26 (Tahoe) refuses to load dylibs with broken signatures at dlopen time (immediate SIGKILL (Code Signature Invalid) with no stdout), so an ad-hoc re-sign is mandatory after each install/upgrade:

brew install startergo/angle/angle
codesign --force --sign - $(brew --prefix)/opt/angle/lib/libEGL.dylib
codesign --force --sign - $(brew --prefix)/opt/angle/lib/libGLESv2.dylib

See Homebrew/brew#19144 for the upstream tracking issue. The release path above avoids this problem entirely.

RUST_LOG=edgefirst_image=debug cargo run --release --example pipeline_demo

Look for ANGLE (Apple, ANGLE Metal Renderer: ...) in the bring-up log. If ANGLE is missing or signatures are still broken you will see a warning and the CPU backend is selected.

If your ANGLE install is not on the default search path, set EDGEFIRST_ANGLE_PATH to the directory containing libEGL.dylib and libGLESv2.dylib (flat siblings — see the note above):

EDGEFIRST_ANGLE_PATH=/path/to/angle/lib cargo run --release ...

The lookup order is: EDGEFIRST_ANGLE_PATH → Homebrew → @loader_path (alongside the binary) → @executable_path → unqualified libEGL.dylib on the dyld search path. For bundled distributions, drop the re-signed ANGLE dylibs next to the executable (or into <App>.app/Contents/Frameworks/) and no env var is needed.

• pip install edgefirst-hal — the macOS wheel ships ANGLE bundled alongside the Python extension; no separate install required.
• EdgeFirst-signed binary distribution — official binary releases bundle ANGLE re-signed under the EdgeFirst Apple Developer ID. Install and run with no additional setup.

These channels exist precisely so end users do not need to deal with the Homebrew install or re-signing step.

The HAL Rust library closure builds for iOS (arm64 device + arm64 simulator) with the default features (including opengl), reusing the same ANGLE-over-Metal GL backend as macOS. The supported targets are:

• aarch64-apple-ios — iOS devices (arm64)
• aarch64-apple-ios-sim — iOS Simulator on Apple-Silicon Macs (arm64)

ANGLE note: iOS GL requires ANGLE xcframeworks. There is no public Homebrew equivalent for iOS (unlike macOS), so fetch them from the public angle-package release with scripts/fetch-angle.sh (no credentials needed). If you would rather not fetch ANGLE at all, you can still build the Rust library for iOS with the opengl feature disabled: cargo build --target aarch64-apple-ios --no-default-features --features ndarray,tracing. The Rust cargo build itself (with opengl) succeeds without ANGLE present — see How the GL backend resolves ANGLE on iOS.

Intel-simulator (x86_64-apple-ios) is not supported — the angle-package distribution ships arm64-only slices (see below).

Xcode + the iOS SDKs (xcode-select --install or a full Xcode), plus the Rust iOS targets:

rustup target add aarch64-apple-ios aarch64-apple-ios-sim

The one-command entry point builds for both targets and validates the link closure against the ANGLE xcframeworks:

scripts/build-ios.sh                 # device + sim, build + link-validate
scripts/build-ios.sh device          # device only
scripts/build-ios.sh --no-validate   # build only, skip link validation

Or build the library closure directly:

cargo build --target aarch64-apple-ios     --release -p edgefirst-hal
cargo build --target aarch64-apple-ios-sim --release -p edgefirst-hal

ANGLE's EGL/GLES symbols are resolved at runtime via libloading, not at link time. On macOS the HAL dlopens libEGL.dylib from the release flat-lib (or Homebrew); on iOS the symbols are already in the process image (the ANGLE xcframeworks are embedded in the app bundle), so the loader resolves them via Library::this() (equivalent to dlopen(NULL)).

Consequence: a standalone cargo build for an iOS target succeeds without the ANGLE frameworks present — the Rust staticlib has no link-time references to eglInitialize etc. The frameworks are only needed at app-link/runtime. The .cargo/config.toml iOS entries therefore carry no rustflags or linker overrides.

iOS GL requires shipping ANGLE as embedded dynamic frameworks in the app bundle. Our integration uses the signed + notarized xcframeworks from the public EdgeFirstAI/angle-package release (EGL.xcframework + GLESv2.xcframework, each with ios-arm64, ios-arm64-simulator, macos-arm64). scripts/fetch-angle.sh downloads and verifies them (default tag v2.1.28252, matching the ANGLE GL_VERSION string):

scripts/fetch-angle.sh       # → target/angle/{EGL,GLESv2}.xcframework

A consuming iOS app target embeds them (Xcode "Embed & Sign", or XcodeGen embed: true):

dependencies:
  - { framework: ../hal/target/angle/EGL.xcframework,   embed: true }
  - { framework: ../hal/target/angle/GLESv2.xcframework, embed: true }

scripts/validate-ios-link.sh builds the production C API staticlib (libedgefirst_hal.a — a Rust staticlib archives the full HAL closure) and links a test executable referencing real C API entry points against the ANGLE xcframeworks + the Apple system frameworks the HAL references via #[link(kind = "framework")] (IOSurface, CoreFoundation, Metal). It also verifies with nm that the archive carries the IOSurface references and that the ANGLE binaries export the EGL entry-point names the runtime loader will look up. This proves the native symbol closure of the shipped artifact is complete.

What is not covered by this effort (future work):

• Swift bindings — a C/Swift API surface and a ship-able HAL .xcframework. The Rust staticlib is the deliverable here.
• Runtime validation — actual EGL initialization on a device or simulator requires the app shell (a future effort). The internal hal-mobile assessment already proved the ANGLE-over-Metal + IOSurface path works on iPhone 17 Pro (GL_EXT_color_buffer_half_float present).

Unlike aarch64-apple-darwin (where +fp16,+dotprod,+i8mm are baked in — every M-series chip is ARMv8.6-A+), the iOS targets carry no target-feature rustflags. The iOS 16 deployment floor still includes A11 (iPhone 8, ARMv8.1-A, no fp16/dotprod/i8mm), so enabling them would SIGILL on older devices. The deployment target matches the angle-package build (IPHONEOS_DEPLOYMENT_TARGET = 16.0).

The HAL builds for Android with the default features (including opengl), using the platform's native OpenGL ES driver directly — unlike macOS/iOS there is no ANGLE translation layer to install, because Android ships a first-class GLES implementation (Adreno, Mali, etc.). Zero-copy buffer interchange uses AHardwareBuffer (the role DMA-BUF plays on Linux and IOSurface on Apple platforms), imported into GL via EGL_ANDROID_image_native_buffer. The supported targets are:

• aarch64-linux-android — Android devices (arm64-v8a)
• x86_64-linux-android — the Android emulator on x86_64 hosts

The minimum supported API level is 26 (Android 8.0) — the floor of the stable AHardwareBuffer NDK ABI.

rustup target add aarch64-linux-android x86_64-linux-android
cargo install cargo-ndk
# Android NDK r26+ (r27c LTS recommended); set ANDROID_NDK_HOME or let
# cargo-ndk auto-detect it under your Android SDK.

# HAL + C API (the future JNI library) for both ABIs at API 26:
scripts/build-android.sh
# or directly:
cargo ndk -t arm64-v8a -t x86_64 -P 26 build --release -p edgefirst-hal

scripts/validate-android-link.sh [arm64|x86_64] builds the production C API staticlib (libedgefirst_hal.a — a Rust staticlib archives the full HAL closure), verifies with llvm-nm that the archive carries the AHardwareBuffer references and that the NDK's API-26 stubs export every EGL/GLES entry point the runtime resolves dynamically, then links a test executable referencing real C API entry points against the NDK system libraries. CI runs this for both ABIs on every PR (build-android lane).

What is not covered by this effort (future work):

• Kotlin bindings — a JNI/Kotlin API surface, like the Swift API on iOS. The edgefirst-hal-capi cdylib (libedgefirst_hal.so) is the deliverable here.
• Runtime validation — on-device GL correctness and performance run via the internal hal-mobile AWS Device Farm harness, which drives the real ImageProcessor through JNI (see TESTING.md § Android On-Device Validation); the Phase-1 assessment already proved the native-GLES + AHardwareBuffer path on a Galaxy S26 Ultra (GL_EXT_color_buffer_half_float present, letterbox 720p→640×640 F16 in 741 µs).
• Deferred zero-copy paths — YUV camera buffers (external-OES sampling) and single-channel Grey/NV imports (R8_UNORM needs API 29); these fall back to CPU conversion today.

The convert destination is a real AHardwareBuffer, so an NPU runtime can consume it directly — no map(), no CPU readback:

// Allocate once, reuse every frame (Rule 1). F16 NCHW model input;
// auto-select yields an AHardwareBuffer when the GL backend is active —
// assert hal_tensor_memory_type(dst) == HAL_TENSOR_MEMORY_DMA at startup.
HalTensor* dst = hal_image_processor_create_image(
    proc, 640, 640, HAL_PIXEL_FORMAT_PLANAR_RGB, HAL_DTYPE_F16, HAL_CPU_ACCESS_NONE);

// One-time: hand the SAME buffer to the NPU runtime.
AHardwareBuffer* ahb = hal_tensor_hardware_buffer_ptr(dst);
ANeuralNetworksMemory* mem;
ANeuralNetworksMemory_createFromAHardwareBuffer(ahb, &mem);   // NNAPI
// (LiteRT: wrap `ahb` via TfLiteAHardwareBufferAttachment instead.)

// Per frame: when convert() returns, the GPU has finished writing and
// the handle contents are safe to execute against.
hal_image_processor_convert(proc, src, dst, HAL_ROTATION_NONE,
                            HAL_FLIP_NONE, &letterbox);
// ... ANeuralNetworksExecution_setInputFromMemory(exec, 0, NULL, mem, 0, bytes);

For a pipelined handoff that skips the blocking GPU sync entirely, use hal_image_processor_convert_fence(): it returns a sync-fence fd (EGL_ANDROID_native_fence_sync) the NPU runtime waits on instead (ANeuralNetworksExecution_startComputeWithDependencies), or -1 with the work already synced on drivers without fence support.

Flatness: gralloc chooses the row pitch and may pad it (observed on the S26 Ultra: 640-px planar F16 → 1536-byte rows, natural 1280). Check hal_tensor_recorded_row_stride(dst):

• 0 — the buffer IS the flat [1, C, H, W] stream; hand it off as-is.
• nonzero — describe the pitch to the runtime, pick a width whose pitch the device does not pad, or fall back to hal_tensor_copy_to_flat(dst, buf, len) (~0.3 ms at 2.4 MB — still cheaper than a full CPU convert, but no longer zero-copy; profile).

INT8 NPUs: allocate the destination as HAL_PIXEL_FORMAT_RGB / HAL_PIXEL_FORMAT_RGBA with HAL_DTYPE_U8 or HAL_DTYPE_I8 (NHWC, zero-copy on Android via the RGBA8888 texel packing) and attach the model's quantization so consumers agree on the scale:

hal_tensor_set_quantization(dst, /*scale=*/1.0f / 255.0f, /*zero_point=*/0);

The I8 path applies the ^0x80 bias in-shader during the convert — the buffer bytes are already signed model input.

Tile compression (bandwidth): a hardware-only destination can additionally request the device's vendor tile layout through the image-descriptor path — the GPU renders into it and Qualcomm's QNN can consume it natively (UBWC data formats declared at context-binary preparation); other NPU stacks take the linear default:

HalImageDesc* desc = hal_image_desc_new(640, 640, HAL_PIXEL_FORMAT_RGBA, HAL_DTYPE_U8);
hal_image_desc_set_compression(desc, HAL_COMPRESSION_ANY);   // linear fallback is counted
HalTensor* dst = hal_image_processor_create_image_desc(proc, desc);
hal_image_desc_free(desc);
// hal_tensor_compression(dst) records the scheme actually allocated
// (HAL_COMPRESSION_UBWC on Adreno, ..._NONE = linear fallback — see
// hal_compression_fallback_count()).

Compression requires HAL_CPU_ACCESS_NONE (CPU mapping pins the layout linear) and a compressed tensor has no meaningful linear row stride — it is a hardware-to-hardware handle only.

macOS parity: the same pattern works with hal_tensor_iosurface_ref() — wrap the IOSurface in a CVPixelBuffer (CVPixelBufferCreateWithIOSurface) for CoreML/ANE input; convert() returning likewise guarantees GPU completion.

Like iOS, the Android targets carry no target-feature rustflags: the API-26 device floor spans ARMv8.0-A cores (Cortex-A53 class, no fp16/dotprod/i8mm), so baking those features in would SIGILL on real older hardware (see .cargo/config.toml).

The workspace builds with standard cargo. The Makefile wraps the common workflows (make test, make bench, make build, make format lint check) with the right flags and gates.

For Python wheels, see crates/python/README.md and crates/python/TESTING.md. For the C library and consumer linking, see crates/capi/README.md.

Per-crate variables and additional detail live in each crate's README.

See TESTING.md for the cross-cutting testing guide (single-threaded rule, on-target gating, cross-compilation, CI matrix, optimization validation). Per-crate testing detail lives in each crate's TESTING.md — links in the Core Components table.

Run on host:

cargo bench -p edgefirst-image --bench pipeline_benchmark -- --bench

# Force a backend
EDGEFIRST_FORCE_BACKEND=cpu cargo bench -p edgefirst-image --bench pipeline_benchmark -- --bench

Cross-compile + deploy to a target (SSH hostnames in ~/.ssh/config: imx8mp-frdm, imx95-frdm, rpi5-hailo, jetson-orin-nano, maivin):

cargo-zigbuild zigbuild --target aarch64-unknown-linux-gnu --release \
  -p edgefirst-image --features opengl --bench pipeline_benchmark

scp target/aarch64-unknown-linux-gnu/release/deps/pipeline_benchmark-* imx8mp-frdm:/tmp/
ssh imx8mp-frdm '/tmp/pipeline_benchmark-* --bench --json /tmp/pipeline.json'

All benchmarks accept --bench --json <path> for structured output. Store results under benchmarks/<platform>/<name>.json. Update BENCHMARKS.md via:

python3 .github/scripts/generate_benchmark_tables.py --data-dir benchmarks/

The HAL captures performance traces across every processing stage. Traces are written in the Chrome JSON format and open directly in the Perfetto UI.

Every HAL library crate emits tracing spans on hot paths. Those spans cost close to nothing when no subscriber is active: each site compiles to a single relaxed atomic load, with no heap allocations, no string formatting, and no function calls on the hot path.

When a session is started via the API, a Chrome JSON subscriber records all span enter/exit events with high-resolution timestamps and structured metadata (detection counts, proto dimensions, format conversions, memory types, etc.) to a file.

The tracing surface covers decode, image conversion, GL multi-pass, mask materialization, tensor lifecycle, tracker association, and the Python entry points. Each span carries structured fields — see the per-crate ARCHITECTURE.md files for the authoritative list of spans and fields per component.

Python:

import edgefirst_hal as hal
with hal.Tracing("/tmp/trace.json"):
    # ... run inference pipeline ...
    pass

Rust:

use edgefirst_hal::trace::{start_tracing, stop_tracing};

start_tracing("/tmp/trace.json").expect("start tracing");
// ... inference pipeline ...
stop_tracing(); // flushes and closes the trace file

C:

#include <edgefirst/hal.h>
hal_start_tracing("/tmp/trace.json");
/* ... inference pipeline ... */
hal_stop_tracing();

1. Open https://ui.perfetto.dev/
2. Drag the generated .json file onto the page
3. Click slices to see structured fields in the Current Selection panel

The tracing infrastructure complements the rules in the Optimization Guide and the data in BENCHMARKS.md:

1. Identify bottlenecks — common findings:
   • decoder.decode_proto.extract_proto_data > 3 ms → model emits NCHW protos but HAL is transposing (check the layout field)
   • image.convert.cpu.format_convert appearing twice → intermediate format conversion (consider matching src/dst formats)
   • tensor.alloc per-frame → tensors not being reused (Rule 1)
   • image.convert.gl.egl_import on every frame → camera tensors re-imported instead of cached (Rule 3)
2. Validate rules — re-run with tracing after applying a rule to confirm the expected spans disappear or shrink.
3. Cross-reference with perf — for CPU-bound spans, combine trace data with perf record for instruction-level hotspots.

• Only one trace session per process lifetime (Rust global subscriber model).
• Rayon worker spans are not automatically parented to the calling span.
• The log::* output (via env_logger / C callback logger) operates independently from trace capture; both can be active simultaneously.

• PyO3 — Python bindings
• ndarray — N-dimensional arrays
• rayon — Data parallelism
• fast_image_resize — CPU image operations
• zune-png — PNG image decoding (JPEG uses custom decoder)
• dma-heap — Linux DMA allocation
• nix — Unix system calls

graph TD
    EF[edgefirst-hal<br/>umbrella]
    Tensor[edgefirst-tensor]
    Image[edgefirst-image]
    Decoder[edgefirst-decoder]
    Tracker[edgefirst-tracker<br/>optional]
    G2D[g2d-sys<br/>optional]

    EF --> Tensor
    EF --> Image
    EF --> Decoder
    Image --> Tensor
    Image --> Decoder
    Image -.optional.-> G2D
    Image -.->|tracker feature| Tracker
    Decoder -.->|tracker feature| Tracker

    Python[edgefirst_hal<br/>PyO3]
    CAPI[edgefirst-hal-capi]

    Python --> EF
    CAPI --> EF
    CAPI --> Tensor
    CAPI --> Image
    CAPI --> Decoder
    CAPI --> Tracker

    style EF fill:#fff4e1
    style Python fill:#e1f5ff
    style CAPI fill:#e1f5ff
    style Tracker fill:#e8f5e9

1. Model HAL — planned abstraction for inference engines (ONNX, TFLite, Kinara)
2. VPI integration — support for NVIDIA Vision Programming Interface
3. Additional trackers — SORT, Deep SORT
4. Async I/O — non-blocking image loading and processing

• GitHub Discussions — questions and ideas
• Issue Tracker — bug reports and feature requests

This project is part of the EdgeFirst Perception stack:

• EdgeFirst Studio — complete MLOps platform
• EdgeFirst Hardware Platforms — NPU/GPU acceleration on NXP i.MX

Au-Zone Technologies supports production deployments with training and workshops, custom development, integration services, enterprise SLAs, and hardware reference designs.

Contact: [email protected] · au-zone.com

Contributions are welcome. See CONTRIBUTING.md for development setup and guidelines. This project follows our Code of Conduct.

For security vulnerabilities, see SECURITY.md or email [email protected] with subject "Security Vulnerability".

• ARCHITECTURE.md — cross-crate architecture story
• TESTING.md — workspace testing rules and CI matrix
• BENCHMARKS.md — empirical performance reference
• CHANGELOG.md — release history
• Per-crate docs (README + ARCHITECTURE + TESTING) — see Core Components table

Apache License 2.0 — see LICENSE for details.

Copyright 2025-2026 Au-Zone Technologies