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

IOMMU Domain Management is implemented

Browse Source
This commit is contained in:
Invisi
2026-04-04 17:50:36 +07:00
parent 14d39ea9ce
commit 1d3c0f5829
8 changed files with 825 additions and 668 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-rknpu2/CMakeLists.txt
-1
View File
@@ -4,7 +4,6 @@ find_package(OpenMP REQUIRED)
# Adding backend files
ggml_add_backend_library(ggml-rknpu2
ggml-rknpu2.cpp
rknpu2-allocation.cpp
rknpu2-quantization.cpp
rknpu2-calibration.cpp
rknpu2-configuration.cpp
ggml/src/ggml-rknpu2/CONTRIBUTING.md
-3
View File
@@ -21,9 +21,6 @@ All backend-specific code is contained within the `ggml/src/ggml-rknpu2` directo
* **`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.
ggml/src/ggml-rknpu2/README.md
+27 -8
View File
@@ -32,17 +32,19 @@ make -j4
3. Run inference
```sh
# For Dense models
./build/bin/llama-cli -m ./gemma-3-1b-it-Q8_0.gguf
# Set new max limit for open files
ulimit -n 65536
# For MoE models
./build/bin/llama-cli -m ./LFM2-8B-A1B-Q4_0.gguf --cpu-moe
# Start chat with the model
./build/bin/llama-cli -m ./gemma-3-1b-it-Q8_0.gguf
```
## 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.
### Dense Models
| Model | Type | Backend | Perplexity | PP (tok/s) | TG (tok/s) | Power (W) |
| :--- | :--- | :--- | :--- | :--- | :--- | :--- |
| **Granite4.0 350M** | F16 | CPU | 🟢 20.73±0.74 | 🔴 154.1±0.1 | 🟢 25.8±0.1 | 🔴 6.4±0.4 |
@@ -61,14 +63,31 @@ The following benchmarks were conducted on an RK3588, comparing the performance,
| | | 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 |
| **Ministral3 3B** | Q8_0 | CPU | 🟢 7.67±0.21 | 🔴 17.5±0.1 | 🟡 6.1±0.1 | 🔴 7.2±0.6 |
| | | NPU | 🟢 7.84±0.22 | 🟢 120.9±0.6 | 🟢 6.9±0.1 | 🟢 3.2±0.2 |
| | Q4_0 | CPU | 🟢 7.93±0.22 | 🟢 15.1±0.1 | 🟢 8.3±0.2 | 🔴 7.8±0.6 |
| | | NPU | 🔴 32.37±1.16 | 🟡 13.8±0.1 | 🟡 7.2±0.1 | 🟢 3.2±0.2 |
| **Qwen3.5 9B** | Q8_0 | CPU | 🟢 6.82±0.19 | 🔴 7.2±0.1 | 🟡 2.5±0.1 | 🔴 7.4±0.6 |
| | | NPU | 🟢 7.14±0.20 | 🟢 44.6±0.1 | 🟢 3.2±0.1 | 🟢 3.2±0.2 |
| | Q4_0 | CPU | 🟢 7.10±0.20 | 🟡 6.5±0.1 | 🟡 3.6±0.1 | 🔴 8.0±0.6 |
| | | NPU | 🟡 14.67±0.48 | 🟢 7.3±0.1 | 🟢 4.4±0.1 | 🟢 4.0±0.2 |
### MoE Models
| Model | Type | Backend | Perplexity | PP (tok/s) | TG (tok/s) | Power (W) |
| :--- | :--- | :--- | :--- | :--- | :--- | :--- |
| **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 |
| | 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 |
| | | NPU | 🟡 26.09±1.06 | 🟡 47.5±0.1 | 🟡 19.0±0.1 | 🟢 5.8±0.4 |
| **Gemma4 26B A4B** | Q8_0 | CPU | 🟢 169.08±10.50 | 🟡 9.8±0.4 | 🟢 4.1±0.2 | 🔴 7.2±0.6 |
| | | NPU | 🟢 215.18±13.84 | 🟢 13.4±0.4 | 🟡 2.9±0.1 | 🟢 4.2±0.4 |
| | Q4_0 | CPU | 🟢 201.44±12.77 | 🟢 11.3±0.4 | 🟢 6.1±0.1 | 🔴 7.4±0.6 |
| | | NPU | 🟡 635.77±44.36 | 🟢 12.9±0.5 | 🟡 5.6±0.1 | 🟢 4.4±0.4 |
**Legend**: 🟢 Excellent | 🟡 Acceptable | 🔴 Poor
@@ -142,8 +161,8 @@ The **Perplexity** metrics were measured on the `Granite-4.0-350M-F16` model to
| `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\*\* |
| `INT8_STANDARD` | INT8xINT8 | 22.50 ± 0.81 | - |
| `INT4_HADAMARD` | INT4xINT4 | 86.16 ± 3.42 | 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>
ggml/src/ggml-rknpu2/ggml-rknpu2.cpp
+754 -508
View File
File diff suppressed because it is too large Load Diff
ggml/src/ggml-rknpu2/rknpu2-allocation.cpp
-77
View File
@@ -1,77 +0,0 @@
#include "rknpu2-allocation.h"
#include <sys/ioctl.h>
#include <sys/mman.h>
#include <fcntl.h>
#include <unistd.h>
#include <string.h>
#include <cerrno>
#include <cstdio>
// --- Anonymous namespace for implementation details ---
namespace {
// --- DMA-Heap specific structures and IOCTL commands ---
struct dma_heap_allocation_data {
uint64_t len;
uint32_t fd;
uint32_t fd_flags;
uint64_t heap_flags;
};
#define DMA_HEAP_IOC_MAGIC 'H'
#define DMA_HEAP_IOCTL_ALLOC _IOWR(DMA_HEAP_IOC_MAGIC, 0x0, struct dma_heap_allocation_data)
} // anonymous namespace
namespace rknpu2_allocation {
DmaBuffer alloc(size_t size) {
DmaBuffer buffer;
buffer.size = size;
const char* path = "/dev/dma_heap/system";
int dma_heap_fd = open(path, O_RDWR);
if (dma_heap_fd < 0) {
fprintf(stderr, "RKNPU_DMA_ALLOC: Failed to open %s: %s\n", path, strerror(errno));
return buffer;
}
dma_heap_allocation_data buf_data;
memset(&buf_data, 0, sizeof(buf_data));
buf_data.len = size;
buf_data.fd_flags = O_CLOEXEC | O_RDWR;
if (ioctl(dma_heap_fd, DMA_HEAP_IOCTL_ALLOC, &buf_data) < 0) {
fprintf(stderr, "RKNPU_DMA_ALLOC: ioctl DMA_HEAP_IOCTL_ALLOC failed: %s\n", strerror(errno));
close(dma_heap_fd);
return buffer;
}
close(dma_heap_fd);
buffer.fd = buf_data.fd;
buffer.virt_addr = mmap(nullptr, size, PROT_READ | PROT_WRITE, MAP_SHARED, buffer.fd, 0);
if (buffer.virt_addr == MAP_FAILED) {
fprintf(stderr, "RKNPU_DMA_ALLOC: mmap failed: %s\n", strerror(errno));
close(buffer.fd);
buffer.fd = -1;
buffer.virt_addr = nullptr;
}
return buffer;
}
void free(const DmaBuffer& buffer) {
if (buffer.virt_addr != nullptr) {
munmap(buffer.virt_addr, buffer.size);
}
if (buffer.fd >= 0) {
close(buffer.fd);
}
}
} // namespace rknpu2_allocation
ggml/src/ggml-rknpu2/rknpu2-allocation.h
-43
View File
@@ -1,43 +0,0 @@
#pragma once
#include <cstddef>
#include <cstdint>
/**
* @brief Provides a simple interface for allocating and freeing physically contiguous
* memory using the Linux DMA-Heap subsystem, which is required for zero-copy
* operations with the Rockchip NPU.
*/
namespace rknpu2_allocation {
/**
* @brief Represents a block of memory allocated from the DMA-Heap.
*
* This struct holds the file descriptor, virtual address, and size of the
* allocated buffer. It is the handle used for all operations.
*/
struct DmaBuffer {
int fd = -1; // File descriptor for the DMA buffer
void* virt_addr = nullptr; // User-space virtual address mapped to the buffer
size_t size = 0; // Size of the allocation in bytes
};
/**
* @brief Allocates a physically contiguous block of memory from the system DMA heap.
*
* @param size The number of bytes to allocate.
* @return A DmaBuffer struct. If allocation fails, `fd` will be -1 and
* `virt_addr` will be nullptr.
*/
DmaBuffer alloc(size_t size);
/**
* @brief Frees a previously allocated DMA buffer.
*
* This function unmaps the virtual address and closes the file descriptor.
*
* @param buffer The DmaBuffer to be freed.
*/
void free(const DmaBuffer& buffer);
} // namespace rknpu2_allocation
ggml/src/ggml-rknpu2/rknpu2-configuration.cpp
+32 -22
View File
@@ -13,25 +13,26 @@ namespace {
// Packing KxN FP16 (row-major: idx [k,n] -> k*N + n) into native RKNN for RK3588: (N/16, K/32, 16, 32)
void pack_B_rk3588_fp16(
uint8_t* dst_u8, const uint8_t* src_u8,
int K, int N_total, int n_offset, int n_segment) {
int K_total, int k_offset, int k_segment,
int N_total, int n_offset, int n_segment) {
auto dst = reinterpret_cast<uint16_t*>(dst_u8);
auto src = reinterpret_cast<const uint16_t*>(src_u8);
GGML_ASSERT(K % 32 == 0 && N_total > 0 && K > 0);
GGML_ASSERT(k_segment % 32 == 0 && N_total > 0 && K_total > 0);
GGML_ASSERT(n_offset % 16 == 0 && n_segment % 16 == 0 && n_offset + n_segment <= N_total);
const size_t s0 = (size_t)(K / 32) * 16 * 32;
const size_t s0 = (size_t)(k_segment / 32) * 16 * 32;
const size_t s1 = 16 * 32;
const size_t s2 = 32;
for (int i = 0; i < n_segment / 16; ++i) {
for (int j = 0; j < K / 32; ++j) {
for (int j = 0; j < k_segment / 32; ++j) {
const size_t dst_block = (size_t) i * s0 + (size_t) j * s1;
for (int ii = 0; ii < 16; ++ii) {
const size_t n_global = (size_t)n_offset + (size_t)i * 16 + (size_t)ii;
const uint16_t * src_ptr = src + n_global * K + j * 32;
const uint16_t * src_ptr = src + n_global * K_total + k_offset + j * 32;
uint16_t * dst_ptr = dst + dst_block + ii * s2;
uint16x8_t d0 = vld1q_u16(src_ptr + 0);
@@ -51,25 +52,26 @@ void pack_B_rk3588_fp16(
// Packing KxN INT8 (row-major) into native RKNN for RK3588: (N/32, K/32, 32, 32)
void pack_B_rk3588_int8(
uint8_t* dst_u8, const uint8_t* src_u8,
int K, int N_total, int n_offset, int n_segment) {
int K_total, int k_offset, int k_segment,
int N_total, int n_offset, int n_segment) {
auto dst = reinterpret_cast<int8_t*>(dst_u8);
auto src = reinterpret_cast<const int8_t*>(src_u8);
GGML_ASSERT(K % 32 == 0 && N_total > 0 && K > 0);
GGML_ASSERT(k_segment % 32 == 0 && N_total > 0 && K_total > 0);
GGML_ASSERT(n_offset % 32 == 0 && n_segment % 32 == 0 && n_offset + n_segment <= N_total);
const size_t s0 = (size_t)(K / 32) * 32 * 32;
const size_t s0 = (size_t)(k_segment / 32) * 32 * 32;
const size_t s1 = 32 * 32;
const size_t s2 = 32;
for (int i = 0; i < n_segment / 32; ++i) {
for (int j = 0; j < K / 32; ++j) {
for (int j = 0; j < k_segment / 32; ++j) {
const size_t dst_block = (size_t) i * s0 + (size_t) j * s1;
for (int ii = 0; ii < 32; ++ii) {
const size_t n_global = (size_t)n_offset + (size_t)i * 32 + (size_t)ii;
const int8_t* src_ptr = src + n_global * K + j * 32;
const int8_t* src_ptr = src + n_global * K_total + k_offset + j * 32;
int8_t* dst_ptr = dst + dst_block + ii * s2;
int8x16_t d0 = vld1q_s8(src_ptr);
@@ -85,24 +87,25 @@ void pack_B_rk3588_int8(
// Packing KxN INT4 (row-major) into native RKNN for RK3588: (N/64, K/32, 64, 32)
void pack_B_rk3588_int4(
uint8_t * dst, const uint8_t * src,
int K, int N_total, int n_offset, int n_segment) {
int K_total, int k_offset, int k_segment,
int N_total, int n_offset, int n_segment) {
GGML_ASSERT(K % 32 == 0 && N_total > 0 && K > 0);
GGML_ASSERT(k_segment % 32 == 0 && N_total > 0 && K_total > 0);
GGML_ASSERT(n_offset % 64 == 0 && n_segment % 64 == 0 && n_offset + n_segment <= N_total);
const size_t s0 = (size_t)(K / 32) * 64 * (32 / 2);
const size_t s0 = (size_t)(k_segment / 32) * 64 * (32 / 2);
const size_t s1 = 64 * (32 / 2);
const size_t s2 = (32 / 2);
const size_t s2 = (32 / 2);
const size_t src_row_stride_bytes = (size_t)K / 2;
const size_t src_row_stride_bytes = (size_t)K_total / 2;
for (int i = 0; i < n_segment / 64; ++i) {
for (int j = 0; j < K / 32; ++j) {
for (int j = 0; j < k_segment / 32; ++j) {
const size_t dst_block = (size_t) i * s0 + (size_t) j * s1;
for (int ii = 0; ii < 64; ++ii) {
const size_t n_global = (size_t)n_offset + (size_t)i * 64 + (size_t)ii;
const uint8_t* src_ptr = src + n_global * src_row_stride_bytes + (j * 32) / 2;
const uint8_t* src_ptr = src + n_global * src_row_stride_bytes + (k_offset + j * 32) / 2;
uint8_t* dst_ptr = dst + dst_block + ii * s2;
uint8x16_t d0 = vld1q_u8(src_ptr);
@@ -139,11 +142,11 @@ const std::vector<std::string>* Rknpu2DeviceConfig::get_active_pattern(int tenso
if (it == default_patterns.end()) {
return nullptr;
}
if (use_custom_pattern && !custom_hybrid_pattern.empty()) {
return &custom_hybrid_pattern;
}
return &it->second;
}
@@ -159,7 +162,7 @@ const Rknpu2HardwarePipeline* Rknpu2DeviceConfig::resolve_op_support(const struc
// 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;
@@ -211,6 +214,7 @@ Rknpu2ConfigManager::Rknpu2ConfigManager() {
Rknpu2DeviceConfig rk3588_config;
rk3588_config.device_name = "RK3588";
rk3588_config.core_count = 3;
rk3588_config.max_k_limit = 8192;
rk3588_config.hardware_pipelines = {
{
/* .pipeline_name = */ "FP16_STANDARD",
@@ -220,6 +224,7 @@ Rknpu2ConfigManager::Rknpu2ConfigManager() {
/* .k_align = */ 32,
/* .n_align = */ 16,
/* .pack_func = */ pack_B_rk3588_fp16,
/* .effective_k = */ 0,
/* .use_hadamard = */ false
},
{
@@ -230,6 +235,7 @@ Rknpu2ConfigManager::Rknpu2ConfigManager() {
/* .k_align = */ 32,
/* .n_align = */ 16,
/* .pack_func = */ pack_B_rk3588_fp16,
/* .effective_k = */ 0,
/* .use_hadamard = */ true
},
{
@@ -240,6 +246,7 @@ Rknpu2ConfigManager::Rknpu2ConfigManager() {
/* .k_align = */ 32,
/* .n_align = */ 32,
/* .pack_func = */ pack_B_rk3588_int8,
/* .effective_k = */ 0,
/* .use_hadamard = */ false
},
{
@@ -250,6 +257,7 @@ Rknpu2ConfigManager::Rknpu2ConfigManager() {
/* .k_align = */ 32,
/* .n_align = */ 32,
/* .pack_func = */ pack_B_rk3588_int8,
/* .effective_k = */ 0,
/* .use_hadamard = */ true
},
{
@@ -260,6 +268,7 @@ Rknpu2ConfigManager::Rknpu2ConfigManager() {
/* .k_align = */ 32,
/* .n_align = */ 64,
/* .pack_func = */ pack_B_rk3588_int4,
/* .effective_k = */ 0,
/* .use_hadamard = */ false
},
{
@@ -270,10 +279,11 @@ Rknpu2ConfigManager::Rknpu2ConfigManager() {
/* .k_align = */ 32,
/* .n_align = */ 64,
/* .pack_func = */ pack_B_rk3588_int4,
/* .effective_k = */ 0,
/* .use_hadamard = */ true
}
};
// Assigning custom variables
rk3588_config.use_custom_pattern = use_custom_pattern;
rk3588_config.custom_hybrid_pattern = custom_pattern;
ggml/src/ggml-rknpu2/rknpu2-configuration.h
+12 -6
View File
@@ -21,7 +21,9 @@ namespace rknpu2_configuration {
*
* @param dst Pointer to the destination buffer for the packed data.
* @param src Pointer to the source data in row-major layout (ggml format).
* @param K The number of columns in the source matrix (inner dimension).
* @param K_total Total original columns in source matrix.
* @param k_offset The column offset from which to start packing.
* @param k_segment The number of columns to pack for this segment.
* @param N_total The total number of rows in the source matrix.
* @param n_offset The row offset from which to start packing.
* @param n_segment The number of rows to pack.
@@ -29,7 +31,9 @@ namespace rknpu2_configuration {
using PackingFunction = std::function<void(
uint8_t* dst,
const uint8_t* src,
int K,
int K_total,
int k_offset,
int k_segment,
int N_total,
int n_offset,
int n_segment
@@ -51,9 +55,9 @@ enum Rknpu2NpuType {
/**
* @brief Describes a single supported hardware pipeline available on the NPU.
*
* 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
* 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 Rknpu2HardwarePipeline {
@@ -67,6 +71,7 @@ struct Rknpu2HardwarePipeline {
int n_align; // Required alignment for the N dimension
PackingFunction pack_func; // Function to pack the weight matrix for this op
int effective_k; // Specific K limit for matrix tiling
bool use_hadamard; // Flag for using Hadamard Transform
};
@@ -79,6 +84,7 @@ struct Rknpu2HardwarePipeline {
struct Rknpu2DeviceConfig {
std::string device_name;
int core_count;
int max_k_limit = 0;
std::vector<Rknpu2HardwarePipeline> hardware_pipelines;
// Type-specific default patterns mapping
@@ -86,7 +92,7 @@ struct Rknpu2DeviceConfig {
// Custom global pattern
bool use_custom_pattern = false;
std::vector<std::string> custom_hybrid_pattern;
std::vector<std::string> custom_hybrid_pattern;
// Tensor-Quantization mapping
mutable std::shared_ptr<std::mutex> pattern_mutex = std::make_shared<std::mutex>();