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, D3D11 texture, POSIX shared memory, OpenGL PBO, and heap, with automatic backend selection
• Zero-copy CUDA tensor mapping — convert() output mapped directly to a CUDA device pointer for TensorRT and other CUDA consumers: a PBO with no host round-trip on Jetson (Orin-series), a D3D11 texture imported as CUDA external memory on Windows. 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, Android, and Windows, over CPU / GPU / zero-copy-buffer tiers

Python — install only the pieces you need:

pip install edgefirst-tensor    # the zero-copy tensor core
pip install edgefirst-codec     # JPEG/PNG decode (needs tensor)
pip install edgefirst-image     # convert/resize/draw (needs tensor, not decoder)
pip install edgefirst-decoder   # YOLO/ModelPack post-processing
pip install edgefirst-tracker   # ByteTrack (standalone)

Each is a self-contained wheel under the edgefirst. namespace; they compose without any one pulling in the others.

Rust:

cargo add edgefirst-codec
cargo add edgefirst-image   # add only what you use

C: download edgefirst-hal-<version>-<target>.tar.gz (Linux) or .zip (Windows, macOS) from GitHub Releases. The archive is relocatable — see packaging/c/INSTALL.txt for pkg-config, runtime search path, and a JPEG→tensor example. Link libedgefirst_tensor plus any of codec / image / decoder / tracker. Headers live under include/edgefirst/ (tensor.h, codec.h, image.h, decoder.h, tracker.h, detect.h). Windows ships bin/*.dll and lib/*.lib import libraries. Before 1.0 the C ABI is stable across patch releases but may break across minors — any 0.N.z is drop-in for any other 0.N.z, and 0.N to 0.N+1 may not be. While pre-1.0, mix the five libraries only within one minor; from 1.0 the boundary becomes the major. See § C ABI Stability and Versioning.

Python:

from edgefirst.codec import Tensor, decode_file_into
from edgefirst.decoder import Decoder
from edgefirst.image import ImageProcessor, PixelFormat

processor = ImageProcessor()

# Decode straight into a DMA/PBO-backed tensor for zero-copy convert(). A
# real pipeline allocates once and reuses the tensor every frame; JPEG
# decodes to its native Nv12, PNG to Rgb/Rgba/Grey. decode_file_into() is
# edgefirst.codec's free function for decoding into a tensor from *another*
# edgefirst.* package -- create_image() tensors can't be a decode_image()
# method's `self`, so this is the entry point that makes them decodable.
info = Tensor.peek_image_info_file("image.jpg")
# info.format is edgefirst.codec's own PixelFormat, and create_image() below
# is edgefirst.image's -- each package has its own type objects, but
# PixelFormat is a value type: it compares/hashes by value and is accepted
# in argument position across packages, so info.format works directly here.
src = processor.create_image(info.width, info.height, info.format, "uint8", "readwrite")
decode_file_into(src, "image.jpg")

model_input = processor.create_image(640, 640, PixelFormat.Rgb, "uint8", "readwrite")

# 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 = Decoder(model_config, score_threshold=0.5, iou_threshold=0.45)
boxes, scores, classes, masks = decoder.decode(outputs)

# Fused decode + draw lives on the decoder so image-only installs stay decoder-free.
decoder.draw_onto(processor, outputs, model_input)

edgefirst.codec.set_output_format(PixelFormat.Rgb) (4:4:4 sources) or set_output_format(PixelFormat.Nv12) (any colour source) opts a JPEG decode into fusing that colour conversion into the decode itself, in place of decode-then-convert(); set_dct_method(DctMethod.Fast) trades a small, bounded accuracy loss for a faster IDCT. See crates/python-codec/README.md.

Rust:

Depend on the crates you use. There is no umbrella edgefirst-hal crate.

use edgefirst_codec::{ImageDecoder, ImageLoad};
use edgefirst_image::{Crop, Flip, ImageProcessor, ImageProcessorTrait, Rotation};
use edgefirst_tensor::{CpuAccess, DType, PixelFormat};

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

decoder.set_output_format(Some(PixelFormat::Rgb)) (4:4:4 sources) or Some(PixelFormat::Nv12) (any colour source) fuses that colour conversion into the decode itself, in place of decode-then-convert(); decoder.set_dct_method(DctMethod::Fast) trades a small, bounded accuracy loss for a faster IDCT. See crates/codec/ARCHITECTURE.md#fused-decode-output-opt-in-pure-cpu.

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/image.h>
#include <edgefirst/tensor.h>

ef_image_processor *proc = ef_image_processor_new();
/* `src` is decoded with libedgefirst_codec (`ef_image_decoder_decode_file_into`)
 * or wrapped from a dma-buf with `ef_tensor_builder_wrap`. */
ef_tensor *src = /* ... */;
ef_tensor *dst = ef_image_processor_create_image(
    proc, 640, 640, /* format/dtype/access: see image.h */);
ef_image_processor_convert(proc, src, dst, /* rotation, flip, crop */);

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_tensor::{is_cuda_available, TensorTrait};

// 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:

from edgefirst.image import ImageProcessor, PixelFormat

proc = ImageProcessor()
dst = proc.create_image(640, 640, 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 a greedy, class-aware pass under the IoS (intersection-over-smaller) metric, because an object split across a tile seam has low IoU with its own fragments but high IoS; by default it keeps the highest-scoring box of each matched group unchanged (MergeMode::KeepBest). 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_image::{ImageProcessor, ImageProcessorTrait, TilingConfig};
use edgefirst_decoder::{DecoderBuilder, DetectBox, Nms, Segmentation};
use edgefirst_decoder::tiling::{MergeConfig, TiledFrameAccumulator};
use edgefirst_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, keep-best, 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:

from edgefirst.image import TilingConfig, PixelFormat
from edgefirst.decoder import TiledFrameAccumulator, MergeConfig

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

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

acc = TiledFrameAccumulator(
    (float(src.width), float(src.height)), len(placements), 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), a final score_threshold (0.0), and the mode (KeepBest by default, or Union). The score_threshold 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 (ef_image_processor_plan_tiles / ef_image_processor_tile_into / ef_image_processor_tile_one on the input side, ef_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, C API, Python.

graph TB
    subgraph "EdgeFirst HAL Ecosystem"
        PyT["edgefirst-tensor<br/>wheel"]
        PyC["edgefirst-codec<br/>wheel"]
        PyI["edgefirst-image<br/>wheel"]
        PyD["edgefirst-decoder<br/>wheel"]
        PyK["edgefirst-tracker<br/>wheel"]

        TensorSo["libedgefirst_tensor.so"]
        CodecSo["libedgefirst_codec.so"]
        ImageSo["libedgefirst_image.so"]
        DecoderSo["libedgefirst_decoder.so"]
        TrackerSo["libedgefirst_tracker.so"]

        PyT --> TensorSo
        PyC --> CodecSo
        PyC --> TensorSo
        PyI --> ImageSo
        PyI --> TensorSo
        PyD --> DecoderSo
        PyD --> TensorSo
        PyK --> TrackerSo

        CodecSo --> TensorSo
        ImageSo --> TensorSo
        DecoderSo --> TensorSo
    end

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

from edgefirst.image import ImageProcessor, PixelFormat

proc = ImageProcessor()
dst = proc.create_image(640, 640, 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 ef_tensor *tensor = lookup_tensor(cache, &key);
if (!tensor) {
    ef_tensor_builder *b = ef_tensor_builder_new();
    if (!b) { perror("ef_tensor_builder_new"); continue; }
    ef_tensor_builder_dtype(b, EF_DTYPE_U8);
    ef_tensor_builder_format(b, "NV12");
    ef_tensor_builder_add_plane(b, fd, plane_offset, 0, 0, 0, 0);
    tensor = ef_tensor_builder_wrap(b);
    ef_tensor_builder_free(b);
    if (!tensor) { perror("ef_tensor_builder_wrap"); continue; }
    insert_tensor(cache, &key, tensor);
}
ef_image_processor_convert(proc, tensor, dst, /* rotation, flip, crop */);

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-python-tensor \
  --exclude edgefirst-python-codec --exclude edgefirst-python-image \
  --exclude edgefirst-python-decoder --exclude edgefirst-python-tracker

# 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:

from edgefirst.image import MaskResolution

# 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=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.

Two version floors are project policy. Both are chosen, not derived from whatever hardware happens to be on hand, and both are deliberately conservative — raising either one drops users.

The MSRV is set once in [workspace.package] and inherited by every crate. The workspace uses resolver = "3", the MSRV-aware resolver, so rust-version also constrains dependency selection: cargo update will not pull a dependency release that requires a newer compiler. Raising the MSRV is therefore a user-visible change, not just a build detail — bump it here, in rust-toolchain.toml, and in the CHANGELOG together.

No source change was needed to support 1.94.0; it records the floor the code already met. Verified against a real 1.94.0 toolchain across the combinations that compile genuinely different code:

--no-default-features always needs static (or dynamic) named alongside it: edgefirst-tensor's two backend features are mutually exclusive and one is required, so --no-default-features with neither is a compile error, not a fallback to a default.

The python-* crates cannot be cross-compiled without PYO3_CROSS_PYTHON_VERSION or an abi3-py3* feature, so they are covered by the host build rather than the Linux one.

The glibc floor is what scripts/on-target-test.sh builds against (cargo-zigbuild --target <arch>-unknown-linux-gnu.2.35), so one binary set runs on every supported target. Building against a newer glibc links symbols an older loader cannot resolve, and the failure appears only at run time on the oldest device you own. Override for a one-off with GLIBC=<version>; change the default in the script and here together.

Published Python wheels use a separate, lower floor — manylinux2014 (glibc 2.17), via maturin --zig --compatibility manylinux2014 — because they are built for the wider PyPI audience rather than for our targets. make wheel applies it.

¹ 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. On Windows the OpenGL backend runs the same engine on ANGLE over Direct3D 11, rendering zero-copy into D3D11 texture tensors — see Windows GPU Acceleration 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 static,ndarray,tracing) to build the CPU-only path, which needs no ANGLE at all. static is required alongside --no-default-features: edgefirst-tensor's two backend features are mutually exclusive and one is required, so naming neither is a compile error, not a fallback.

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.

• EdgeFirst-signed binary distribution — official binary releases bundle ANGLE re-signed under the EdgeFirst Apple Developer ID. Install and run with no additional setup.
• Windows wheels and C archives bundle ANGLE (see Windows GPU Acceleration); the macOS wheels do not yet — set EDGEFIRST_ANGLE_PATH there as above.

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

On Windows the HAL runs the same OpenGL ES engine on Google's ANGLE translating to Direct3D 11. The zero-copy buffer kind there is a D3D11 texture: TensorMemory::DmaBuf is an ID3D11Texture2D on a HAL-owned D3D11 device that ANGLE renders straight into. If ANGLE cannot be loaded the HAL logs a warning and falls back to the CPU backend.

ImageProcessor::create_image() allocates a texture first, then a PBO, then heap memory, so a GPU destination reports TensorMemory.DMABUF on Windows the way it does on Linux and macOS. Sources are uploaded by GL as before. map() on a texture goes through a staging copy whose pitch becomes the tensor's row_stride(), so a padded destination must be read row by row (or through the buffer protocol) rather than as one flat run. try_map() returns WouldBlock while that copy is still in flight instead of stalling; a retry loop must yield or sleep between attempts, because on the WARP software adapter the CPU threads are the GPU and a tight loop starves the copy it is waiting for.

Semi-planar textures (NV12, NV16, NV24) are allocated as wide as the driver's staging row pitch, always an even width, and that width is the tensor's row_stride(). An external semi-planar texture handed to from_d3d11_texture / from_d3d11_shared_handle (or the C and Python constructors) is accepted only when its width is even and equal to its own staging pitch; anything else is refused with an error naming both numbers. A D3D12 or CUDA producer that needs one should allocate through the HAL, or match that pitch.

Pre-built libEGL.dll + libGLESv2.dll (Direct3D 11 backend, static CRT, no VC++ redistributable needed) are published from the same public EdgeFirstAI/angle-package release tag as the Apple xcframeworks, as angle-windows-x64-<tag>.zip. Fetch, sha256-verify and extract them with the same helper, from Git Bash (the --windows flag is implied on a Windows host):

bash scripts/fetch-angle.sh                # → target/angle/windows-x64/{bin,lib,include}

The HAL looks for libEGL.dll in this order, loading it by absolute path with LOAD_LIBRARY_SEARCH_DLL_LOAD_DIR so ANGLE's libEGL.dll can find its sibling libGLESv2.dll (the two must stay in the same directory):

1. %EDGEFIRST_ANGLE_PATH%\libEGL.dll
2. next to the module that contains the HAL (edgefirst_image.dll from the C archive, or _image.pyd from the wheel) — this is how the bundled distributions work with no configuration;
3. next to the executable;
4. libEGL.dll on the default DLL search path.

Everything else (d3d11.dll, dxgi.dll, d3dcompiler_47.dll) comes from System32 on Windows 10/11.

$env:EDGEFIRST_ANGLE_PATH = "$PWD\target\angle\windows-x64\bin"
$env:RUST_LOG = 'edgefirst_image=debug'
cargo run --release -p edgefirst-image --example pipeline_demo

Look for ANGLE D3D11 adapter: NVIDIA GeForce RTX 3070 ... and ANGLE (NVIDIA, NVIDIA GeForce RTX 3070 ... Direct3D11 ...) in the bring-up log, and GLConverter created (transfer=D3d11Texture). create_image() then returns TensorMemory::DmaBuf destinations.

EDGEFIRST_D3D11_ADAPTER selects the Direct3D 11 adapter the device is created on. EDGEFIRST_ANGLE_ADAPTER is read as an alias and both names stay valid; one variable drives both the device and ANGLE's display, because ANGLE builds its display on the HAL's device instead of creating a second one. When both are set and they differ, EDGEFIRST_D3D11_ADAPTER wins and the HAL logs the disagreement once, so a stale deployment variable cannot quietly pick the adapter. The device and the display are each created once per process, so set the variable before the first allocation or the first ImageProcessor:

RUST_LOG=edgefirst_image=debug lists every DXGI adapter with its LUID. Under a Remote Desktop session the hardware adapter is normally still enumerated; if only the Basic Render Driver is, the HAL warns up front and falls back to CPU unless EDGEFIRST_ALLOW_SOFTWARE_GL=1.

Tensor::cuda_map() / cuda_map_mut() (ef_tensor_cuda_map, ef_tensor_cuda_map_mut, Tensor.cuda_map() in Python) import the texture's shared handle as CUDA external memory and hand back a linear device pointer. A texture is not linear, so each map costs one device copy, ordered behind the fence value the last GPU write recorded. The mapping is tight rows of width * bytes_per_texel and its len() is their sum — not row_stride(), which is the D3D11 staging pitch a CPU map sees and is the larger number on a padded texture. Striding the device pointer by row_stride() reads past the end of the mapping and, on a write, scrambles the image rather than failing. Releasing a writable map copies the buffer back into the texture and synchronizes before it returns.

DirectML runs on D3D12, so a DirectML consumer is a D3D12 consumer: Tensor::d3d11_shared_handle() (ef_tensor_d3d11_shared_handle, Tensor.d3d11_shared_handle()) duplicates the texture's shared NT handle for ID3D12Device::OpenSharedHandle or any other device that opens D3D11 shared resources. The caller owns the handle and closes it with CloseHandle. Tensor::gpu_completion() (ef_tensor_gpu_completion, Tensor.gpu_completion()) reports the fence handle and the value to wait for before reading, which is cheaper than blocking the GL thread. Tensor::gpu_write_value() (ef_tensor_gpu_write_value, Tensor.gpu_write_value) is the value alone, with no handle to close and no DuplicateHandle behind it, for a consumer that already holds the process fence and only needs the number to wait for. A consumer that writes the texture itself records its own completion with Tensor::set_gpu_write() (ef_tensor_set_gpu_write, Tensor.set_gpu_write()); the recorded value is a maximum, so recording one a consumer already holds is safe. In Python a destination's gpu_completion() reflects the convert that just wrote it — the binding writes the recorded value back onto the object the caller passed. The C API needs no such step.

ef_tensor_export / ef_tensor_import (edgefirst_tensor::blob) carry a texture across a process boundary: the blob names the producer's pid and the importer duplicates the handles out of that process, so the producer must still be alive and the importer must be able to open it with PROCESS_DUP_HANDLE. No file descriptors travel on Windows, so the fds arrays stay empty. Independently linked packages inside one process exchange textures through the D3D11_TEXTURE kind of the tensor descriptor protocol (crates/python-common/INTEROP.md).

A host that already has an ID3D11Device — a Media Foundation pipeline, a game engine, a D3D12 application with an interop device — installs it with edgefirst_tensor::d3d11::use_external_device (ef_d3d11_use_external_device, edgefirst.tensor.d3d11_use_external_device) so HAL tensors are allocated on it and ANGLE renders on it. The reference stays the host's. This must run before anything that creates the device as a side effect: the first allocation, d3d11_device(), or is_gpu_buffer_available().

• pip install edgefirst-image — the Windows wheel bundles libEGL.dll
  • libGLESv2.dll next to the extension module.
• The Windows C archive (edgefirst-hal-<version>-x86_64-windows.zip) ships them in bin/ next to edgefirst_image.dll, with ANGLE's BSD licence under share/licenses/angle/.

Both are built by CI with EDGEFIRST_ANGLE_PATH set at build time; a local maturin build / scripts/package-capi.sh bundles ANGLE the same way when the variable is set, and produces a CPU-only artifact when it is not.

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 static,ndarray,tracing (static is required alongside --no-default-features — edgefirst-tensor's two backend features are mutually exclusive and one is required). 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

This repo's mobile responsibility is the native Rust API compiling (and linting clean) on the iOS targets — not a C artifact, not Swift bindings, not app packaging. mobile-sdk binds to these crates directly via boltffi and owns everything above that line. Build the sibling crates directly:

cargo build --target aarch64-apple-ios     --release \
  -p edgefirst-tensor -p edgefirst-image -p edgefirst-codec -p edgefirst-decoder -p edgefirst-tracker
cargo build --target aarch64-apple-ios-sim --release \
  -p edgefirst-tensor -p edgefirst-image -p edgefirst-codec -p edgefirst-decoder -p edgefirst-tracker

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 }

CI's build-ios job proves the Rust API — including the iOS-only IOSurface backend and the GL platform seam — compiles and lints clean for aarch64-apple-ios and aarch64-apple-ios-sim. It links nothing and produces no artifact for shipping; that this repo's mobile responsibility stops there is deliberate, not an oversight.

What is not covered here, because it belongs to mobile-sdk:

• boltffi bindings and Swift/Kotlin packaging — whether the result is a monolith .xcframework or a modular one is mobile-sdk's call, not this repo's.
• Link-closure validation against the ANGLE xcframeworks + Apple system frameworks (IOSurface, CoreFoundation, Metal) — that is a property of whatever mobile-sdk ships, not of this crate's .rlib.
• Runtime validation — actual EGL initialization on a device or simulator requires the app shell. 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.

This repo's mobile responsibility is the native Rust API compiling (and linting clean) on the Android targets — not a C artifact, not Kotlin bindings, not app packaging. mobile-sdk binds to these crates directly via boltffi and owns everything above that line. Build the sibling crates directly, both ABIs at API 26:

cargo ndk -t arm64-v8a -t x86_64 -P 26 build --release \
  -p edgefirst-tensor -p edgefirst-image -p edgefirst-codec -p edgefirst-decoder -p edgefirst-tracker

CI's build-android job proves the Rust API — including the Android-only AHardwareBuffer backend and the GL platform seam — compiles and lints clean for aarch64-linux-android and x86_64-linux-android. It links nothing and produces no artifact for shipping.

What is not covered here, because it belongs elsewhere:

• boltffi bindings and Kotlin packaging — mobile-sdk's responsibility, not this repo's.
• Link-closure validation against the NDK system libraries — a property of whatever mobile-sdk ships, not of this crate's .rlib.
• 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 ef_tensor_storage_kind(dst) == EF_STORAGE_KIND_DMA_BUF at startup.
ef_tensor *dst = ef_image_processor_create_image(
    proc, 640, 640, "rgb8_planar", EF_DTYPE_F16, EF_STORAGE_KIND_DMA_BUF, EF_CPU_ACCESS_NONE);

// One-time: hand the SAME buffer to the NPU runtime.
AHardwareBuffer* ahb = ef_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.
ef_image_processor_convert(proc, src, dst, 0 /* rotation none */,
                            0 /* flip none */, &letterbox);
// ... ANeuralNetworksExecution_setInputFromMemory(exec, 0, NULL, mem, 0, bytes);

For a pipelined handoff that skips the blocking GPU sync entirely, use ef_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 ef_tensor_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 ef_tensor_copy_to(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 "rgb8" / "rgba8" with EF_DTYPE_U8 or EF_DTYPE_I8 (NHWC, zero-copy on Android via the RGBA8888 texel packing) and attach the model's quantization so consumers agree on the scale:

float scale = 1.0f / 255.0f;
int32_t zp = 0;
ef_tensor_quantization_set(dst, /*axis=*/-1, &scale, &zp, 1);

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:

ef_tensor_image_desc *desc = ef_tensor_image_desc_new(640, 640, "rgba8", EF_DTYPE_U8);
ef_tensor_image_desc_set_compression(desc, 1);   // 1 = any scheme; linear fallback is counted
ef_tensor *dst = ef_image_processor_create_image_desc(proc, desc);
ef_tensor_image_desc_free(desc);
// ef_tensor_compression(dst) records the scheme actually allocated
// (EF_COMPRESSION_UBWC on Adreno, EF_COMPRESSION_NONE = linear fallback — see
// ef_compression_fallback_count()).

Compression requires EF_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 ef_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-tensor/README.md and crates/python-common/TESTING.md. For the C libraries and consumer linking, see crates/tensor-capi/README.md and the sibling *-capi READMEs.

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:

from edgefirst.tensor import Tracing
with Tracing("/tmp/trace.json"):
    # ... run inference pipeline ...
    pass

Rust:

use edgefirst_tensor::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/tensor.h>
ef_start_tracing("/tmp/trace.json");
/* ... inference pipeline ... */
ef_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
    Tensor[edgefirst-tensor]
    DetectAbi[edgefirst-decoder-abi<br/>header-only values]
    Codec[edgefirst-codec]
    Image[edgefirst-image]
    Decoder[edgefirst-decoder]
    Tracker[edgefirst-tracker]
    G2D[g2d-sys<br/>Linux only]

    Image --> Tensor
    Image --> DetectAbi
    Image -.->|codec feature, on by default| Codec
    Image -.->|decode feature, off by default| Decoder
    Image -.->|Linux target only| G2D
    Image -.->|tracker feature| Tracker
    Decoder --> Tensor
    Decoder --> DetectAbi
    Decoder -.->|tracker feature| Tracker
    Codec --> Tensor
    Tracker -.->|static or tensor-boxes feature| Tensor

    Python[edgefirst.{tensor,codec,image,decoder,tracker}<br/>PyO3]
    TensorC[libedgefirst_tensor]
    ImageC[libedgefirst_image]
    CodecC[libedgefirst_codec]
    DecoderC[libedgefirst_decoder]
    TrackerC[libedgefirst_tracker]

    Python --> Tensor
    Python --> Codec
    Python --> Image
    Python --> Decoder
    Python --> Tracker
    TensorC --> Tensor
    ImageC --> Image
    CodecC --> Codec
    DecoderC --> Decoder
    TrackerC --> Tracker

    style Python fill:#e1f5ff
    style TensorC fill:#e1f5ff
    style ImageC fill:#e1f5ff
    style CodecC fill:#e1f5ff
    style DecoderC fill:#e1f5ff
    style TrackerC 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