marfrit/rk-llama.cpp
1
0
Fork 0
Code Issues Pull Requests Actions 33 Packages Projects Releases Wiki Activity

Hardware Pipelines & Hybrid Quantization are implemented

Browse Source
This commit is contained in:
Invisi
2026-03-22 23:36:58 +07:00
parent 5bfeb50726
commit 14d39ea9ce
5 changed files with 685 additions and 497 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-rknpu2/CONTRIBUTING.md
+52
View File
@@ -0,0 +1,52 @@
# Contributing to rk-llama.cpp
This document outlines the project's structure, current development priorities, and workflow guidelines.
## Project Structure
All backend-specific code is contained within the `ggml/src/ggml-rknpu2` directory. Here is a brief overview of the key files and their responsibilities:
* **`ggml-rknpu2.cpp`**
The core GGML backend implementation. It bridges the GGML framework with the Rockchip NPU.
* `ggml_backend_rknpu_buffer_type_alloc_buffer`: Handles the allocation of DMA-Heap memory blocks.
* `ggml_backend_rknpu_buffer_set_tensor`: Dequantizes incoming GGUF weights, requantizes them to the target NPU format, packs them into native RKNN layout, and writes them to the DMA buffer.
* `ggml_backend_rknpu_graph_compute`: Orchestrates the matrix multiplication, assigns workload segments to different NPU cores, and handles activation transformations.
* **`rknpu2-configuration.cpp`**
Defines chip-specific configurations, available hardware pipelines (data types, alignment rules, packing functions), and manages the hybrid quantization patterns.
* **`rknpu2-quantization.cpp`**
Contains mathematical utilities for scaling and symmetric quantization from FP32 to target NPU integer formats (INT8, INT4), as well as dequantization routines.
* **`rknpu2-calibration.cpp`**
Provides statistical methods for finding optimal quantization parameters (e.g., KL-Divergence, Min-MSE) and matrix transformations (like the Fast Walsh-Hadamard Transform).
* **`rknpu2-allocation.cpp`**
A lightweight wrapper for allocating and freeing physically contiguous, zero-copy memory via the Linux DMA-Heap subsystem.
## Future Directions
This section outlines the roadmap for future contributions and improvements.
### Ongoing Needs
Contributions in these areas are always welcome:
* **Bug Fixes:** Any reproducible issues confirmed by the community.
* **Optimizations:** Improvements to calculation accuracy, memory usage reduction, or inference speedups with experimental benchmark data.
### Hot Topics
Here are several "cutting-edge" ideas and active development goals:
* **Support for Other Chipsets:** Expanding configurations in `rknpu2-configuration.cpp` to support other Rockchip SoCs, including RISC-V variants.
* **Advanced Low-Bit Optimizations:** Currently, pure `INT4_STANDARD` produces garbage output. The Hadamard Transform solves the accuracy issue but introduces significant $O(K \log K)$ CPU overhead. Architectural optimizations, faster math routines or advanced algorithms are welcomed to solve the problem.
* **Split Quantization:** Implementing a system where a single weight matrix is split into two: a sparse matrix containing outliers (computed in high-bit pipeline) and a dense matrix for the rest (computed in low-bit pipeline). Orchestrating this efficiently on the NPU is a major milestone.
* **Smart Hybrid Quantization:** Currently, hybrid quantization applies a cyclical pattern across layers. Adding support for "smart" quantization-such as using Regex patterns to target specific, sensitive neural network layers with high-bit pipelines-would vastly improve the performance/accuracy ratio.
## Development Workflow
**Important:** This fork undergoes frequent `rebase` and `push --force` operations against the upstream `llama.cpp` repository to ensure compatibility with the newest architectural updates from the community.
To prevent merge conflicts:
* If you are working on a feature, **please open a Draft Pull Request as early as possible.** This avoids breaking your work during an upstream sync.
* Ensure your code matches the existing C/C++ style of the project.
* Keep your commits clean and descriptive.
ggml/src/ggml-rknpu2/README.md
+102 -63
View File
@@ -26,78 +26,19 @@ cd rk-llama.cpp
```sh
mkdir build && cd build
cmake .. -DLLAMA_RKNPU2=ON
make -j8
make -j4
```
3. Run inference
```sh
# For Dense models
./build/bin/llama-cli -m ~/Projects/gemma-3-1b-it-Q8_0.gguf
./build/bin/llama-cli -m ./gemma-3-1b-it-Q8_0.gguf
# For MoE models
./build/bin/llama-cli -m ~/Projects/LFM2-8B-A1B-Q4_0.gguf --cpu-moe
./build/bin/llama-cli -m ./LFM2-8B-A1B-Q4_0.gguf --cpu-moe
```
## Quantizations
### Weights
Weights are converted based on the input type. Implemented types:
`FP16`
Input **F16** weights are directly used in native **FP16** format.
`INT8`
Input **Q8_0** weights are dequantized to FP32, then re-quantized to a uniform per-tensor **INT8** format.
`INT4`
Input **Q4_0** weights are dequantized, rotated using a randomized Hadamard transform (see [2404.00456](https://arxiv.org/abs/2404.00456)), calibrated using a KL-Divergence (see [2411.02530](https://arxiv.org/abs/2411.02530)), and then re-quantized to per-tensor **INT4**.
### Activations
Activations are converted based on the operation type. Implemented types:
`FP16`
Input **F32** activations are converted to **FP16** format.
`INT8`
Input **F32** activations are quantized to **INT8** using per-channel scaling.
`INT4`
Input **F32** activations are rotated using a Hadamard transform and then quantized to **INT4** using per-channel scaling.
### Results
Results (of a matrix multiplication) are converted based on the output type. Implemented types:
`FP32`
**FP32** results from the NPU are already in **F32** and used directly.
`INT32`
**INT32** results from the NPU are dequantized to **F32** using combined weight and activation scales.
`INT16`
**INT16** results from the NPU are dequantized to **F32** with an additional normalization factor for rotated computations.
## Chipsets
### RK3588
The backend supports the following computation types:
* **W16A16**: FP16 weights & FP16 activations
* **W8A8**: INT8 weights & INT8 activations
* **W4A4**: INT4 weights & INT4 activations
## Benchmarks
The following benchmarks were conducted on an RK3588, comparing the performance, accuracy, and power consumption of the NPU backend against the standard CPU (NEON) backend.
@@ -108,21 +49,119 @@ The following benchmarks were conducted on an RK3588, comparing the performance,
| | | NPU | 🟢 20.74±0.74 | 🟢 432.3±0.6 | 🟡 20.2±0.4 | 🟢 3.2±0.2 |
| | Q8_0 | CPU | 🟢 20.71±0.74 | 🔴 163.4±0.2 | 🟢 40.6±0.1 | 🔴 6.4±0.4 |
| | | NPU | 🟡 22.68±0.82 | 🟢 311.8±2.1 | 🔴 25.4±0.4 | 🟢 3.6±0.2 |
| | Q6_K | CPU | 🟢 21.55±0.77 | 🔴 78.7±0.1 | 🟢 42.5±0.1 | 🔴 6.8±0.4 |
| | | NPU | 🔴 29.52±1.07 | 🟢 142.4±1.3 | 🔴 25.8±1.2 | 🟢 3.8±0.2 |
| | Q4_0 | CPU | 🟢 24.46±0.88 | 🟢 340.4±0.9 | 🟢 55.2±0.1 | 🔴 6.2±0.4 |
| | | NPU | 🔴 74.09±2.87 | 🔴 163.6±0.2 | 🔴 26.7±0.5 | 🟢 4.0±0.2 |
| **Gemma3 1B** | F16 | CPU | 🟢 26.20±1.08 | 🔴 68.5±0.1 | 🟢 11.1±0.1 | 🔴 6.8±0.4 |
| | | NPU | 🟢 26.18±1.07 | 🟢 249.6±0.2 | 🟢 10.8±0.2 | 🟢 2.8±0.2 |
| | Q8_0 | CPU | 🟢 26.08±1.07 | 🔴 73.3±0.1 | 🟢 19.5±0.1 | 🔴 7.4±0.4 |
| | | NPU | 🟡 29.15±1.22 | 🟢 378.6±0.4 | 🟡 16.5±0.3 | 🟢 3.0±0.2 |
| | Q6_K | CPU | 🟢 25.94±1.06 | 🔴 51.9±0.1 | 🟢 18.7±0.1 | 🔴 7.2±0.4 |
| | | NPU | 🟡 38.46±1.68 | 🟢 209.4±3.2 | 🟡 16.6±0.2 | 🟢 3.0±0.2 |
| | Q4_0 | CPU | 🟢 30.77±1.31 | 🟢 164.7±0.2 | 🟢 28.3±0.1 | 🔴 7.0±0.4 |
| | | NPU | 🔴 55.53±2.30 | 🔴 51.4±0.1 | 🔴 16.7±0.3 | 🟢 3.0±0.2 |
| **LFM2 8B A1B** | F16 | CPU | 🟢 15.79±0.58 | 🟡 31.1±2.9 | 🟢 6.8±0.2 | 🟡 7.0±0.6 |
| | | NPU | 🟢 15.82±0.58 | 🟢 38.3±3.2 | 🟢 6.3±0.4 | 🟢 5.8±0.4 |
| | Q8_0 | CPU | 🟢 15.92±0.59 | 🟡 31.7±0.1 | 🟢 12.9±0.1 | 🟡 7.4±0.6 |
| | | NPU | 🟡 16.76±0.62 | 🟢 40.7±0.7 | 🟢 12.5±0.3 | 🟢 5.8±0.4 |
| | Q6_K | CPU | 🟢 15.91±0.58 | 🟡 16.4±0.1 | 🟡 12.9±0.1 | 🟡 7.4±0.6 |
| | | NPU | 🟡 21.16±0.83 | 🟢 21.4±0.1 | 🟢 13.7±0.1 | 🟢 5.8±0.4 |
| | Q4_0 | CPU | 🟢 18.24±0.53 | 🟢 62.2±0.1 | 🟢 22.7±0.1 | 🟡 7.4±0.6 |
| | | NPU | 🟡 26.09±1.06 | 🟡 47.5±0.1 | 🟢 19.0±0.1 | 🟢 5.8±0.4 |
**Legend**: 🟢 Excellent | 🟡 Acceptable | 🔴 Poor
### Methodology
* **Perplexity:** Evaluates accuracy. Measured over 32 chunks of 512 tokens using `wiki.test.raw`. Lower is better.
```sh
taskset -c 4-7 ./build/bin/llama-perplexity -m ./model.gguf -f ./wiki.test.raw -t 4 -b 512 --chunks 32
```
* **PP / TG (tok/s):** Prompt Processing and Text Generation speeds. Evaluated using standard pp512 and tg128. Higher is better.
```sh
taskset -c 4-7 ./build/bin/llama-bench -m ./model.gguf -t 4
```
* **Power (W):** Represents relative active power consumption. Calculated as `Power(TextGeneration) - Power(Idle)`. Lower is better.
## Hybrid Quantization
Hybrid quantization is a technique designed to strike an optimal balance between memory consumption, inference speed, and generation quality.
Instead of forcing the entire model into a single format, the backend processes the model layer by layer, cyclically applying a predefined sequence of hardware pipelines. This approach is entirely transparent and works with any input GGUF weight type. You can find the default pipeline sequences for each chip in the **Chipsets** section below.
### Custom Patterns
Building on this layer-by-layer approach, you can easily override the default behavior to experiment with your own quantization strategies.
By setting the `HYBRID_PATTERN` environment variable, you define a custom sequence of hardware pipelines. The backend will cyclically iterate through this sequence as it loads the model's layers. For example, if you specify two pipelines, Layer 1 will use the first, Layer 2 will use the second, Layer 3 will revert to the first, and so on.
```sh
# Alternates model layers between FP16, INT8 and INT4 NPU pipelines
# This pipeline results in (16 + 8 + 4) / 3 ≈ 9.3 BPW
HYBRID_PATTERN="FP16_STANDARD,INT8_STANDARD,INT4_HADAMARD" ./build/bin/llama-cli -m ./model.gguf
# When the backend encounters a wrong pipeline name, it offloads layer to the default CPU backend
# This pipeline will offload to the NPU around half of the model's layers
HYBRID_PATTERN="CPU_STANDARD,INT8_STANDARD" ./build/bin/llama-cli -m ./model.gguf
```
### Weights Requantizations
It is important to understand how weights are loaded into the NPU. The backend does not execute GGUF formats (like `Q4_0` or `Q8_0`) natively. Instead, input weights are first dequantized to `FP32` on the CPU, calibrated, and then requantized into the NPU's native hardware formats.
Because of this double conversion process, you must carefully pair your input GGUF model with your target NPU pipelines:
1. **Ideal Case (GGUF Precision > NPU Precision)**<br>
Using an `F16` or `Q6_K` model to run `INT8` or `INT4` NPU pipelines. The NPU quantization algorithm receives highly accurate FP32 reference data, allowing it to calculate optimal scales and minimize the final quantization error.
2. **Recommended Case (GGUF Precision ≈ NPU Precision)**<br>
Using `Q8_0` for an `INT8` pipeline, or `Q4_0` for an `INT4` pipeline. The precision levels match closely, keeping the initial GGUF file size small on disk while minimizing further information loss during NPU requantization.
3. **Terrible Case (GGUF Precision < NPU Precision)**<br>
Upscaling a `Q4_0` model to run on an `FP16` or `INT8` NPU pipeline. Information was irreversibly lost when the model was originally compressed to `Q4_0`. Requantizing it to a higher bit-depth cannot restore the lost accuracy, meaning you will get the poor quality of a 4-bit model while wasting the memory bandwidth and compute resources of an 8-bit or 16-bit model.
## Chipsets
The backend configures operations based on hardware pipelines-specific hardware-accelerated paths mapping mathematical operations to native NPU types. Each supported chipset defines its own set of pipelines and default quantization behaviors.
### RK3588
#### Available Pipelines
Below is a comparison of all available hardware pipelines on the RK3588.
The **Perplexity** metrics were measured on the `Granite-4.0-350M-F16` model to isolate the accuracy impact of the NPU's internal quantization algorithms.
| Name | Operation | Perplexity | Notes |
| :--- | :--- | :--- | :--- |
| `FP16_HADAMARD` | FP16xFP16 | 20.74 ± 0.74 | Hadamard Transform\* |
| `FP16_STANDARD` | FP16xFP16 | 20.74 ± 0.74 | - |
| `INT8_HADAMARD` | INT8xINT8 | 20.85 ± 0.74 | Hadamard Transform\* |
| `INT8_STANDARD` | INT8xINT8 | 22.91 ± 0.83 | - |
| `INT4_HADAMARD` | INT4xINT4 | 109.15 ± 4.46 | Hadamard\*, KL-Div\*\* |
| `INT4_STANDARD` | INT4xINT4 | 240048.97 ± 9261.03 | KL-Divergence\*\* |
\* **Hadamard Transform:** Applies a randomized Fast Walsh-Hadamard Transform to smooth out activation outliers before quantization (see [2404.00456](https://arxiv.org/abs/2404.00456)).<br>
\*\* **KL-Divergence:** Uses entropy-based calibration to find the optimal scaling factor for 4-bit weights by minimizing information loss (see [2411.02530](https://arxiv.org/abs/2411.02530)).
> **Note:** As seen in the table, pure `INT4_STANDARD` produces completely broken outputs (extremely high perplexity). This is due to massive outliers being heavily clipped in the limited 4-bit range. The `INT4_HADAMARD` pipeline mitigates this by mathematically distributing the outliers across all channels, making 4-bit inference actually usable, though it introduces some CPU overhead.
#### Default Mappings
If no custom `HYBRID_PATTERN` is provided, the RK3588 backend will automatically map your input GGUF model types to the following default NPU pipelines to provide the best out-of-the-box balance of speed and accuracy:
| Input Weight Type | Default Hardware Pipeline | Bits Per Weight |
| :--- | :--- | :--- |
| `F16` | [`FP16_STANDARD`] | 16 |
| `Q8_0` | [`INT8_STANDARD`] | 8 |
| `Q6_K` | [`INT8_STANDARD`, `INT4_HADAMARD`] | 6 |
| `Q4_0` | [`INT4_HADAMARD`] | 4 |
## Contributing
Feel free to open an issue to discuss a bug or feature, or submit a pull request with your improvements.
Feel free to open an issue to discuss a bug or feature, or submit a pull request with your improvements. Please refer to CONTRIBUTING.md for a breakdown of the project structure, current development goals, and workflow guidelines.
ggml/src/ggml-rknpu2/ggml-rknpu2.cpp
+341 -384
View File
@@ -35,7 +35,7 @@
do { \
int ret = (stmt); \
if (ret < 0) { \
fprintf(stderr,"RKNN error %d at %s:%d: %s\n", ret, \
fprintf(stderr,"RKNN error %d at %s:%d: %s\n", ret, \
__FILE__, __LINE__, msg); \
assert(false); \
} \
@@ -163,7 +163,7 @@ struct ggml_backend_rknpu_context {
// A- and C-matrices cache (from create_mem)
std::unordered_map<std::tuple<int, int, int>, std::shared_ptr<rknn_tensor_mem>, TupleHasher> a_buffer_cache;
std::unordered_map<std::tuple<int, int, int>, std::shared_ptr<rknn_tensor_mem>, TupleHasher> c_buffer_cache;
std::unordered_map<std::tuple<int, int, int, int>, std::shared_ptr<rknn_tensor_mem>, TupleHasher> c_buffer_cache;
std::shared_ptr<rknpu_matmul_context> get_matmul_ctx(int M, int K, int N, int core_id, rknn_matmul_type type) {
std::lock_guard<std::mutex> lock(mutex);
@@ -247,17 +247,20 @@ static void ggml_backend_rknpu_free(ggml_backend_t backend) {
}
// Function for getting buffer from cache or creating new one
template <typename CacheKeyType>
static std::shared_ptr<rknn_tensor_mem> get_or_create_npu_buffer(
ggml_backend_rknpu_context* backend_ctx,
rknn_matmul_ctx matmul_ctx,
size_t size,
const std::tuple<int, int, int>& key,
std::unordered_map<std::tuple<int, int, int>, std::shared_ptr<rknn_tensor_mem>, TupleHasher>& cache
const CacheKeyType& key,
std::unordered_map<CacheKeyType, std::shared_ptr<rknn_tensor_mem>, TupleHasher>& cache
) {
std::lock_guard<std::mutex> lock(backend_ctx->mutex);
auto it = cache.find(key);
if (it != cache.end()) {
return it->second;
if (it->second->size >= size) {
return it->second;
}
}
rknn_tensor_mem* mem = rknn_create_mem(matmul_ctx, size);
@@ -277,36 +280,35 @@ static std::shared_ptr<rknn_tensor_mem> get_or_create_npu_buffer(
static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend, struct ggml_cgraph* cgraph) {
auto* backend_ctx = (ggml_backend_rknpu_context*)backend->context;
// Getting the current device configuration once
const auto& config = rknpu2_configuration::Rknpu2ConfigManager::get_instance().get_current_config();
for (int i = 0; i < cgraph->n_nodes; i++) {
struct ggml_tensor* node = cgraph->nodes[i];
for (int node_i = 0; node_i < cgraph->n_nodes; node_i++) {
struct ggml_tensor* node = cgraph->nodes[node_i];
if (node->op != GGML_OP_MUL_MAT) continue;
const struct ggml_tensor* src0 = node->src[0]; // Weights : (K x N)
const struct ggml_tensor* src1 = node->src[1]; // Activations : (M x K)
struct ggml_tensor* dst = node;
const ggml_type w_type = src0->type;
const int M = (int)src1->ne[1];
const int K = (int)src0->ne[0];
const int N = (int)src0->ne[1];
// Skip zero-dimension matmuls
// Skipping zero-dimension matmuls
if (M == 0 || K == 0 || N == 0) {
continue;
}
const bool is_q4_hadamard = (src0->type == GGML_TYPE_Q4_0);
const int K_op = is_q4_hadamard ? rknpu2_calibration::next_power_of_two(K) : K;
const auto* pipeline = config.resolve_op_support(src0);
if (!pipeline) continue;
const auto* op_support = config.find_op_support(src0->type);
if (!op_support) return GGML_STATUS_FAILED;
const bool is_hadamard = (pipeline->use_hadamard);
const int K_op = is_hadamard ? rknpu2_calibration::next_power_of_two(K) : K;
const int alignment = op_support->n_align;
const rknn_matmul_type matmul_type = pipeline->mm_type;
const int alignment = pipeline->n_align;
auto all_segments = compute_matrix_segments(N, config.core_count, alignment);
@@ -328,18 +330,14 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
// ===========================================
// ========== 1. Preparing contexts ==========
// ===========================================
for (size_t i = 0; i < num_active_segments; ++i) {
const auto& seg = active_segments[i];
matmul_ctxs[i] = backend_ctx->get_matmul_ctx(M, K_op, seg.size_n, seg.core_id, op_support->mm_type);
if (!matmul_ctxs[i] || matmul_ctxs[i]->ctx == 0) return GGML_STATUS_FAILED;
// Skip RKNN matmuls with zero segments
if (matmul_ctxs[i]->info.M == 0 || matmul_ctxs[i]->info.K == 0 || matmul_ctxs[i]->info.N == 0) {
matmul_ctxs[i] = nullptr;
{
for (size_t idx = 0; idx < num_active_segments; ++idx) {
const auto& seg = active_segments[idx];
matmul_ctxs[idx] = backend_ctx->get_matmul_ctx(M, K_op, seg.size_n, seg.core_id, matmul_type);
if (!matmul_ctxs[idx] || matmul_ctxs[idx]->ctx == 0) return GGML_STATUS_FAILED;
}
}
// ===========================================
// ========== 2. Preparing B-matrix ==========
// ===========================================
@@ -348,33 +346,38 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
auto* src0_buf_ctx = (ggml_backend_rknpu_buffer_context*)src0_buffer->context;
size_t src0_base_offset_in_dma = (uintptr_t)src0->data - (uintptr_t)ggml_backend_buffer_get_base(src0_buffer);
size_t type_size_packed;
if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_FP16) type_size_packed = 2;
else if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT8) type_size_packed = 1;
else type_size_packed = 0;
size_t current_offset_in_tensor = 0;
for (size_t i = 0; i < num_active_segments; ++i) {
// const auto& seg = active_segments[i];
size_t segment_size_bytes = matmul_ctxs[i]->io_attr.B.size;
size_t total_offset = src0_base_offset_in_dma + current_offset_in_tensor;
for (const auto& seg : all_segments) {
for (size_t idx = 0; idx < num_active_segments; ++idx) {
if (active_segments[idx].offset_n == seg.offset_n) {
auto& matmul_ctx = matmul_ctxs[idx];
size_t segment_size_bytes = matmul_ctx->io_attr.B.size;
size_t total_offset = src0_base_offset_in_dma + current_offset_in_tensor;
auto cache_key = std::make_pair(src0_buffer, total_offset);
std::lock_guard<std::mutex> lock(backend_ctx->mutex);
auto it = backend_ctx->b_mem_handle_cache.find(cache_key);
auto cache_key = std::make_pair(src0_buffer, total_offset);
{
std::lock_guard<std::mutex> lock(backend_ctx->mutex);
auto it = backend_ctx->b_mem_handle_cache.find(cache_key);
if (it != backend_ctx->b_mem_handle_cache.end()) {
mem_B_segments[i] = it->second;
} else {
rknn_tensor_mem* mem = rknn_create_mem_from_fd(
matmul_ctxs[i]->ctx, src0_buf_ctx->dma_buf.fd,
src0_buf_ctx->dma_buf.virt_addr, segment_size_bytes, total_offset
);
if (!mem) return GGML_STATUS_FAILED;
mem_B_segments[i] = std::shared_ptr<rknn_tensor_mem>(mem, [ctx=matmul_ctxs[i]->ctx](rknn_tensor_mem* m){ if (m) rknn_destroy_mem(ctx,m); });
backend_ctx->b_mem_handle_cache[cache_key] = mem_B_segments[i];
if (it != backend_ctx->b_mem_handle_cache.end()) {
mem_B_segments[idx] = it->second;
} else {
rknn_tensor_mem* mem = rknn_create_mem_from_fd(matmul_ctx->ctx, src0_buf_ctx->dma_buf.fd, src0_buf_ctx->dma_buf.virt_addr, segment_size_bytes, total_offset);
if (!mem) return GGML_STATUS_FAILED;
auto deleter = [ctx = matmul_ctx->ctx](rknn_tensor_mem* m) { if (m) rknn_destroy_mem(ctx, m); };
mem_B_segments[idx] = std::shared_ptr<rknn_tensor_mem>(mem, deleter);
backend_ctx->b_mem_handle_cache[cache_key] = mem_B_segments[idx];
}
RKNN_CHECK(rknn_matmul_set_io_mem(matmul_ctx->ctx, mem_B_segments[idx].get(), &matmul_ctx->io_attr.B), "set_io_mem B segment");
break;
}
RKNN_CHECK(rknn_matmul_set_io_mem(matmul_ctxs[i]->ctx, mem_B_segments[i].get(), &matmul_ctxs[i]->io_attr.B), "set_io_mem B segment");
}
current_offset_in_tensor += segment_size_bytes;
current_offset_in_tensor += type_size_packed > 0 ? (size_t)seg.size_n * K_op * type_size_packed : (size_t)seg.size_n * K_op / 2;
}
}
// ===========================================
@@ -383,7 +386,7 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
std::vector<float> scales_A(M);
float scale_B = 1.0f;
{
auto cache_key = std::make_tuple(M, K, (int)w_type);
auto cache_key = std::make_tuple(M, K_op, (int)pipeline->npu_type_a);
auto& matmul_ctx_0 = matmul_ctxs[0];
mem_A_shared = get_or_create_npu_buffer(backend_ctx, matmul_ctx_0->ctx, matmul_ctx_0->io_attr.A.size, cache_key, backend_ctx->a_buffer_cache);
@@ -393,69 +396,54 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
const int row_stride = (int)(src1->nb[1] / sizeof(float));
void* dst_base = mem_A_shared->virt_addr;
switch (op_support->npu_type_a) {
case rknpu2_configuration::NPU_TYPE_FP16: {
uint16_t* dst_ptr = (uint16_t*)dst_base;
#pragma omp parallel for
for (int m = 0; m < M; ++m) {
const float* src_row = x + (size_t)m * row_stride;
uint16_t* dst_row = dst_ptr + (size_t)m * K;
rknpu2_quantization::convert_fp32_to_fp16(src_row, dst_row, K);
}
break;
}
case rknpu2_configuration::NPU_TYPE_INT8: {
int8_t* dst_ptr = (int8_t*)dst_base;
#pragma omp parallel for
for (int m = 0; m < M; ++m) {
const float* src_row = x + (size_t)m * row_stride;
float amax_m = 0.0f;
for (int k = 0; k < K; ++k) {
amax_m = std::max(amax_m, std::abs(src_row[k]));
}
scales_A[m] = amax_m / 127.0f;
int8_t* dst_row = dst_ptr + (size_t)m * K;
rknpu2_quantization::quantize_fp32_to_int8(src_row, dst_row, K, scales_A[m]);
}
break;
}
case rknpu2_configuration::NPU_TYPE_INT4: {
auto* src0_buf_ctx = (ggml_backend_rknpu_buffer_context*)src0->buffer->context;
std::vector<float> s_vec;
{
std::lock_guard<std::mutex> lock(src0_buf_ctx->mutex);
auto it = src0_buf_ctx->hadamard_s_vectors.find(src0);
GGML_ASSERT(it != src0_buf_ctx->hadamard_s_vectors.end() && "Hadamard 's' vector not found");
s_vec = it->second;
}
uint8_t* dst_ptr = (uint8_t*)dst_base;
#pragma omp parallel for
for (int m = 0; m < M; ++m) {
const float* src_row = x + (size_t)m * row_stride;
std::vector<float> signed_row(K);
for(int k=0; k<K; ++k) signed_row[k] = src_row[k] * s_vec[k];
std::vector<float> rotated_row(K_op);
rknpu2_calibration::hadamard_transform(rotated_row.data(), signed_row.data(), K, K_op);
float amax_m = 0.0f;
for (int k = 0; k < K_op; ++k) amax_m = std::max(amax_m, std::abs(rotated_row[k]));
scales_A[m] = amax_m / 7.0f;
uint8_t* dst_row = dst_ptr + (size_t)m * (K_op / 2);
rknpu2_quantization::quantize_fp32_to_int4_packed(rotated_row.data(), dst_row, K_op, scales_A[m]);
}
break;
}
default:
// This should not be reached if config is correct
return GGML_STATUS_FAILED;
std::vector<float> s_vec;
if (is_hadamard) {
auto* src0_buf_ctx = (ggml_backend_rknpu_buffer_context*)src0->buffer->context;
std::lock_guard<std::mutex> lock(src0_buf_ctx->mutex);
auto it = src0_buf_ctx->hadamard_s_vectors.find(src0);
GGML_ASSERT(it != src0_buf_ctx->hadamard_s_vectors.end() && "Hadamard 's' vector not found");
s_vec = it->second;
}
if (op_support->npu_type_a == rknpu2_configuration::NPU_TYPE_INT8 || op_support->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
#pragma omp parallel for
for (int m = 0; m < M; ++m) {
const float* src_row = x + (size_t)m * row_stride;
std::vector<float> ready_row(K_op);
if (is_hadamard) {
std::vector<float> signed_row(K);
for(int k=0; k<K; ++k) signed_row[k] = src_row[k] * s_vec[k];
rknpu2_calibration::hadamard_transform(ready_row.data(), signed_row.data(), K, K_op);
} else {
memcpy(ready_row.data(), src_row, K * sizeof(float));
}
if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_FP16) {
uint16_t* dst_ptr = (uint16_t*)dst_base;
uint16_t* dst_row = dst_ptr + (size_t)m * K_op;
rknpu2_quantization::convert_fp32_to_fp16(ready_row.data(), dst_row, K_op);
}
else if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT8) {
float amax_m = 0.0f;
for (int k = 0; k < K_op; ++k) amax_m = std::max(amax_m, std::abs(ready_row[k]));
scales_A[m] = amax_m / 127.0f;
int8_t* dst_ptr = (int8_t*)dst_base;
int8_t* dst_row = dst_ptr + (size_t)m * K_op;
rknpu2_quantization::quantize_fp32_to_int8(ready_row.data(), dst_row, K_op, scales_A[m]);
}
else if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
float amax_m = 0.0f;
for (int k = 0; k < K_op; ++k) amax_m = std::max(amax_m, std::abs(ready_row[k]));
scales_A[m] = amax_m / 7.0f;
uint8_t* dst_ptr = (uint8_t*)dst_base;
uint8_t* dst_row = dst_ptr + (size_t)m * (K_op / 2);
rknpu2_quantization::quantize_fp32_to_int4_packed(ready_row.data(), dst_row, K_op, scales_A[m]);
}
}
if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT8 || pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
ggml_backend_buffer_t src0_buffer = src0->buffer;
auto* src0_buf_ctx = (ggml_backend_rknpu_buffer_context*)src0_buffer->context;
{
@@ -468,8 +456,8 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
RKNN_CHECK(rknn_mem_sync(matmul_ctx_0->ctx, mem_A_shared.get(), RKNN_MEMORY_SYNC_TO_DEVICE), "sync A TO_DEVICE");
for (size_t i = 0; i < num_active_segments; i++) {
RKNN_CHECK(rknn_matmul_set_io_mem(matmul_ctxs[i]->ctx, mem_A_shared.get(), &matmul_ctxs[i]->io_attr.A), "set_io_mem A for core");
for (size_t idx = 0; idx < num_active_segments; idx++) {
RKNN_CHECK(rknn_matmul_set_io_mem(matmul_ctxs[idx]->ctx, mem_A_shared.get(), &matmul_ctxs[idx]->io_attr.A), "set_io_mem A for core");
}
}
@@ -477,12 +465,12 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
// ========== 4. Preparing C-matrix ==========
// ===========================================
{
for (size_t i = 0; i < num_active_segments; i++) {
auto& matmul_ctx = matmul_ctxs[i];
auto cache_key = std::make_tuple(M, active_segments[i].size_n, active_segments[i].core_id);
mem_C_segments[i] = get_or_create_npu_buffer(backend_ctx, matmul_ctx->ctx, matmul_ctx->io_attr.C.size, cache_key, backend_ctx->c_buffer_cache);
if (!mem_C_segments[i]) return GGML_STATUS_FAILED;
RKNN_CHECK(rknn_matmul_set_io_mem(matmul_ctx->ctx, mem_C_segments[i].get(), &matmul_ctx->io_attr.C), "set_io_mem C");
for (size_t idx = 0; idx < num_active_segments; idx++) {
auto& matmul_ctx = matmul_ctxs[idx];
auto cache_key = std::make_tuple(M, active_segments[idx].size_n, active_segments[idx].core_id, (int)pipeline->npu_type_c);
mem_C_segments[idx] = get_or_create_npu_buffer(backend_ctx, matmul_ctx->ctx, matmul_ctx->io_attr.C.size, cache_key, backend_ctx->c_buffer_cache);
if (!mem_C_segments[idx]) return GGML_STATUS_FAILED;
RKNN_CHECK(rknn_matmul_set_io_mem(matmul_ctx->ctx, mem_C_segments[idx].get(), &matmul_ctx->io_attr.C), "set_io_mem C");
}
}
@@ -491,8 +479,8 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
// ==========================================
{
#pragma omp parallel for num_threads(num_active_segments)
for (size_t i = 0; i < num_active_segments; i++) {
int ret = rknn_matmul_run(matmul_ctxs[i]->ctx);
for (size_t idx = 0; idx < num_active_segments; idx++) {
int ret = rknn_matmul_run(matmul_ctxs[idx]->ctx);
if (ret != RKNN_SUCC) {
// Handle error
}
@@ -504,33 +492,41 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
// ===========================================
{
float* dst_data = (float*)dst->data;
for (size_t i = 0; i < num_active_segments; i++) {
RKNN_CHECK(rknn_mem_sync(matmul_ctxs[i]->ctx, mem_C_segments[i].get(), RKNN_MEMORY_SYNC_FROM_DEVICE), "sync C FROM_DEVICE");
for (size_t idx = 0; idx < num_active_segments; idx++) {
RKNN_CHECK(rknn_mem_sync(matmul_ctxs[idx]->ctx, mem_C_segments[idx].get(), RKNN_MEMORY_SYNC_FROM_DEVICE), "sync C FROM_DEVICE");
}
const float hadamard_divisor = pipeline->use_hadamard ? (float)K_op : 1.0f;
#pragma omp parallel for
for (int m = 0; m < M; m++) {
switch (op_support->npu_type_c) {
switch (pipeline->npu_type_c) {
case rknpu2_configuration::NPU_TYPE_FP32: {
for (size_t i = 0; i < num_active_segments; i++) {
int N_offset = active_segments[i].offset_n;
int N_segment = active_segments[i].size_n;
float* src_segment_base = (float*)mem_C_segments[i]->virt_addr;
memcpy(dst_data + (size_t)m * N + N_offset,
src_segment_base + (size_t)m * N_segment,
N_segment * sizeof(float));
for (size_t idx = 0; idx < num_active_segments; idx++) {
int N_offset = active_segments[idx].offset_n;
int N_segment = active_segments[idx].size_n;
float* src_segment_base = (float*)mem_C_segments[idx]->virt_addr;
float* dst_ptr = dst_data + (size_t)m * N + N_offset;
float* src_ptr = src_segment_base + (size_t)m * N_segment;
if (pipeline->use_hadamard) {
for(int n=0; n<N_segment; ++n) dst_ptr[n] = src_ptr[n] / hadamard_divisor;
} else {
memcpy(dst_ptr, src_ptr, N_segment * sizeof(float));
}
}
break;
}
case rknpu2_configuration::NPU_TYPE_INT32: {
float dequant_scale = scales_A[m] * scale_B;
for (size_t i = 0; i < num_active_segments; i++) {
int N_offset = active_segments[i].offset_n;
int N_segment = active_segments[i].size_n;
float dequant_scale = (scales_A[m] * scale_B) / hadamard_divisor;
for (size_t idx = 0; idx < num_active_segments; idx++) {
int N_offset = active_segments[idx].offset_n;
int N_segment = active_segments[idx].size_n;
float* dst_ptr = dst_data + (size_t)m * N + N_offset;
int32_t* src_segment_base = (int32_t*)mem_C_segments[i]->virt_addr;
int32_t* src_segment_base = (int32_t*)mem_C_segments[idx]->virt_addr;
int32_t* src_ptr = src_segment_base + (size_t)m * N_segment;
rknpu2_quantization::dequantize_int32_to_fp32(src_ptr, dst_ptr, N_segment, dequant_scale);
}
@@ -538,22 +534,19 @@ static enum ggml_status ggml_backend_rknpu_graph_compute(ggml_backend_t backend,
}
case rknpu2_configuration::NPU_TYPE_INT16: {
float dequant_scale = scales_A[m] * scale_B;
if (op_support->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
dequant_scale /= (float)K_op;
}
float dequant_scale = (scales_A[m] * scale_B) / hadamard_divisor;
for (size_t i = 0; i < num_active_segments; i++) {
int N_offset = active_segments[i].offset_n;
int N_segment = active_segments[i].size_n;
for (size_t idx = 0; idx < num_active_segments; idx++) {
int N_offset = active_segments[idx].offset_n;
int N_segment = active_segments[idx].size_n;
float* dst_ptr = dst_data + (size_t)m * N + N_offset;
int16_t* src_segment_base = (int16_t*)mem_C_segments[i]->virt_addr;
int16_t* src_segment_base = (int16_t*)mem_C_segments[idx]->virt_addr;
int16_t* src_ptr = src_segment_base + (size_t)m * N_segment;
rknpu2_quantization::dequantize_int16_to_fp32(src_ptr, dst_ptr, N_segment, dequant_scale);
}
break;
}
default:
// This should not be reached if config is correct
break;
@@ -587,229 +580,188 @@ static enum ggml_status ggml_backend_rknpu_buffer_init_tensor(ggml_backend_buffe
return GGML_STATUS_SUCCESS;
}
// Function for dequantizing GGUF format to FP32 and optionally applying Hadamard transform
static std::vector<float> dequantize_tensor(
const struct ggml_tensor * tensor,
ggml_backend_rknpu_buffer_context * ctx,
const void * raw_data,
int K, int N, int K_op, bool use_hadamard
) {
std::vector<float> fp32_matrix((size_t)N * K_op);
auto dequantize_row = [&](int n, float* row_out) {
if (tensor->type == GGML_TYPE_F32) {
const float* src = (const float*)raw_data;
memcpy(row_out, src + (size_t)n * K, K * sizeof(float));
} else if (tensor->type == GGML_TYPE_F16) {
const ggml_fp16_t* src = (const ggml_fp16_t*)raw_data;
const ggml_fp16_t* src_row = src + (size_t)n * K;
for (int k = 0; k < K; ++k) row_out[k] = ggml_fp16_to_fp32(src_row[k]);
} else if (tensor->type == GGML_TYPE_Q8_0) {
const block_q8_0* src = (const block_q8_0*)raw_data;
dequantize_row_q8_0(src + (size_t)n * (K / QK8_0), row_out, K);
} else if (tensor->type == GGML_TYPE_Q6_K) {
const block_q6_K* src = (const block_q6_K*)raw_data;
dequantize_row_q6_K(src + (size_t)n * (K / QK_K), row_out, K);
} else if (tensor->type == GGML_TYPE_Q4_0) {
const block_q4_0* src = (const block_q4_0*)raw_data;
dequantize_row_q4_0(src + (size_t)n * (K / QK4_0), row_out, K);
} else {
GGML_ASSERT(false && "Unsupported weight type for NPU pipeline");
}
};
if (use_hadamard) {
std::vector<float> s_vec(K_op, 1.0f);
std::mt19937 gen(reinterpret_cast<uintptr_t>(tensor));
std::uniform_int_distribution<int> distrib(0, 1);
for(int k = 0; k < K_op; ++k) {
s_vec[k] = (distrib(gen) == 0) ? -1.0f : 1.0f;
}
{
std::lock_guard<std::mutex> lock(ctx->mutex);
ctx->hadamard_s_vectors[tensor] = s_vec;
}
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
std::vector<float> raw_row(K);
dequantize_row(n, raw_row.data());
std::vector<float> signed_row(K);
for(int k=0; k<K; ++k) signed_row[k] = raw_row[k] * s_vec[k];
rknpu2_calibration::hadamard_transform(fp32_matrix.data() + (size_t)n * K_op, signed_row.data(), K, K_op);
}
} else {
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
dequantize_row(n, fp32_matrix.data() + (size_t)n * K);
}
}
return fp32_matrix;
}
// Function for quantizing FP32 matrix to target NPU format
static std::vector<uint8_t> quantize_tensor(
const struct ggml_tensor * tensor,
ggml_backend_rknpu_buffer_context * ctx,
const std::vector<float>& fp32_matrix,
int K_op, int N,
rknpu2_configuration::Rknpu2NpuType npu_type
) {
size_t n_elements = (size_t)N * K_op;
// FP16
if (npu_type == rknpu2_configuration::NPU_TYPE_FP16) {
std::vector<uint8_t> npu_bytes(n_elements * sizeof(uint16_t));
uint16_t* fp16_ptr = (uint16_t*)npu_bytes.data();
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
rknpu2_quantization::convert_fp32_to_fp16(
fp32_matrix.data() + (size_t)n * K_op,
fp16_ptr + (size_t)n * K_op,
K_op);
}
return npu_bytes;
}
float amax = 0.0f;
if (npu_type == rknpu2_configuration::NPU_TYPE_INT4) {
amax = rknpu2_calibration::calculate_entropy_amax(fp32_matrix.data(), n_elements);
} else {
#pragma omp parallel for reduction(max:amax)
for (size_t i = 0; i < n_elements; ++i) {
amax = std::max(amax, std::abs(fp32_matrix[i]));
}
}
float quant_divisor = (npu_type == rknpu2_configuration::NPU_TYPE_INT4) ? 7.0f : 127.0f;
float global_scale_b = amax / quant_divisor;
{
std::lock_guard<std::mutex> lock(ctx->mutex);
ctx->quantized_tensor_scales[tensor] = global_scale_b;
}
// INT8
if (npu_type == rknpu2_configuration::NPU_TYPE_INT8) {
std::vector<uint8_t> npu_bytes(n_elements);
int8_t* int8_ptr = (int8_t*)npu_bytes.data();
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
rknpu2_quantization::quantize_fp32_to_int8(
fp32_matrix.data() + (size_t)n * K_op,
int8_ptr + (size_t)n * K_op,
K_op, global_scale_b);
}
return npu_bytes;
}
// INT4
std::vector<uint8_t> npu_bytes(n_elements / 2);
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
rknpu2_quantization::quantize_fp32_to_int4_packed(
fp32_matrix.data() + (size_t)n * K_op,
npu_bytes.data() + (size_t)n * (K_op / 2),
K_op, global_scale_b);
}
return npu_bytes;
}
// Function for splitting into NPU-native layout segments and writing to DMA buffer
static void pack_tensor(
const uint8_t* src_data,
uint8_t* dst_dma_ptr,
int K_op, int N, int core_count,
const rknpu2_configuration::Rknpu2HardwarePipeline * pipeline
) {
auto segments = compute_matrix_segments(N, core_count, pipeline->n_align);
uint8_t* current_write_ptr = dst_dma_ptr;
std::vector<uint8_t> packed_temp;
for (const auto& seg : segments) {
if (seg.size_n == 0) continue;
size_t segment_packed_size = 0;
if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_FP16) {
segment_packed_size = (size_t)seg.size_n * K_op * 2;
} else if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT8) {
segment_packed_size = (size_t)seg.size_n * K_op;
} else if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
segment_packed_size = (size_t)seg.size_n * K_op / 2;
}
packed_temp.resize(segment_packed_size);
pipeline->pack_func(packed_temp.data(), src_data, K_op, N, seg.offset_n, seg.size_n);
memcpy(current_write_ptr, packed_temp.data(), segment_packed_size);
current_write_ptr += segment_packed_size;
}
}
static void ggml_backend_rknpu_buffer_set_tensor(ggml_backend_buffer_t buffer, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
auto * ctx = (ggml_backend_rknpu_buffer_context *) buffer->context;
uint8_t* dma_base = (uint8_t*)ctx->dma_buf.virt_addr;
uint8_t* tensor_dma_ptr = dma_base + ((uintptr_t)tensor->data - (uintptr_t)ggml_backend_buffer_get_base(buffer));
// Getting the current device configuration to drive the packing logic
const auto& config = rknpu2_configuration::Rknpu2ConfigManager::get_instance().get_current_config();
const auto* op_support = config.find_op_support(tensor->type);
const auto* pipeline = config.resolve_op_support(tensor);
// If there is a specific packing function defined for this tensor type, it's a weight matrix
if (op_support && op_support->pack_func) {
if (pipeline && pipeline->pack_func) {
const int K = (int)tensor->ne[0];
const int N = (int)tensor->ne[1];
if (tensor->type == GGML_TYPE_F16) {
auto segments = compute_matrix_segments(N, config.core_count, op_support->n_align);
const int K_op = pipeline->use_hadamard ? rknpu2_calibration::next_power_of_two(K) : K;
uint8_t* current_write_ptr = tensor_dma_ptr;
std::vector<uint8_t> packed_data_temp;
for (const auto& seg : segments) {
if (seg.size_n == 0) continue;
size_t segment_packed_size_bytes = (size_t)seg.size_n * K * sizeof(uint16_t);
packed_data_temp.resize(segment_packed_size_bytes);
op_support->pack_func(
packed_data_temp.data(),
(const uint8_t*)data,
K, N, seg.offset_n, seg.size_n
);
memcpy(current_write_ptr, packed_data_temp.data(), segment_packed_size_bytes);
current_write_ptr += segment_packed_size_bytes;
}
// GGML_TYPE_Q8_0
} else if (tensor->type == GGML_TYPE_Q8_0) {
const block_q8_0* src_blocks = (const block_q8_0*)data;
const size_t n_elements = (size_t)K * N;
// Finding the global scale without storing the full dequantized matrix.
float amax = 0.0f;
#pragma omp parallel
{
std::vector<float> tmp_row(K);
#pragma omp for reduction(max:amax)
for (int n = 0; n < N; ++n) {
dequantize_row_q8_0(src_blocks + (size_t)n * (K / QK8_0), tmp_row.data(), K);
for (int k = 0; k < K; ++k) amax = std::max(amax, std::abs(tmp_row[k]));
}
}
const float global_scale_b = amax / 127.0f;
// Storing it in the buffer context cache
{
std::lock_guard<std::mutex> lock(ctx->mutex);
ctx->quantized_tensor_scales[tensor] = global_scale_b;
}
// Dequantizing and re-quantizing directly row-by-row
std::vector<int8_t> requantized_data(n_elements);
#pragma omp parallel
{
std::vector<float> tmp_row(K);
#pragma omp for
for (int n = 0; n < N; ++n) {
dequantize_row_q8_0(src_blocks + (size_t)n * (K / QK8_0), tmp_row.data(), K);
rknpu2_quantization::quantize_fp32_to_int8(tmp_row.data(), requantized_data.data() + (size_t)n * K, K, global_scale_b);
}
}
// Packing into native format by segments
auto segments = compute_matrix_segments(N, config.core_count, op_support->n_align);
uint8_t* current_write_ptr = tensor_dma_ptr;
std::vector<uint8_t> packed_data_temp;
for (const auto& seg : segments) {
if (seg.size_n == 0) continue;
size_t segment_packed_size = (size_t)seg.size_n * K;
packed_data_temp.resize(segment_packed_size);
op_support->pack_func(
packed_data_temp.data(),
(const uint8_t*)requantized_data.data(),
K, N, seg.offset_n, seg.size_n
);
memcpy(current_write_ptr, packed_data_temp.data(), segment_packed_size);
current_write_ptr += segment_packed_size;
}
// GGML_TYPE_Q4_0
} else if (tensor->type == GGML_TYPE_Q4_0) {
if (op_support->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
const block_q4_0* src_blocks = (const block_q4_0*)data;
const int K = (int)tensor->ne[0];
const int N = (int)tensor->ne[1];
const int padded_K = rknpu2_calibration::next_power_of_two(K);
std::vector<float> s_vec(padded_K);
std::mt19937 gen(reinterpret_cast<uintptr_t>(tensor));
std::uniform_int_distribution<int> distrib(0, 1);
for(int k = 0; k < padded_K; ++k) {
s_vec[k] = (distrib(gen) == 0) ? -1.0f : 1.0f;
}
// Storing the vector for use during activation processing
{
std::lock_guard<std::mutex> lock(ctx->mutex);
ctx->hadamard_s_vectors[tensor] = s_vec;
}
// Dequantize original weights
std::vector<float> dequantized_data((size_t)N * K);
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
dequantize_row_q4_0(src_blocks + (size_t)n * (K / QK4_0), dequantized_data.data() + (size_t)n * K, K);
}
// Apply Hadamard transform using precomputed s vector
const size_t n_elements_padded = (size_t)N * padded_K;
std::vector<float> rotated_b_fp32(n_elements_padded);
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
const float* src_row = dequantized_data.data() + (size_t)n * K;
float* dst_row_rot = rotated_b_fp32.data() + (size_t)n * padded_K;
std::vector<float> signed_row(K);
for(int k=0; k<K; ++k) signed_row[k] = src_row[k] * s_vec[k];
rknpu2_calibration::hadamard_transform(dst_row_rot, signed_row.data(), K, padded_K);
}
// Finding scale on the rotated data
const float amax = rknpu2_calibration::calculate_entropy_amax(rotated_b_fp32.data(), n_elements_padded);
const float global_scale_b = amax / 7.0f;
{
std::lock_guard<std::mutex> lock(ctx->mutex);
ctx->quantized_tensor_scales[tensor] = global_scale_b;
}
// Quantizing data to INT4
std::vector<uint8_t> packed_int4_data(n_elements_padded / 2);
#pragma omp parallel for
for (int n = 0; n < N; ++n) {
rknpu2_quantization::quantize_fp32_to_int4_packed(
rotated_b_fp32.data() + (size_t)n * padded_K,
packed_int4_data.data() + (size_t)n * (padded_K/2),
padded_K, global_scale_b
);
}
// Packing into native RKNN format with padded_K
uint8_t* current_write_ptr = tensor_dma_ptr;
std::vector<uint8_t> packed_data_temp;
for (const auto& seg : compute_matrix_segments(N, config.core_count, op_support->n_align)) {
if (seg.size_n == 0) continue;
size_t segment_packed_size = (size_t)seg.size_n * padded_K / 2;
packed_data_temp.resize(segment_packed_size);
op_support->pack_func(
packed_data_temp.data(),
packed_int4_data.data(),
padded_K, N, seg.offset_n, seg.size_n
);
memcpy(current_write_ptr, packed_data_temp.data(), segment_packed_size);
current_write_ptr += segment_packed_size;
}
} else {
const block_q4_0* src_blocks = (const block_q4_0*)data;
const size_t n_elements = (size_t)K * N;
// Finding the global scale without storing the full dequantized matrix.
float amax = 0.0f;
#pragma omp parallel
{
std::vector<float> tmp_row(K);
#pragma omp for reduction(max:amax)
for (int n = 0; n < N; ++n) {
dequantize_row_q4_0(src_blocks + (size_t)n * (K / QK4_0), tmp_row.data(), K);
for (int k = 0; k < K; ++k) amax = std::max(amax, std::abs(tmp_row[k]));
}
}
const float global_scale_b = amax / 7.0f;
// Storing it in the buffer context cache
{
std::lock_guard<std::mutex> lock(ctx->mutex);
ctx->quantized_tensor_scales[tensor] = global_scale_b;
}
// Dequantizing and re-quantizing directly row-by-row
std::vector<uint8_t> requantized_data(n_elements / 2);
#pragma omp parallel
{
std::vector<float> tmp_row(K);
#pragma omp for
for (int n = 0; n < N; ++n) {
dequantize_row_q4_0(src_blocks + (size_t)n * (K / QK4_0), tmp_row.data(), K);
rknpu2_quantization::quantize_fp32_to_int4_packed(tmp_row.data(), requantized_data.data() + (size_t)n * (K/2), K, global_scale_b);
}
}
// Packing into native format by segments
auto segments = compute_matrix_segments(N, config.core_count, op_support->n_align);
uint8_t* current_write_ptr = tensor_dma_ptr;
std::vector<uint8_t> packed_data_temp;
for (const auto& seg : segments) {
if (seg.size_n == 0) continue;
size_t segment_packed_size = (size_t)seg.size_n * K / 2;
packed_data_temp.resize(segment_packed_size);
op_support->pack_func(
packed_data_temp.data(), requantized_data.data(),
K, N, seg.offset_n, seg.size_n
);
memcpy(current_write_ptr, packed_data_temp.data(), segment_packed_size);
current_write_ptr += segment_packed_size;
}
}
}
// Other tensor types
std::vector<float> fp32_matrix = dequantize_tensor(tensor, ctx, data, K, N, K_op, pipeline->use_hadamard);
std::vector<uint8_t> npu_matrix = quantize_tensor(tensor, ctx, fp32_matrix, K_op, N, pipeline->npu_type_a);
pack_tensor(npu_matrix.data(), tensor_dma_ptr, K_op, N, config.core_count, pipeline);
} else {
memcpy(tensor_dma_ptr + offset, data, size);
}
@@ -874,25 +826,30 @@ static size_t ggml_backend_rknpu_buffer_type_get_alloc_size(ggml_backend_buffer_
// Getting the current device configuration
const auto& config = rknpu2_configuration::Rknpu2ConfigManager::get_instance().get_current_config();
const auto* op_support = config.find_op_support(tensor->type);
// Padding Q4_0 weights to the next power of two for Hadamard Transform.
if (tensor->type == GGML_TYPE_Q4_0) {
if (op_support && op_support->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
const int K = (int)tensor->ne[0];
const int N = (int)tensor->ne[1];
// Defining hardware pipeline for the tensor
const auto* pipeline = config.resolve_op_support(tensor);
const int padded_K = rknpu2_calibration::next_power_of_two(K);
if (pipeline) {
const int K = (int)tensor->ne[0];
const int N = (int)tensor->ne[1];
auto segments = compute_matrix_segments(N, config.core_count, pipeline->n_align);
auto segments = compute_matrix_segments(N, config.core_count, op_support->n_align);
size_t total_size = 0;
for (const auto& seg : segments) {
if (seg.size_n > 0) {
total_size += (size_t)seg.size_n * padded_K / 2;
const int K_op = pipeline->use_hadamard ? rknpu2_calibration::next_power_of_two(K) : K;
size_t total_size = 0;
for (const auto& seg : segments) {
if (seg.size_n > 0) {
if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT4) {
total_size += (size_t)seg.size_n * K_op / 2;
} else if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_INT8) {
total_size += (size_t)seg.size_n * K_op;
} else if (pipeline->npu_type_a == rknpu2_configuration::NPU_TYPE_FP16) {
total_size += (size_t)seg.size_n * K_op * 2;
}
}
return total_size;
}
return total_size;
}
// Fallback to default size calculation for other types.
@@ -951,31 +908,31 @@ static bool ggml_backend_rknpu_device_supports_op(ggml_backend_dev_t dev, const
case GGML_OP_MUL_MAT: {
const struct ggml_tensor * src0 = op->src[0]; // Weights
const struct ggml_tensor * src1 = op->src[1]; // Activations
// Reject zero-dimension ops
// Searching for available hardware pipeline for this tensor
const auto* pipeline = config.resolve_op_support(src0);
if (!pipeline) {
return false;
}
// Rejecting zero-dimension ops
if (src0->ne[0] == 0 || src0->ne[1] == 0 ||
src1->ne[0] == 0 || src1->ne[1] == 0) {
return false;
}
// Finding if there is a supported operation for the given weight type
const auto* op_support = config.find_op_support(src0->type);
if (!op_support) {
return false;
}
// Checking if activation type matches the supported operation
if (src1->type != op_support->type_a) {
if (src1->type != GGML_TYPE_F32) {
return false;
}
// Checking for K alignment
if (src0->ne[0] % op_support->k_align != 0) {
if (src0->ne[0] % pipeline->k_align != 0) {
return false;
}
// Checking for N alignment
if (src0->ne[1] % op_support->n_align != 0) {
if (src0->ne[1] % pipeline->n_align != 0) {
return false;
}
@@ -1004,7 +961,7 @@ static ggml_backend_t ggml_backend_rknpu_device_init_backend(ggml_backend_dev_t
if (!rknpu2_configuration::Rknpu2ConfigManager::get_instance().select_device("RK3588")) return NULL;
ggml_backend_rknpu_context * ctx = new ggml_backend_rknpu_context();
static const struct ggml_backend_i rknpu_backend_interface = {
/* .get_name = */ ggml_backend_rknpu_name,
/* .free = */ ggml_backend_rknpu_free,
@@ -1052,7 +1009,7 @@ static ggml_backend_dev_t ggml_backend_rknpu_reg_get_device(ggml_backend_reg_t r
if (index != 0) {
return NULL;
}
static const struct ggml_backend_buffer_type_i rknpu_buffer_type_interface = {
/* .get_name = */ ggml_backend_rknpu_buffer_type_get_name,
/* .alloc_buffer = */ ggml_backend_rknpu_buffer_type_alloc_buffer,
@@ -1091,7 +1048,7 @@ static ggml_backend_dev_t ggml_backend_rknpu_reg_get_device(ggml_backend_reg_t r
/* .reg = */ reg,
/* .context = */ NULL,
};
if (rknpu_buffer_type.device == NULL) {
rknpu_buffer_type.device = &rknpu_device;
}
ggml/src/ggml-rknpu2/rknpu2-configuration.cpp
+152 -25
View File
@@ -3,6 +3,8 @@
#include "rknpu2-configuration.h"
#include <arm_neon.h>
#include <cstdlib>
#include <sstream>
// --- Anonymous namespace for chip-specific packing functions ---
@@ -112,6 +114,18 @@ void pack_B_rk3588_int4(
} // anonymous namespace
namespace {
// Function for parsing ENV variable
std::vector<std::string> split_string(const std::string& str, char delimiter) {
std::vector<std::string> tokens;
std::string token;
std::istringstream tokenStream(str);
while (std::getline(tokenStream, token, delimiter)) {
if(!token.empty()) tokens.push_back(token);
}
return tokens;
}
} // anonymous namespace
namespace rknpu2_configuration {
@@ -120,43 +134,156 @@ Rknpu2ConfigManager& Rknpu2ConfigManager::get_instance() {
return instance;
}
const std::vector<std::string>* Rknpu2DeviceConfig::get_active_pattern(int tensor_type) const {
auto it = default_patterns.find(tensor_type);
if (it == default_patterns.end()) {
return nullptr;
}
if (use_custom_pattern && !custom_hybrid_pattern.empty()) {
return &custom_hybrid_pattern;
}
return &it->second;
}
const Rknpu2HardwarePipeline* Rknpu2DeviceConfig::resolve_op_support(const struct ggml_tensor* w_tensor) const {
if (!w_tensor) return nullptr;
auto find_pipeline = [this](const std::string& name) -> const Rknpu2HardwarePipeline* {
for (const auto& pipe : hardware_pipelines) {
if (pipe.pipeline_name == name) return &pipe;
}
return nullptr;
};
// Retrieve active quantization pattern based on tensor type (or custom ENV variable)
const std::vector<std::string>* pattern_ptr = get_active_pattern((int)w_tensor->type);
// If no pattern is registered for this type and no global override exists, reject operation
if (!pattern_ptr || pattern_ptr->empty()) {
return nullptr;
}
const auto& pattern = *pattern_ptr;
// Acquiring the lock on the pattern mutex for thread-safe tensor tracking
std::lock_guard<std::mutex> lock(*pattern_mutex);
// Retrieving the unique tensor name
std::string name = w_tensor->name;
if (name.empty()) {
name = "ptr_" + std::to_string(reinterpret_cast<uintptr_t>(w_tensor));
}
// Assigning the next sequence number if this tensor is seen for the first time
if (tensor_sequence_map.find(name) == tensor_sequence_map.end()) {
tensor_sequence_map[name] = global_tensor_counter++;
}
// Selecting the pipeline cyclically based on the defined pattern
int seq_id = tensor_sequence_map[name];
size_t pattern_idx = seq_id % pattern.size();
const std::string& selected_pipeline = pattern[pattern_idx];
const auto* pipeline = find_pipeline(selected_pipeline);
// If no hardware pipeline exists with this name, reject operation
if (!pipeline) {
return nullptr;
}
return pipeline;
}
Rknpu2ConfigManager::Rknpu2ConfigManager() {
// Reading custom hybrid pattern ENV variable
const char* env_pattern = std::getenv("HYBRID_PATTERN");
bool use_custom_pattern = false;
std::vector<std::string> custom_pattern;
if (env_pattern != nullptr) {
custom_pattern = split_string(env_pattern, ',');
use_custom_pattern = true;
}
// --- Define RK3588 Configuration ---
Rknpu2DeviceConfig rk3588_config;
rk3588_config.device_name = "RK3588";
rk3588_config.core_count = 3;
rk3588_config.supported_ops = {
rk3588_config.hardware_pipelines = {
{
/* .type_w = */ GGML_TYPE_F16, // Weights must be converted from F16
/* .type_a = */ GGML_TYPE_F32,
/* .npu_type_a = */ NPU_TYPE_FP16, // Activations must be converted to FP16
/* .npu_type_c = */ NPU_TYPE_FP32, // Result is already in FP32
/* .mm_type = */ RKNN_FLOAT16_MM_FLOAT16_TO_FLOAT32,
/* .k_align = */ 32,
/* .n_align = */ 16,
/* .pack_func = */ pack_B_rk3588_fp16
/* .pipeline_name = */ "FP16_STANDARD",
/* .npu_type_a = */ NPU_TYPE_FP16,
/* .npu_type_c = */ NPU_TYPE_FP32,
/* .mm_type = */ RKNN_FLOAT16_MM_FLOAT16_TO_FLOAT32,
/* .k_align = */ 32,
/* .n_align = */ 16,
/* .pack_func = */ pack_B_rk3588_fp16,
/* .use_hadamard = */ false
},
{
/* .type_w = */ GGML_TYPE_Q8_0, // Weights must be converted from Q8_0
/* .type_a = */ GGML_TYPE_F32,
/* .npu_type_a = */ NPU_TYPE_INT8, // Activations must be converted to INT8
/* .npu_type_c = */ NPU_TYPE_INT32, // Result must be converted from INT32
/* .mm_type = */ RKNN_INT8_MM_INT8_TO_INT32,
/* .k_align = */ 32,
/* .n_align = */ 32,
/* .pack_func = */ pack_B_rk3588_int8
/* .pipeline_name = */ "FP16_HADAMARD",
/* .npu_type_a = */ NPU_TYPE_FP16,
/* .npu_type_c = */ NPU_TYPE_FP32,
/* .mm_type = */ RKNN_FLOAT16_MM_FLOAT16_TO_FLOAT32,
/* .k_align = */ 32,
/* .n_align = */ 16,
/* .pack_func = */ pack_B_rk3588_fp16,
/* .use_hadamard = */ true
},
{
/* .type_w = */ GGML_TYPE_Q4_0, // Weights must be converted from Q4_0
/* .type_a = */ GGML_TYPE_F32,
/* .npu_type_a = */ NPU_TYPE_INT4, // Activations must be converted to INT4
/* .npu_type_c = */ NPU_TYPE_INT16, // Result must be converted from INT16
/* .mm_type = */ RKNN_INT4_MM_INT4_TO_INT16,
/* .k_align = */ 32,
/* .n_align = */ 64,
/* .pack_func = */ pack_B_rk3588_int4
/* .pipeline_name = */ "INT8_STANDARD",
/* .npu_type_a = */ NPU_TYPE_INT8,
/* .npu_type_c = */ NPU_TYPE_INT32,
/* .mm_type = */ RKNN_INT8_MM_INT8_TO_INT32,
/* .k_align = */ 32,
/* .n_align = */ 32,
/* .pack_func = */ pack_B_rk3588_int8,
/* .use_hadamard = */ false
},
{
/* .pipeline_name = */ "INT8_HADAMARD",
/* .npu_type_a = */ NPU_TYPE_INT8,
/* .npu_type_c = */ NPU_TYPE_INT32,
/* .mm_type = */ RKNN_INT8_MM_INT8_TO_INT32,
/* .k_align = */ 32,
/* .n_align = */ 32,
/* .pack_func = */ pack_B_rk3588_int8,
/* .use_hadamard = */ true
},
{
/* .pipeline_name = */ "INT4_STANDARD",
/* .npu_type_a = */ NPU_TYPE_INT4,
/* .npu_type_c = */ NPU_TYPE_INT16,
/* .mm_type = */ RKNN_INT4_MM_INT4_TO_INT16,
/* .k_align = */ 32,
/* .n_align = */ 64,
/* .pack_func = */ pack_B_rk3588_int4,
/* .use_hadamard = */ false
},
{
/* .pipeline_name = */ "INT4_HADAMARD",
/* .npu_type_a = */ NPU_TYPE_INT4,
/* .npu_type_c = */ NPU_TYPE_INT16,
/* .mm_type = */ RKNN_INT4_MM_INT4_TO_INT16,
/* .k_align = */ 32,
/* .n_align = */ 64,
/* .pack_func = */ pack_B_rk3588_int4,
/* .use_hadamard = */ true
}
};
// Assigning custom variables
rk3588_config.use_custom_pattern = use_custom_pattern;
rk3588_config.custom_hybrid_pattern = custom_pattern;
// Defining default quantization sequences for each supported ggml_type
rk3588_config.default_patterns[(int)GGML_TYPE_F16] = {"FP16_STANDARD"};
rk3588_config.default_patterns[(int)GGML_TYPE_Q8_0] = {"INT8_STANDARD"};
rk3588_config.default_patterns[(int)GGML_TYPE_Q6_K] = {"INT8_STANDARD", "INT4_HADAMARD"};
rk3588_config.default_patterns[(int)GGML_TYPE_Q4_0] = {"INT4_HADAMARD"};
device_configs["RK3588"] = rk3588_config;
// --- Define RK3576 Configuration (Placeholder) ---
ggml/src/ggml-rknpu2/rknpu2-configuration.h
+38 -25
View File
@@ -7,14 +7,14 @@
#include <string>
#include <vector>
#include <map>
#include <unordered_map>
#include <functional>
#include <cstdint>
#include <mutex>
#include <memory>
namespace rknpu2_configuration {
// Forward declaration for packing function pointer
struct ggml_tensor;
/**
* @brief Defines a function signature for packing a segment of the weight matrix (B)
* from a standard row-major layout into the device-specific native layout.
@@ -49,49 +49,62 @@ enum Rknpu2NpuType {
};
/**
* @brief Describes a single supported matrix multiplication operation on the NPU.
* @brief Describes a single supported hardware pipeline available on the NPU.
*
* This struct maps a combination of ggml tensor types for weights (src0) and
* activations (src1) to a specific RKNN matmul type. It also specifies the
* hardware alignment requirements for the K and N dimensions of the matrices.
* This struct encapsulates the specific constraints and capabilities of a
* hardware execution path, including required data types, memory alignment
* rules, and the specific packing function needed to prepare tensor for
* the underlying matrix multiplication operation.
*/
struct Rknpu2Operation {
ggml_type type_w; // Weight tensor type
ggml_type type_a; // Activation tensor type
struct Rknpu2HardwarePipeline {
std::string pipeline_name; // The unique hardware pipeline name
Rknpu2NpuType npu_type_a; // The target data type for activations on the NPU
Rknpu2NpuType npu_type_c; // The data type of the result from the NPU
rknn_matmul_type mm_type; // Corresponding RKNN operation type
int k_align; // Required alignment for the K dimension
int n_align; // Required alignment for the N dimension
PackingFunction pack_func; // Function to pack the weight matrix for this op
bool use_hadamard; // Flag for using Hadamard Transform
};
/**
* @brief Holds the complete hardware configuration for a specific Rockchip NPU.
*
* This includes the device name, number of available NPU cores, and a list of
* all matrix multiplication operations supported by the hardware.
* all available hardware pipelines. It also manages layer-by-layer quantization strategies.
*/
struct Rknpu2DeviceConfig {
std::string device_name;
int core_count;
std::vector<Rknpu2Operation> supported_ops;
std::vector<Rknpu2HardwarePipeline> hardware_pipelines;
// Type-specific default patterns mapping
std::map<int, std::vector<std::string>> default_patterns;
// Custom global pattern
bool use_custom_pattern = false;
std::vector<std::string> custom_hybrid_pattern;
// Tensor-Quantization mapping
mutable std::shared_ptr<std::mutex> pattern_mutex = std::make_shared<std::mutex>();
mutable std::unordered_map<std::string, int> tensor_sequence_map;
mutable int global_tensor_counter = 0;
/**
* @brief Finds the corresponding operation configuration for the given weight type.
* @param w_type The ggml_type of the weight tensor.
* @return A pointer to the Rknpu2Operation if found, otherwise nullptr.
* @brief Resolves the appropriate hardware pipeline for a given tensor.
* Applies layer-by-layer cyclical selection if the active pattern has multiple entries.
*/
const Rknpu2Operation* find_op_support(ggml_type w_type) const {
for (const auto& op : supported_ops) {
if (op.type_w == w_type) {
return &op;
}
}
return nullptr;
}
const Rknpu2HardwarePipeline* resolve_op_support(const struct ggml_tensor* w_tensor) const;
private:
/**
* @brief Retrieves the active pattern for a given tensor type.
* Prioritizes custom ENV variable pattern over default mapping.
*/
const std::vector<std::string>* get_active_pattern(int tensor_type) const;
};
/**