diff --git a/ggml/src/ggml-rknpu2/CONTRIBUTING.md b/ggml/src/ggml-rknpu2/CONTRIBUTING.md new file mode 100644 index 000000000..3836d17da --- /dev/null +++ b/ggml/src/ggml-rknpu2/CONTRIBUTING.md @@ -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. diff --git a/ggml/src/ggml-rknpu2/README.md b/ggml/src/ggml-rknpu2/README.md index f1ef70c0f..6133c3840 100644 --- a/ggml/src/ggml-rknpu2/README.md +++ b/ggml/src/ggml-rknpu2/README.md @@ -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)**
+ 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)**
+ 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)**
+ 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)).
+\*\* **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. diff --git a/ggml/src/ggml-rknpu2/ggml-rknpu2.cpp b/ggml/src/ggml-rknpu2/ggml-rknpu2.cpp index 02383e198..8275a6030 100644 --- a/ggml/src/ggml-rknpu2/ggml-rknpu2.cpp +++ b/ggml/src/ggml-rknpu2/ggml-rknpu2.cpp @@ -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::shared_ptr, TupleHasher> a_buffer_cache; - std::unordered_map, std::shared_ptr, TupleHasher> c_buffer_cache; + std::unordered_map, std::shared_ptr, TupleHasher> c_buffer_cache; std::shared_ptr get_matmul_ctx(int M, int K, int N, int core_id, rknn_matmul_type type) { std::lock_guard 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 static std::shared_ptr get_or_create_npu_buffer( ggml_backend_rknpu_context* backend_ctx, rknn_matmul_ctx matmul_ctx, size_t size, - const std::tuple& key, - std::unordered_map, std::shared_ptr, TupleHasher>& cache + const CacheKeyType& key, + std::unordered_map, TupleHasher>& cache ) { std::lock_guard 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 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 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 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(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(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 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 s_vec; - { - std::lock_guard 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 signed_row(K); - for(int k=0; k 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 s_vec; + if (is_hadamard) { + auto* src0_buf_ctx = (ggml_backend_rknpu_buffer_context*)src0->buffer->context; + std::lock_guard 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 ready_row(K_op); + + if (is_hadamard) { + std::vector signed_row(K); + for(int k=0; knpu_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; nvirt_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 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 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 s_vec(K_op, 1.0f); + std::mt19937 gen(reinterpret_cast(tensor)); + std::uniform_int_distribution distrib(0, 1); + for(int k = 0; k < K_op; ++k) { + s_vec[k] = (distrib(gen) == 0) ? -1.0f : 1.0f; + } + + { + std::lock_guard lock(ctx->mutex); + ctx->hadamard_s_vectors[tensor] = s_vec; + } + + #pragma omp parallel for + for (int n = 0; n < N; ++n) { + std::vector raw_row(K); + dequantize_row(n, raw_row.data()); + + std::vector signed_row(K); + for(int k=0; k quantize_tensor( + const struct ggml_tensor * tensor, + ggml_backend_rknpu_buffer_context * ctx, + const std::vector& 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 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 lock(ctx->mutex); + ctx->quantized_tensor_scales[tensor] = global_scale_b; + } + + // INT8 + if (npu_type == rknpu2_configuration::NPU_TYPE_INT8) { + std::vector 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 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 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 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 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 lock(ctx->mutex); - ctx->quantized_tensor_scales[tensor] = global_scale_b; - } - - // Dequantizing and re-quantizing directly row-by-row - std::vector requantized_data(n_elements); - #pragma omp parallel - { - std::vector 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 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 s_vec(padded_K); - std::mt19937 gen(reinterpret_cast(tensor)); - std::uniform_int_distribution 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 lock(ctx->mutex); - ctx->hadamard_s_vectors[tensor] = s_vec; - } - - // Dequantize original weights - std::vector 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 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 signed_row(K); - for(int k=0; k lock(ctx->mutex); - ctx->quantized_tensor_scales[tensor] = global_scale_b; - } - - // Quantizing data to INT4 - std::vector 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 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 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 lock(ctx->mutex); - ctx->quantized_tensor_scales[tensor] = global_scale_b; - } - - // Dequantizing and re-quantizing directly row-by-row - std::vector requantized_data(n_elements / 2); - #pragma omp parallel - { - std::vector 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 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 fp32_matrix = dequantize_tensor(tensor, ctx, data, K, N, K_op, pipeline->use_hadamard); + std::vector 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; } diff --git a/ggml/src/ggml-rknpu2/rknpu2-configuration.cpp b/ggml/src/ggml-rknpu2/rknpu2-configuration.cpp index 25a0dc4d2..727e745e0 100644 --- a/ggml/src/ggml-rknpu2/rknpu2-configuration.cpp +++ b/ggml/src/ggml-rknpu2/rknpu2-configuration.cpp @@ -3,6 +3,8 @@ #include "rknpu2-configuration.h" #include +#include +#include // --- 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 split_string(const std::string& str, char delimiter) { + std::vector 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* 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* 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 lock(*pattern_mutex); + + // Retrieving the unique tensor name + std::string name = w_tensor->name; + if (name.empty()) { + name = "ptr_" + std::to_string(reinterpret_cast(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 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) --- diff --git a/ggml/src/ggml-rknpu2/rknpu2-configuration.h b/ggml/src/ggml-rknpu2/rknpu2-configuration.h index c58d92110..bd1188955 100644 --- a/ggml/src/ggml-rknpu2/rknpu2-configuration.h +++ b/ggml/src/ggml-rknpu2/rknpu2-configuration.h @@ -7,14 +7,14 @@ #include #include #include +#include #include #include +#include +#include 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 supported_ops; + std::vector hardware_pipelines; + + // Type-specific default patterns mapping + std::map> default_patterns; + + // Custom global pattern + bool use_custom_pattern = false; + std::vector custom_hybrid_pattern; + + // Tensor-Quantization mapping + mutable std::shared_ptr pattern_mutex = std::make_shared(); + mutable std::unordered_map 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* get_active_pattern(int tensor_type) const; }; /**