Tiny-OpenCL Framework for TinyAI Edge Devices

• Tiny-OpenCL is a minimalist, header-only OpenCL-compliant runtime that integrates with RISC-V e-GPUs to support TinyAI workloads on ultra-constrained systems.
• It employs a streamlined architecture with a static host driver, inlined device-side scheduler, and unified memory model to ensure efficient, synchronous compute operations.
• Co-designed with e-GPU hardware, the framework achieves measurable speed-ups and energy savings within strict area and power budgets, proving its practical impact for embedded computing.

The Tiny-OpenCL framework is a minimalist, header-only OpenCL-compliant runtime tightly coupled with the e-GPU platform—a configurable RISC-V-based GPU for TinyAI workloads. It addresses the challenge of enabling OpenCL-like accelerator programmability on ultra-constrained systems by trading full API generality for compactness, low runtime overhead, and seamless integration with bare-metal RISC-V microcontroller environments. Designed as an embedded, statically linked library for systems with stringent area and power budgets, Tiny-OpenCL provides a subset of the OpenCL API through a minimal host driver, an inlined device-side scheduler, and a unified memory model mapped directly onto host DRAM. The framework’s co-design with e-GPU hardware enables end-to-end speed-ups and energy reductions within sub-0.4 mm² and <30 mW constraints, while maintaining a negligible scheduling overhead for practical problem sizes (Machetti et al., 13 May 2025).

1. Architectural Components and Programming Model

Tiny-OpenCL’s architecture is organized around three primary entities: a host-side driver, a device-side scheduler, and a memory manager, each streamlined for resource-limited deployment [(Machetti et al., 13 May 2025), §5, §6].

• Host Driver and Command Queue: Implemented as a static library linked into Newlib-based RISC-V firmware. The host exports only a subset of the OpenCL API (clCreateBuffer, clSetKernelArg, clEnqueueNDRangeKernel, clFinish), each mapped to straightforward memory-mapped register accesses on the e-GPU controller.
• Device Scheduler: Embedded as “startup” and “scheduler” routines, precompiled with kernel images into the e-GPU instruction cache. At kernel launch, compute units (CUs) execute scheduler code in single-threaded mode until the main SIMT workload commences.
• Memory Manager: Directly exposes the host DRAM as the e-GPU’s unified address space. Buffers are statically allocated at link time by the host and passed to the device without explicit clEnqueueRead/WriteBuffer operations. Cache parameters are hardware-configurable, with minimum alignment at 4 bytes and buffer sizes enforced to cache-line multiples.

The programming model diverges sharply from full OpenCL: no dynamic object management, no explicit event or out-of-order execution, and no host-side callbacks. Command queues are in-order and strictly synchronous, blocking on hardware interrupts. Work-items and work-groups are linearly mapped to hardware threads, with all out-of-bounds checks hoisted into the scheduler, allowing kernels to be written without explicit bounds guards [(Machetti et al., 13 May 2025), §5.2].

2. Scheduling Mechanism and Runtime Overhead

Scheduling in Tiny-OpenCL is managed centrally by the device-side scheduler, which is responsible for decomposing the global workload and assigning it to available hardware threads [(Machetti et al., 13 May 2025), Fig. 10]. The scheduler reads global and local sizes from a reserved DRAM region, computes the work-group grid, and iteratively assigns each work-item to a hardware thread. The mapping process involves computing the number of work-groups, assigning thread IDs, and issuing SLEEP_REQ instructions at completion to support fine-grained CU power gating.

Formal runtime decomposition of a kernel launch for matrix dimensions $M \times N$ is:

$$T_{\text{total}}(M,N) = T_{\text{transfer}}(M,N) + T_{\text{sched}}(M,N) + T_{\text{compute}}(M,N)$$

• $$T_{\text{transfer}}(M,N) = \alpha + \beta \cdot 4 M N$$, where $\alpha$ models setup stalls and $\beta$ reflects per-byte transfer on a 32b bus at 300 MHz.
• $$T_{\text{sched}}(M,N) \approx T_0 \approx 25\ \mu\text{s}$$, observed experimentally to be constant across problem sizes and thread configurations.

For $N \geq 256$, $T_{\text{sched}}/T_{\text{total}} < 1\%$. Compute time dominates for realistic workloads, rendering scheduling costs negligible at practical scales [(Machetti et al., 13 May 2025), Fig. 10].

Pseudocode for the scheduler:

void tinyocl_schedule() { Gx, Gy = arg->global_size[0,1]; Lx, Ly = arg->local_size[0,1]; numWG = ceil(Gx/Lx)*ceil(Gy/Ly); thread_id = 0; for (wg=0; wg<numWG; wg++) { for (wi=0; wi<Lx*Ly; wi++) { if (thread_id >= HW_THREADS) break; launch_workitem(thread_id++, wg, wi); } } // issue SLEEP_REQ to signal end of scheduling }

Editor's term: "Flat thread mapping" refers to this work-item flattening.

3. Memory Hierarchy and Data Transfer Model

The unified memory model eliminates explicit data movement primitives from the API, instead exposing a single global address space mapped to host DRAM [(Machetti et al., 13 May 2025), §4.2]. Buffers managed by the host runtime are simply pointers in this space. The e-GPU’s instruction and data caches mediate all coherency requirements, parameterized by line size, associativity, and bank count.

Transfer time modeling leverages the bus characteristics:

$$T_{\text{transfer}} = \alpha + \beta \times \text{DataSize},\quad \text{with}\ \text{DataSize} = 4 \cdot M \cdot N\ \text{bytes}$$

Alignment requirements: buffers must be 4-byte aligned and sized to at least one cache line to prevent bank conflicts.

No clEnqueueRead/WriteBuffer or host-side memory commands are present in the runtime; all buffer residency and validity is enforced statically.

4. Host Integration and ISA Extensions

Tiny-OpenCL achieves system-level integration by extending the host’s standard C library (Newlib) with a minimal driver layer [(Machetti et al., 13 May 2025), §6]. The host sets up kernel launches by writing base addresses and argument pointers to e-GPU memory-mapped registers, specifying grid/work-group sizes, issuing kernel start commands, and synchronizing via interrupts. The completion interrupt signals to the host when execution has completed, minimizing CPU polling.

On the device side, the RISC-V CUs implement an ISA extension: the SLEEP_REQ instruction. This low-level primitive allows CUs to signal “done” events after all instructions retire, supporting dynamic power gating at the compute unit granularity.

Integration with X-HEEP APU uses the eXtendible Heterogeneous Energy-Efficient Platform interconnect, with the e-GPU attached via separate OBI (One-Bus Interface) slave and master ports for control and data respectively.

5. End-to-End Performance and Efficiency

Tiny-OpenCL’s practical efficiency has been demonstrated through two benchmarks:

• Scheduling Overhead (GeMM): For matrix sizes $32 \times 32$ up to $256 \times 256$ and e-GPU configurations with 4, 8, and 16 threads, scheduling costs plateau at approximately $25\ \mu\text{s}$ and become insignificant compared to total execution time for large problem sizes [(Machetti et al., 13 May 2025), Fig. 10].
• Application Benchmarks (TinyBio): For FIR, peak delineation, Stockham FFT, and SVM algorithms, a 16-thread e-GPU achieves speed-ups between $3.6\times$ and $15.1\times$ over the host, with energy consumption reduced by up to $3.1\times$. The overall bio-signal pipeline attains up to $14.3\times$ speed-up and $2.9\times$ energy savings [(Machetti et al., 13 May 2025), Fig. 11].

Area and power analysis (TSMC 16 nm SVT CMOS, 300 MHz, 0.8 V):

Configuration	Area (mm²)	Leakage (μW)
Host only	0.15	29.5
e-GPU 4T	0.24	130.1
e-GPU 8T	0.31	210.7
e-GPU 16T	0.38	305.3

Area overhead is 1.6×–2.5× vs. the host alone, but total system power remains within the 28 mW ceiling required for TinyAI platforms.

6. Toolchain, Synthesis, and Comparative Perspective

Tiny-OpenCL leverages the standard RISC-V GNU toolchain (gcc 12.2, binutils, Newlib) plus a Python script to transpile OpenCL C kernels into plain C augmented with SIMT intrinsics for the e-GPU. No custom LLVM backend modifications are necessary; the device datapath and memory hierarchy are synthesized (SystemVerilog, SRAM macros per Table 2) for TSMC 16 nm SVT CMOS at 300 MHz [(Machetti et al., 13 May 2025), §6].

In comparison to more generic OpenCL development frameworks such as cf4ocl—which emphasize host-side productivity, profiling, and extensibility for heterogeneous high-end compute devices (Fachada et al., 2016)—Tiny-OpenCL’s distinguishing features are its minimal footprint, direct hardware mapping, and static resource model. Where cf4ocl abstracts OpenCL resource management and profiling while introducing small overhead that vanishes for large kernels, Tiny-OpenCL dispenses entirely with dynamic resource objects and event bookkeeping, making it more suitable for ultra-low-end, bare-metal deployment.

7. Significance for TinyAI and Embedded Heterogeneous Computing

Tiny-OpenCL exemplifies a hardware–software co-design approach tailored for the idiosyncratic constraints of TinyAI-class edge devices. Its design demonstrates that a minimal subset of the OpenCL execution model, when mapped to a unified memory and interrupt-driven schedule, suffices to unlock substantial speed-ups and energy efficiency when paired with a dedicated RISC-V GPU coprocessor. The framework establishes a foundation for further research in embedded accelerator runtime systems operating at the intersection of compute efficiency, memory hierarchy simplification, and predictable scheduling costs under severe resource constraints (Machetti et al., 13 May 2025).