NPU Stability Issues on RadxaOS (Kernel 6.1.84-vendor-rk35xx)

[KHAEntertainment]

Environment

• Board: Rock 5C (RK3588)
• OS: RadxaOS (Kernel 6.1.84-8-rk2410)
• RKNPU2 Package: 2.3.0-1
• RKNN Runtime: 2.3.2
• RKNPU Driver: 0.9.8

Problem Description

Experiencing severe NPU instability when running llama.cpp with RKNPU2 backend on RadxaOS. Issues manifest as:

1. EFAULT errors (failed to allocate handle, errno: 14) during inference after ~60s of model loading
2. Board crashes/reboots frequently during NPU operations
3. CPU-only builds also crash the board, indicating the issue is not purely RKNPU-related

Error Messages

E RKNN: failed to allocate handle, errno: 14 (EFAULT)
E RKNN: failed to submit matmul task

Known Driver Bugs

1. Spinlock recursion bug in RKNPU 0.9.x with 4+ contexts in auto core mode

   • Reference: RKNPU 0.9.2 Driver has Spinlock Recursion Bug for RK3588  rockchip-linux/kernel#329
   • Workaround: Force single core (RKNN_NPU_CORE_0)

2. Kernel panic issues with NPU operations on 6.1.84-vendor-rk35xx

   • Reference: https://forum.armbian.com/topic/48399-6184-vendor-rk35xx-kernel-panic/

Investigation Summary

Test	Result
Standalone SDK matmul test	PASS
SDK stress test (50 contexts)	PASS
CPU-only llama-cli build	FAIL (board crashes)
NPU-accelerated llama-cli	FAIL (EFAULT errors)

Conclusion: Hardware/system instability that predates our fork's changes.

Workarounds Applied

1. Force single NPU core (RKNN_NPU_CORE_0)
2. Use static memory context for B-matrix allocations
3. Map quantization to power of 2 for context caching

Recommendation

Switch to Armbian for Rock 5C, which provides more stable NPU driver support.

Expected Behavior

NPU should operate stably for sustained inference sessions without board crashes or EFAULT errors.

Related

• Original driver issue: RKNPU 0.9.2 Driver has Spinlock Recursion Bug for RK3588  rockchip-linux/kernel#329
• Forum discussion: https://forum.armbian.com/topic/48399-6184-vendor-rk35xx-kernel-panic/

🤖 Generated with Claude Code

[traycerai]

Plan

Observations

The three workarounds described in the issue are already partially implemented in file:ggml/src/ggml-rknpu2/ggml-rknpu2.cpp:

1. Single-core forced — RKNN_NPU_CORE_0 is hardcoded in get_matmul_ctx() (lines 183–195), with the multi-core path commented out.
2. Static memory context for B-matrix — get_rknpu_memory_context() singleton is used for B-matrix allocations (lines 374–392).
3. Power-of-2 M quantization — rknpu2_calibration::next_power_of_two(M) is applied to M_ctx (line 306).

However, the issue also reports CPU-only builds crashing, and the error handling in ggml_backend_rknpu_graph_compute is incomplete (e.g., rknn_matmul_run failure is silently swallowed at line 519). The device selection is also hardcoded to "RK3588" with a TODO comment.

Approach

The plan focuses on three concrete improvements: (1) surface rknn_matmul_run errors properly so failures are visible rather than silent, (2) make the NPU core mask and device selection configurable via environment variables to allow runtime tuning without recompilation, and (3) add a graceful fallback path when the memory context fails to initialize, so the backend degrades cleanly instead of crashing.

---

Implementation Steps

1. Propagate rknn_matmul_run errors correctly

In ggml_backend_rknpu_graph_compute (lines 514–522 of file:ggml/src/ggml-rknpu2/ggml-rknpu2.cpp), the rknn_matmul_run return value is checked but the error branch is empty. This means EFAULT errors are silently ignored and the function returns GGML_STATUS_SUCCESS even after a failed NPU submission.

• Replace the empty // Handle error comment with a fprintf(stderr, ...) that prints the segment index and the RKNN error code, then set a shared atomic/flag so the outer loop can return GGML_STATUS_FAILED after the parallel section completes.
• Since the loop uses #pragma omp parallel for, use a local std::atomic<bool> had_error{false} declared before the loop, set it inside the error branch, and check it after the loop exits before proceeding to result collection.

2. Make NPU core selection configurable via environment variable

Currently get_matmul_ctx() unconditionally sets core_mask = RKNN_NPU_CORE_0 (line 183). The multi-core path is commented out.

• Read an environment variable (e.g., RKNN_CORE_MASK) at backend initialization time inside ggml_backend_rknpu_device_init_backend.
• Store the resolved rknn_core_mask value in ggml_backend_rknpu_context as a new field (e.g., rknn_core_mask core_mask).
• In get_matmul_ctx(), replace the hardcoded RKNN_NPU_CORE_0 with this->core_mask.
• Valid string values to parse: "0" → RKNN_NPU_CORE_0, "1" → RKNN_NPU_CORE_1, "2" → RKNN_NPU_CORE_2, "auto" → RKNN_NPU_CORE_AUTO. Default to RKNN_NPU_CORE_0 when the variable is absent or unrecognized (preserving the current safe default for driver 0.9.x).
• Log the selected core mask to stderr during initialization so users can confirm the setting.

3. Make device selection configurable via environment variable

ggml_backend_rknpu_device_init_backend (line 1020) hardcodes "RK3588" with a TODO comment.

• Read an environment variable (e.g., RKNN_DEVICE) at the start of ggml_backend_rknpu_device_init_backend.
• Pass the value (defaulting to "RK3588" when absent) to Rknpu2ConfigManager::select_device().
• If select_device() returns false, print a descriptive error to stderr listing the available device names from device_configs and return NULL.

4. Add a graceful fallback when the global memory context fails

ggml_backend_rknpu_reg_get_device_count (line 1059) already returns 0 when get_rknpu_memory_context().get_ctx() == 0, which prevents the backend from registering. However, get_rknpu_memory_context() itself (lines 207–218) silently sets mem_ctx = 0 on failure without any log output.

• In the rknpu_memory_context constructor, add a fprintf(stderr, "RKNPU2: Failed to create memory context (ret=%d). NPU backend will be unavailable.\n", ret) when rknn_matmul_create returns a negative value.
• This ensures users on RadxaOS with a broken driver see a clear message rather than a silent no-op.

5. Guard B-matrix memcpy size against buffer overrun

In the B-matrix preparation block (lines 386–388), memcpy(mem->virt_addr, src_ptr, segment_size_bytes) copies segment_size_bytes bytes from the DMA buffer. There is no assertion that total_offset + segment_size_bytes stays within the DMA buffer's allocated size.

• Add a GGML_ASSERT(total_offset + segment_size_bytes <= src0_buf_ctx->dma_buf.size) before the memcpy call to catch any miscalculation early with a clear message rather than a silent memory corruption or kernel fault.

---

sequenceDiagram
    participant User
    participant Backend as ggml_backend_rknpu_device_init_backend
    participant MemCtx as rknpu_memory_context
    participant Compute as ggml_backend_rknpu_graph_compute

    User->>Backend: init (env: RKNN_DEVICE, RKNN_CORE_MASK)
    Backend->>MemCtx: get_rknpu_memory_context()
    alt mem_ctx == 0
        MemCtx-->>Backend: log error, return 0 devices
    else mem_ctx OK
        Backend->>Backend: parse RKNN_DEVICE → select_device()
        Backend->>Backend: parse RKNN_CORE_MASK → store in ctx.core_mask
        Backend-->>User: backend ready
    end

    User->>Compute: graph_compute(cgraph)
    Compute->>Compute: get_matmul_ctx(M,K,N, ctx.core_mask, type)
    Compute->>Compute: rknn_matmul_run() × segments
    alt run fails
        Compute-->>User: log error + return GGML_STATUS_FAILED
    else run OK
        Compute-->>User: GGML_STATUS_SUCCESS
    end

Loading

Import In IDE

🤖 Prompt for AI Agents

## Observations

The three workarounds described in the issue are already partially implemented in `file:ggml/src/ggml-rknpu2/ggml-rknpu2.cpp`:
1. **Single-core forced** — `RKNN_NPU_CORE_0` is hardcoded in `get_matmul_ctx()` (lines 183–195), with the multi-core path commented out.
2. **Static memory context for B-matrix** — `get_rknpu_memory_context()` singleton is used for B-matrix allocations (lines 374–392).
3. **Power-of-2 M quantization** — `rknpu2_calibration::next_power_of_two(M)` is applied to `M_ctx` (line 306).

However, the issue also reports **CPU-only builds crashing**, and the error handling in `ggml_backend_rknpu_graph_compute` is incomplete (e.g., `rknn_matmul_run` failure is silently swallowed at line 519). The device selection is also hardcoded to `"RK3588"` with a `TODO` comment.

## Approach

The plan focuses on three concrete improvements: (1) surface `rknn_matmul_run` errors properly so failures are visible rather than silent, (2) make the NPU core mask and device selection configurable via environment variables to allow runtime tuning without recompilation, and (3) add a graceful fallback path when the memory context fails to initialize, so the backend degrades cleanly instead of crashing.

---

## Implementation Steps

### 1. Propagate `rknn_matmul_run` errors correctly

In `ggml_backend_rknpu_graph_compute` (lines 514–522 of `file:ggml/src/ggml-rknpu2/ggml-rknpu2.cpp`), the `rknn_matmul_run` return value is checked but the error branch is empty. This means EFAULT errors are silently ignored and the function returns `GGML_STATUS_SUCCESS` even after a failed NPU submission.

- Replace the empty `// Handle error` comment with a `fprintf(stderr, ...)` that prints the segment index and the RKNN error code, then set a shared atomic/flag so the outer loop can return `GGML_STATUS_FAILED` after the parallel section completes.
- Since the loop uses `#pragma omp parallel for`, use a local `std::atomic<bool> had_error{false}` declared before the loop, set it inside the error branch, and check it after the loop exits before proceeding to result collection.

### 2. Make NPU core selection configurable via environment variable

Currently `get_matmul_ctx()` unconditionally sets `core_mask = RKNN_NPU_CORE_0` (line 183). The multi-core path is commented out.

- Read an environment variable (e.g., `RKNN_CORE_MASK`) at backend initialization time inside `ggml_backend_rknpu_device_init_backend`.
- Store the resolved `rknn_core_mask` value in `ggml_backend_rknpu_context` as a new field (e.g., `rknn_core_mask core_mask`).
- In `get_matmul_ctx()`, replace the hardcoded `RKNN_NPU_CORE_0` with `this->core_mask`.
- Valid string values to parse: `"0"` → `RKNN_NPU_CORE_0`, `"1"` → `RKNN_NPU_CORE_1`, `"2"` → `RKNN_NPU_CORE_2`, `"auto"` → `RKNN_NPU_CORE_AUTO`. Default to `RKNN_NPU_CORE_0` when the variable is absent or unrecognized (preserving the current safe default for driver 0.9.x).
- Log the selected core mask to `stderr` during initialization so users can confirm the setting.

### 3. Make device selection configurable via environment variable

`ggml_backend_rknpu_device_init_backend` (line 1020) hardcodes `"RK3588"` with a `TODO` comment.

- Read an environment variable (e.g., `RKNN_DEVICE`) at the start of `ggml_backend_rknpu_device_init_backend`.
- Pass the value (defaulting to `"RK3588"` when absent) to `Rknpu2ConfigManager::select_device()`.
- If `select_device()` returns `false`, print a descriptive error to `stderr` listing the available device names from `device_configs` and return `NULL`.

### 4. Add a graceful fallback when the global memory context fails

`ggml_backend_rknpu_reg_get_device_count` (line 1059) already returns `0` when `get_rknpu_memory_context().get_ctx() == 0`, which prevents the backend from registering. However, `get_rknpu_memory_context()` itself (lines 207–218) silently sets `mem_ctx = 0` on failure without any log output.

- In the `rknpu_memory_context` constructor, add a `fprintf(stderr, "RKNPU2: Failed to create memory context (ret=%d). NPU backend will be unavailable.\n", ret)` when `rknn_matmul_create` returns a negative value.
- This ensures users on RadxaOS with a broken driver see a clear message rather than a silent no-op.

### 5. Guard B-matrix `memcpy` size against buffer overrun

In the B-matrix preparation block (lines 386–388), `memcpy(mem->virt_addr, src_ptr, segment_size_bytes)` copies `segment_size_bytes` bytes from the DMA buffer. There is no assertion that `total_offset + segment_size_bytes` stays within the DMA buffer's allocated size.

- Add a `GGML_ASSERT(total_offset + segment_size_bytes <= src0_buf_ctx->dma_buf.size)` before the `memcpy` call to catch any miscalculation early with a clear message rather than a silent memory corruption or kernel fault.

---

```mermaid
sequenceDiagram
    participant User
    participant Backend as ggml_backend_rknpu_device_init_backend
    participant MemCtx as rknpu_memory_context
    participant Compute as ggml_backend_rknpu_graph_compute

    User->>Backend: init (env: RKNN_DEVICE, RKNN_CORE_MASK)
    Backend->>MemCtx: get_rknpu_memory_context()
    alt mem_ctx == 0
        MemCtx-->>Backend: log error, return 0 devices
    else mem_ctx OK
        Backend->>Backend: parse RKNN_DEVICE → select_device()
        Backend->>Backend: parse RKNN_CORE_MASK → store in ctx.core_mask
        Backend-->>User: backend ready
    end

    User->>Compute: graph_compute(cgraph)
    Compute->>Compute: get_matmul_ctx(M,K,N, ctx.core_mask, type)
    Compute->>Compute: rknn_matmul_run() × segments
    alt run fails
        Compute-->>User: log error + return GGML_STATUS_FAILED
    else run OK
        Compute-->>User: GGML_STATUS_SUCCESS
    end
```

---

Execution Information

Branch: rknpu2
Commit: 2806dc3

---

💡 Tips

Supported Commands (Inside Comments)

• Use @traycerai generate to iterate on the previous version of the implementation plan.

Supported Commands (Inside Description)

• Add @traycerai ignore anywhere in the ticket description to prevent this ticket from being processed.
• Add @traycerai branch:<branch-name> anywhere in the ticket description to specify the target branch for the implementation plan.

Community

• Join our Discord Community to get help, request features, and share feedback.
• Follow us on X/Twitter for updates and announcements.