Skip to content

Upstream sync: Qwen3.5-MoE support + llama.cpp merge - #1

Merged
KHAEntertainment merged 1304 commits into
rknpu2from
qwen3-support
Mar 18, 2026
Merged

Upstream sync: Qwen3.5-MoE support + llama.cpp merge#1
KHAEntertainment merged 1304 commits into
rknpu2from
qwen3-support

Conversation

@KHAEntertainment

Copy link
Copy Markdown
Owner

Summary

  • Synced with upstream ggml-org/llama.cpp (1,299 commits)
  • Added qwen35moe architecture for Qwen3.5-35B-A3B model
  • Preserved RKNPU2 backend with API compatibility updates
  • Updated documentation (README, CHANGELOG)

Known Issues

  • RKNPU2 backend compiles but has SDK 2.3.x memory allocation incompatibility
  • Error: 'failed to convert fd to handle' during NPU inference
  • Fix planned for follow-up PR

Test Status

  • Mac build: ✅ Compiles successfully
  • Rock 5C build: ✅ Compiles with RKNPU2
  • NPU inference: ⚠️ Blocked by SDK compatibility (to be fixed)

🤖 Generated with Claude Code

jeffbolznv and others added 30 commits February 25, 2026 18:25
* vulkan: check for memory overlap before doing fusion

* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp

* address feedback
Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com>
…org#19826)

* WIP: Add EuroBERT support with autoformatting changes

This commit includes:
- EuroBERT model implementation for GGUF conversion
- C++ backend support for EuroBERT architecture
- Unintended autoformatting changes to Python files

Saving before reverting formatting-only changes.

* feat: add back eos assert when not last token pooling

* feat: removed duplicated code and cleanup

* feat: removed not working architectures and unnecessary check

* fix: typo

* fix: dynamic pooling config

* feat: added an example model for eurobert

* feat: proper llama-vocab implementation for jina-v5

* fix: removed unnecessary comments
Co-authored-by: Roman Marchenko <r.marchenko@ideco.ru>
* ggml-virtgpu-backend: validate the consistency of the received objects

This patch adds consistency checks in the
ggml-virtgpu-backend (running on the host side) to ensure that the
data received from the guest is consistent (valid pointers, valid
sizes and offsets).

* ggml-virtgpu-backend: add fallback/skips for optional ggml backend methods

```
  1. bck->iface.synchronize(bck)
  2. buft->iface.get_alloc_size(buft, op)
  3. buft->iface.get_max_size(buft)
```

these three methods are optional in the GGML interface. `get_max_size`
was already properly defaulted, but `backend sychronize` and `butf
get_max_size` would have segfaulted the backend if not implemented.

* ggml-virtgpu-backend: fix log format missing argument

* ggml-virtgpu-backend: improve the abort message

* ggml-virtgpu-backend: more safety checks

* ggml-virtgpu-backend: new error code

* ggml-virtgpu-backend: initialize all the error codes

* ggml-virtgpu: add a missing comment generated by the code generator

* ggml-virtgpu: add the '[virtgpu]' prefix to the device/buffer names

* ggml-virtgpu: apir_device_buffer_from_ptr: improve the error message

* ggml-virtgpu: shared: make it match the latest api_remoting.h of Virglrenderer APIR

(still unmerged)

* ggml-virtgpu: update the code generator to have dispatch_command_name in a host/guest shared file

* ggml-virtgpu: REMOTE_CALL: fail if the backend returns an error

* docs/backend/VirtGPU.md: indicate that the RAM+VRAM size is limed to 64 GB with libkrun

* ggml-virtgpu: turn off clang-format header ordering for some of the files

Compilation breaks when ordered alphabetically.

* ggml-virtgpu: clang-format

* ggml-virtgpu/backend/shared/api_remoting: better comments for the APIR return codes
* llama: Add option to merge gate and exp weights

* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>

* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>

* update constants.py

* add gate_up for the all MoE models

* convert: simplify merge tensor condition

* update constants.py

* reduce number of models, add create_tensor_gate_up helper

---------

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
llama-perplexity -hf ggml-org/Qwen3-0.6B-GGUF:Q4_0 -f wikitext-2-raw/wiki.test.raw -c 2048 -b 2048 --chunks 2

before this commit:

```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 2.31 seconds per pass - ETA 0.07 minutes
[1]17.3868,[2]22.2199,
Final estimate: PPL = 22.2199 +/- 1.59692

llama_perf_context_print:        load time =     878.56 ms
llama_perf_context_print: prompt eval time =    2037.82 ms /  4096 tokens (    0.50 ms per token,  2009.99 tokens per second)
llama_perf_context_print:        eval time =       0.00 ms /     1 runs   (    0.00 ms per token,      inf tokens per second)
llama_perf_context_print:       total time =    6403.17 ms /  4097 tokens
llama_perf_context_print:    graphs reused =          0
llama_memory_breakdown_print: | memory breakdown [MiB] | total   free    self   model   context   compute    unaccounted |
llama_memory_breakdown_print: |   - Host               |                  845 =   318 +     224 +     302                |
llama_memory_breakdown_print: |   - CPU_REPACK         |                  288 =   288 +       0 +       0                |
llama_memory_breakdown_print: |   - AMX                |                   31 =    31 +       0 +       0                |
```

after this commit:

```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 1.98 seconds per pass - ETA 0.05 minutes
[1]17.2005,[2]21.8220,
Final estimate: PPL = 21.8220 +/- 1.56485

llama_perf_context_print:        load time =     719.23 ms
llama_perf_context_print: prompt eval time =    1676.23 ms /  4096 tokens (    0.41 ms per token,  2443.58 tokens per second)
llama_perf_context_print:        eval time =       0.00 ms /     1 runs   (    0.00 ms per token,      inf tokens per second)
llama_perf_context_print:       total time =    4258.74 ms /  4097 tokens
llama_perf_context_print:    graphs reused =          0
llama_memory_breakdown_print: | memory breakdown [MiB] | total   free    self   model   context   compute    unaccounted |
llama_memory_breakdown_print: |   - Host               |                  845 =   318 +     224 +     302                |
llama_memory_breakdown_print: |   - AMX                |                  319 =   319 +       0 +       0                |
```
(no more CPU_REPACK)

after this commit, disabling amx:

```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 2.34 seconds per pass - ETA 0.07 minutes
[1]17.2005,[2]21.8220,
Final estimate: PPL = 21.8220 +/- 1.56485

llama_perf_context_print:        load time =     841.91 ms
llama_perf_context_print: prompt eval time =    2057.28 ms /  4096 tokens (    0.50 ms per token,  1990.98 tokens per second)
llama_perf_context_print:        eval time =       0.00 ms /     1 runs   (    0.00 ms per token,      inf tokens per second)
llama_perf_context_print:       total time =    6454.51 ms /  4097 tokens
llama_perf_context_print:    graphs reused =          0
llama_memory_breakdown_print: | memory breakdown [MiB] | total   free    self   model   context   compute    unaccounted |
llama_memory_breakdown_print: |   - Host               |                  845 =   318 +     224 +     302                |
llama_memory_breakdown_print: |   - CPU_REPACK         |                  319 =   319 +       0 +       0                |
```
=> same perplexity.

Signed-off-by: Adrien Gallouët <angt@huggingface.co>
- adapt ggml-zendnn.cpp to the new lowoha::matmul interface
- update the ZenDNN git tag in CMake to the latest release (ZenDNN‑2026‑WW08)
- add static lib support in CMake
The binary relies on model files that it tries to find. However, when
configuring the build directory to be parallel to the source tree those
heuristics fail.

This sets the working directory for the test executable to be the
source-tree which resolves this issue.
…gml-org#19926)

* server : support multiple model aliases via comma-separated --alias

* server : update --alias description and regenerate docs

* server : multiple model aliases and tags

- address review feedback from ngxson
- --alias accepts comma-separated values (std::set, no duplicates)
- --tags for informational metadata (not used for routing)
- aliases resolve transparently in router via get_meta/has_model
- /v1/models exposes aliases and tags fields

* regenerate docs

* nits

* server : use first alias as model_name for backward compat

address review feedback from ngxson

* server : add single-model test for aliases and tags
This commit updates the gguf-py package version to 0.18.0 in preperation
of a new release to PyPI.

Refs: ggml-org#19948
This commit changes the runner for the gguf-publish workflow from
ubuntu-slim back to ubuntu-latest, which was updated in Commit
142cbe2 ("ci : use new 1vCPU runner for
lightweight jobs (ggml-org#19107)").

The motivation for this is that the action used in the workflow depends
on the docker daemon, which does not seem not available in the
ubuntu-slim runner. This is currently causing an error in the workflow
and preventing the gguf-publish workflow from running successfully.
Today was the the first time since the original change (I think) that
publish task has been run which may be why the issue was not noticed
before.

Refs: https://github.com/ggml-org/llama.cpp/actions/runs/22481900566
…#19806)

* CUDA: add CDNA3 MFMA support for flash attention MMA kernel

Add MI300X (gfx942) MFMA tensor core flash attention using
v_mfma_f32_16x16x16_f16 (FP16 in, FP32 accumulate).

- Add FATTN_WARP_SIZE=64 for CDNA wavefront64
- Add CDNA config for head sizes 64, 80, 96, 112, 128
- Add FP16 MFMA intrinsic path in mma.cuh
- Add manual V transpose load for MFMA register layout
- Route CDNA to MMA for prompt processing, VEC for token generation
- Fix Q loading and combine stride granularity for non-power-of-2 heads

Benchmarks (Qwen2.5-1.5B Q4_K_M, MI300X):
  pp512  +7%,  pp1024 +13%,  pp2048 +23%,  pp4096 +39%
  tg128  -10% (FA overhead, VEC used for both)

All 2480 flash attention tests pass.

Ref: ggml-org#17917

* address review: replace FATTN_WARP_SIZE with constexpr, improve dispatch

- Replace #define FATTN_WARP_SIZE with constexpr int warp_size =
  ggml_cuda_get_physical_warp_size() in each device function
- Use ne[1]*gqa_ratio threshold for MMA vs tile dispatch. Benchmarked
  crossover on MI300X @ d32768 with power-of-2 GQA models:
    hsk=64  (Llama 1B, gqa=4): MMA wins at eff >= 128 (+11%)
    hsk=128 (Llama 3B, gqa=4): MMA wins at eff >= 128 (+4%)
  Unified threshold: eff_nq >= 128 for all head sizes.
- Remove VEC fallback; small batches fall through to tile kernel

* Update ggml/src/ggml-cuda/fattn.cu

* use ggml_cuda_info().devices warp_size instead of hardcoded check

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Add model metadata loading from huggingface for use with other tests

* Add incremental chunking instead of full redownload, fix caching issue and add warning when it fails

* Add support for split models, load metadata from each individual split file, also avoid mmproj

* Code cleanup, revert incremental downloading

* Only compile when cpp-httplib has SSL support

* Fix formatting
Signed-off-by: Adrien Gallouët <adrien@gallouet.fr>
* vulkan: fix and enable cpy_tensor_async function

* use transfer_queue for async transfers on AMD, synchronize with timeline semaphore

* update offload_op logic

* fix missing transfer submission

* disable async transfer queue on AMD GCN

* revert op batch size change

* fix cpy_tensor_async checks
KHAEntertainment and others added 3 commits March 17, 2026 20:35
- Add GGML_RKNPU2 option to ggml/CMakeLists.txt
- Add ggml_add_backend(RKNPU2) to ggml/src/CMakeLists.txt
- Add RKNPU2 include and registration to ggml-backend-reg.cpp
- Add LLAMA_RKNPU2 deprecation forwarder to root CMakeLists.txt

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- Merged 1,299 commits from ggml-org/llama.cpp
- Added qwen35moe architecture (commit f211220a)
- Preserved RKNPU2 backend with struct field updates
- Updated README and CHANGELOG with fork documentation

RKNPU2 backend compiles but has SDK 2.3.x compatibility issue
(memory allocation to be fixed in follow-up PR)

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
@coderabbitai

coderabbitai Bot commented Mar 18, 2026

Copy link
Copy Markdown

Important

Review skipped

Too many files!

This PR contains 288 files, which is 138 over the limit of 150.

⚙️ Run configuration

Configuration used: Organization UI

Review profile: CHILL

Plan: Pro

Run ID: cc90bb0b-1969-4bce-ba09-d178d2786b0a

📥 Commits

Reviewing files that changed from the base of the PR and between 1b41e28 and cc11ff5.

⛔ Files ignored due to path filters (12)
  • docs/android/imported-into-android-studio.jpg is excluded by !**/*.jpg
  • docs/ops/BLAS.csv is excluded by !**/*.csv
  • docs/ops/CANN.csv is excluded by !**/*.csv
  • docs/ops/CPU.csv is excluded by !**/*.csv
  • docs/ops/CUDA.csv is excluded by !**/*.csv
  • docs/ops/Metal.csv is excluded by !**/*.csv
  • docs/ops/OpenCL.csv is excluded by !**/*.csv
  • docs/ops/SYCL.csv is excluded by !**/*.csv
  • docs/ops/Vulkan.csv is excluded by !**/*.csv
  • docs/ops/WebGPU.csv is excluded by !**/*.csv
  • docs/ops/ZenDNN.csv is excluded by !**/*.csv
  • docs/ops/zDNN.csv is excluded by !**/*.csv
📒 Files selected for processing (288)
  • .devops/cann.Dockerfile
  • .devops/cpu.Dockerfile
  • .devops/cuda-new.Dockerfile
  • .devops/cuda.Dockerfile
  • .devops/intel.Dockerfile
  • .devops/llama-cli-cann.Dockerfile
  • .devops/llama-cpp-cuda.srpm.spec
  • .devops/llama-cpp.srpm.spec
  • .devops/musa.Dockerfile
  • .devops/nix/nixpkgs-instances.nix
  • .devops/nix/package-gguf-py.nix
  • .devops/nix/package.nix
  • .devops/nix/scope.nix
  • .devops/openvino.Dockerfile
  • .devops/rocm.Dockerfile
  • .devops/s390x.Dockerfile
  • .devops/tools.sh
  • .devops/vulkan.Dockerfile
  • .gemini/settings.json
  • .github/ISSUE_TEMPLATE/010-bug-compilation.yml
  • .github/ISSUE_TEMPLATE/011-bug-results.yml
  • .github/ISSUE_TEMPLATE/019-bug-misc.yml
  • .github/actions/linux-setup-openvino/action.yml
  • .github/actions/windows-setup-cuda/action.yml
  • .github/actions/windows-setup-curl/action.yml
  • .github/actions/windows-setup-rocm/action.yml
  • .github/copilot-instructions.md
  • .github/labeler.yml
  • .github/workflows/build-3rd-party.yml
  • .github/workflows/build-android.yml
  • .github/workflows/build-apple.yml
  • .github/workflows/build-cache.yml
  • .github/workflows/build-cann.yml
  • .github/workflows/build-cmake-pkg.yml
  • .github/workflows/build-cross.yml
  • .github/workflows/build-msys.yml
  • .github/workflows/build-riscv-native.yml
  • .github/workflows/build-riscv.yml
  • .github/workflows/build-sanitize.yml
  • .github/workflows/build-self-hosted.yml
  • .github/workflows/build-vulkan.yml
  • .github/workflows/build.yml
  • .github/workflows/check-vendor.yml
  • .github/workflows/close-issue.yml
  • .github/workflows/copilot-setup-steps.yml
  • .github/workflows/docker.yml
  • .github/workflows/editorconfig.yml
  • .github/workflows/gguf-publish.yml
  • .github/workflows/labeler.yml
  • .github/workflows/pre-tokenizer-hashes.yml
  • .github/workflows/python-check-requirements.yml
  • .github/workflows/python-lint.yml
  • .github/workflows/python-type-check.yml
  • .github/workflows/release.yml
  • .github/workflows/server-sanitize.yml
  • .github/workflows/server-self-hosted.yml
  • .github/workflows/server-webui.yml
  • .github/workflows/server.yml
  • .github/workflows/update-ops-docs.yml
  • .github/workflows/winget.yml
  • .gitignore
  • AGENTS.md
  • AUTHORS
  • CHANGELOG.md
  • CLAUDE.md
  • CMakeLists.txt
  • CODEOWNERS
  • CONTRIBUTING.md
  • LICENSE
  • README.md
  • SECURITY.md
  • benches/dgx-spark/dgx-spark.md
  • benches/mac-m2-ultra/mac-m2-ultra.md
  • benches/nemotron/nemotron-dgx-spark.md
  • build-xcframework.sh
  • ci/run.sh
  • cmake/build-info.cmake
  • cmake/common.cmake
  • cmake/download-models.cmake
  • cmake/license.cmake
  • common/CMakeLists.txt
  • common/arg.cpp
  • common/arg.h
  • common/chat-auto-parser-generator.cpp
  • common/chat-auto-parser-helpers.cpp
  • common/chat-auto-parser-helpers.h
  • common/chat-auto-parser.h
  • common/chat-diff-analyzer.cpp
  • common/chat-parser-xml-toolcall.cpp
  • common/chat-parser-xml-toolcall.h
  • common/chat-parser.cpp
  • common/chat-parser.h
  • common/chat-peg-parser.cpp
  • common/chat-peg-parser.h
  • common/chat.cpp
  • common/chat.h
  • common/common.cpp
  • common/common.h
  • common/console.cpp
  • common/console.h
  • common/debug.cpp
  • common/debug.h
  • common/download.cpp
  • common/download.h
  • common/http.h
  • common/jinja/README.md
  • common/jinja/caps.cpp
  • common/jinja/caps.h
  • common/jinja/lexer.cpp
  • common/jinja/lexer.h
  • common/jinja/parser.cpp
  • common/jinja/parser.h
  • common/jinja/runtime.cpp
  • common/jinja/runtime.h
  • common/jinja/string.cpp
  • common/jinja/string.h
  • common/jinja/utils.h
  • common/jinja/value.cpp
  • common/jinja/value.h
  • common/json-partial.h
  • common/json-schema-to-grammar.cpp
  • common/json-schema-to-grammar.h
  • common/llguidance.cpp
  • common/log.cpp
  • common/log.h
  • common/ngram-cache.cpp
  • common/ngram-cache.h
  • common/ngram-map.cpp
  • common/ngram-map.h
  • common/ngram-mod.cpp
  • common/ngram-mod.h
  • common/peg-parser.cpp
  • common/peg-parser.h
  • common/preset.cpp
  • common/preset.h
  • common/reasoning-budget.cpp
  • common/reasoning-budget.h
  • common/regex-partial.cpp
  • common/sampling.cpp
  • common/sampling.h
  • common/speculative.cpp
  • common/speculative.h
  • common/unicode.cpp
  • common/unicode.h
  • convert_hf_to_gguf.py
  • convert_hf_to_gguf_update.py
  • convert_lora_to_gguf.py
  • docs/android.md
  • docs/autoparser.md
  • docs/backend/CANN.md
  • docs/backend/OPENCL.md
  • docs/backend/OPENVINO.md
  • docs/backend/SYCL.md
  • docs/backend/VirtGPU.md
  • docs/backend/VirtGPU/configuration.md
  • docs/backend/VirtGPU/development.md
  • docs/backend/ZenDNN.md
  • docs/backend/snapdragon/CMakeUserPresets.json
  • docs/backend/snapdragon/README.md
  • docs/backend/snapdragon/developer.md
  • docs/backend/snapdragon/windows.md
  • docs/backend/zDNN.md
  • docs/build-riscv64-spacemit.md
  • docs/build-s390x.md
  • docs/build.md
  • docs/development/HOWTO-add-model.md
  • docs/development/parsing.md
  • docs/docker.md
  • docs/function-calling.md
  • docs/multimodal/MobileVLM.md
  • docs/multimodal/minicpmo2.6.md
  • docs/multimodal/minicpmo4.0.md
  • docs/multimodal/minicpmv2.5.md
  • docs/multimodal/minicpmv2.6.md
  • docs/multimodal/minicpmv4.0.md
  • docs/multimodal/minicpmv4.5.md
  • docs/ops.md
  • docs/preset.md
  • docs/speculative.md
  • examples/CMakeLists.txt
  • examples/batched/README.md
  • examples/batched/batched.cpp
  • examples/convert-llama2c-to-ggml/convert-llama2c-to-ggml.cpp
  • examples/debug/CMakeLists.txt
  • examples/debug/README.md
  • examples/debug/debug.cpp
  • examples/deprecation-warning/README.md
  • examples/deprecation-warning/deprecation-warning.cpp
  • examples/diffusion/README.md
  • examples/diffusion/diffusion-cli.cpp
  • examples/embedding/embedding.cpp
  • examples/eval-callback/CMakeLists.txt
  • examples/eval-callback/eval-callback.cpp
  • examples/gen-docs/gen-docs.cpp
  • examples/gguf-hash/gguf-hash.cpp
  • examples/gguf/gguf.cpp
  • examples/idle/CMakeLists.txt
  • examples/idle/README.md
  • examples/idle/idle.cpp
  • examples/json_schema_to_grammar.py
  • examples/llama.android/app/build.gradle.kts
  • examples/llama.android/app/proguard-rules.pro
  • examples/llama.android/app/src/main/AndroidManifest.xml
  • examples/llama.android/app/src/main/java/com/example/llama/Downloadable.kt
  • examples/llama.android/app/src/main/java/com/example/llama/MainActivity.kt
  • examples/llama.android/app/src/main/java/com/example/llama/MainViewModel.kt
  • examples/llama.android/app/src/main/java/com/example/llama/MessageAdapter.kt
  • examples/llama.android/app/src/main/java/com/example/llama/ui/theme/Color.kt
  • examples/llama.android/app/src/main/java/com/example/llama/ui/theme/Theme.kt
  • examples/llama.android/app/src/main/java/com/example/llama/ui/theme/Type.kt
  • examples/llama.android/app/src/main/res/drawable/bg_assistant_message.xml
  • examples/llama.android/app/src/main/res/drawable/bg_user_message.xml
  • examples/llama.android/app/src/main/res/drawable/outline_folder_open_24.xml
  • examples/llama.android/app/src/main/res/drawable/outline_send_24.xml
  • examples/llama.android/app/src/main/res/layout/activity_main.xml
  • examples/llama.android/app/src/main/res/layout/item_message_assistant.xml
  • examples/llama.android/app/src/main/res/layout/item_message_user.xml
  • examples/llama.android/app/src/main/res/values/strings.xml
  • examples/llama.android/app/src/main/res/values/themes.xml
  • examples/llama.android/build.gradle.kts
  • examples/llama.android/gradle.properties
  • examples/llama.android/gradle/libs.versions.toml
  • examples/llama.android/gradle/wrapper/gradle-wrapper.properties
  • examples/llama.android/lib/.gitignore
  • examples/llama.android/lib/build.gradle.kts
  • examples/llama.android/lib/consumer-rules.pro
  • examples/llama.android/lib/proguard-rules.pro
  • examples/llama.android/lib/src/androidTest/java/android/llama/cpp/ExampleInstrumentedTest.kt
  • examples/llama.android/lib/src/main/AndroidManifest.xml
  • examples/llama.android/lib/src/main/cpp/CMakeLists.txt
  • examples/llama.android/lib/src/main/cpp/ai_chat.cpp
  • examples/llama.android/lib/src/main/cpp/logging.h
  • examples/llama.android/lib/src/main/java/com/arm/aichat/AiChat.kt
  • examples/llama.android/lib/src/main/java/com/arm/aichat/InferenceEngine.kt
  • examples/llama.android/lib/src/main/java/com/arm/aichat/gguf/FileType.kt
  • examples/llama.android/lib/src/main/java/com/arm/aichat/gguf/GgufMetadata.kt
  • examples/llama.android/lib/src/main/java/com/arm/aichat/gguf/GgufMetadataReader.kt
  • examples/llama.android/lib/src/main/java/com/arm/aichat/internal/InferenceEngineImpl.kt
  • examples/llama.android/lib/src/main/java/com/arm/aichat/internal/gguf/GgufMetadataReaderImpl.kt
  • examples/llama.android/lib/src/test/java/android/llama/cpp/ExampleUnitTest.kt
  • examples/llama.android/llama/build.gradle.kts
  • examples/llama.android/llama/consumer-rules.pro
  • examples/llama.android/llama/src/main/cpp/CMakeLists.txt
  • examples/llama.android/llama/src/main/cpp/llama-android.cpp
  • examples/llama.android/llama/src/main/java/android/llama/cpp/LLamaAndroid.kt
  • examples/llama.android/settings.gradle.kts
  • examples/llama.vim
  • examples/lookahead/lookahead.cpp
  • examples/lookup/lookup-create.cpp
  • examples/lookup/lookup-merge.cpp
  • examples/lookup/lookup-stats.cpp
  • examples/lookup/lookup.cpp
  • examples/model-conversion/Makefile
  • examples/model-conversion/README.md
  • examples/model-conversion/logits.cpp
  • examples/model-conversion/scripts/causal/compare-embeddings-logits.sh
  • examples/model-conversion/scripts/causal/compare-logits.py
  • examples/model-conversion/scripts/causal/convert-model.sh
  • examples/model-conversion/scripts/causal/modelcard.template
  • examples/model-conversion/scripts/causal/run-casual-gen-embeddings-org.py
  • examples/model-conversion/scripts/causal/run-converted-model-embeddings-logits.sh
  • examples/model-conversion/scripts/causal/run-converted-model.sh
  • examples/model-conversion/scripts/causal/run-org-model.py
  • examples/model-conversion/scripts/embedding/compare-embeddings-logits.sh
  • examples/model-conversion/scripts/embedding/run-converted-model.sh
  • examples/model-conversion/scripts/embedding/run-original-model.py
  • examples/model-conversion/scripts/utils/__init__.py
  • examples/model-conversion/scripts/utils/check-nmse.py
  • examples/model-conversion/scripts/utils/common.py
  • examples/model-conversion/scripts/utils/compare_tokens.py
  • examples/model-conversion/scripts/utils/inspect-org-model.py
  • examples/model-conversion/scripts/utils/perplexity-gen.sh
  • examples/model-conversion/scripts/utils/perplexity-run-simple.sh
  • examples/model-conversion/scripts/utils/perplexity-run.sh
  • examples/model-conversion/scripts/utils/quantize.sh
  • examples/model-conversion/scripts/utils/run-embedding-server.sh
  • examples/model-conversion/scripts/utils/semantic_check.py
  • examples/parallel/parallel.cpp
  • examples/passkey/passkey.cpp
  • examples/retrieval/retrieval.cpp
  • examples/save-load-state/save-load-state.cpp
  • examples/simple-chat/simple-chat.cpp
  • examples/simple-cmake-pkg/README.md
  • examples/simple/simple.cpp
  • examples/speculative-simple/speculative-simple.cpp
  • examples/speculative/speculative.cpp
  • examples/sycl/README.md
  • examples/sycl/build.sh

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

  • 🔍 Trigger review
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests
  • Commit unit tests in branch qwen3-support
📝 Coding Plan
  • Generate coding plan for human review comments

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands and usage tips.

Tip

You can disable the changed files summary in the walkthrough.

Disable the reviews.changed_files_summary setting to disable the changed files summary in the walkthrough.

Changes:
- Replace rknn_create_mem_from_fd with rknn_create_mem for B-matrix
- Add CMA heap fallback in DMA allocation
- Document known SDK 2.3.x GEM handle allocation issue

Investigation findings:
- Driver 0.9.8 and runtime 2.3.2 detected and working
- Memory allocation succeeds during model load
- Allocation fails during inference with errno 14 (EFAULT)
- Issue appears when matmul contexts are destroyed/recreated
  with different batch sizes (M parameter)

Root cause under investigation - likely driver/SDK state issue.

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- Force single NPU core (CORE_0) to avoid spinlock recursion bug
- Quantize M to power of 2 for context caching
- Use static memory context for B-matrix allocations

These workarounds address GEM handle exhaustion (EFAULT errno 14)
but the underlying kernel driver bug remains unresolved.

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
@KHAEntertainment
KHAEntertainment merged commit 2806dc3 into rknpu2 Mar 18, 2026
56 of 73 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.