From 3a1db5260128893662c82fe9ed828697d1200990 Mon Sep 17 00:00:00 2001 From: Michal Harakal Date: Thu, 27 Aug 2026 22:48:24 +0200 Subject: [PATCH 1/3] feat(#1191/#1192/#1193): all GGML block formats mapped-servable on Android and JVM; fallbacks visible MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The 0.50.0 hardening slice. SmolLM2-135M (mixed-quant: k-quant fallback made most of it Q8_0) under MAPPED without prepack went from 48,771 ms/step — silent reference fallback — to 33 ms/step, the fastest configuration measured, with the whole model mapped (98 MB off-heap, 137 KB weight heap). - C: row-major variants for Q8_0/Q4_0/Q5_0/Q5_1/Q5_K, refactored like Q4_K/Q6_K (shared block-term, threaded via skainet_row_threads, #1195). Cross-order results are numerically equivalent, not bit-exact: -ffast-math may contract float accumulation differently per loop shape (q5_1 measured 2 ULP); integer-dot formats happen to match exactly. Docs and the parity suite say so (tolerant cross-order asserts; threaded-vs-solo stays exact). - JNI (#1193): direct-buffer entries for the new formats AND heap-array _rm entries for all seven — JniRowMajorMatmulKernel serves BLOCKED_ROW_MAJOR from mapped, direct, or heap storage with the same C. The 48 s trap was a heap canonical weight with no kernel on its key. - FFM (#1191): FfmRowMajorKernelPack — the same _rm symbols on the JVM, MemorySegment.ofBuffer zero-copy for mapped/direct weights, arena-staged for heap. End-to-end dispatch test proves selection via the trace. - Visibility (#1193): ViewKernel gains a sink-aware run overload; every packed-bridge fallback to the decoding reference emits a KernelRun trace event naming the kernel and reason. The M2-A5 harness reports the count — both device runs now print "reference fallbacks: 0". - Loader/planner: all seven formats in MAPPED_SERVABLE_DEFAULT, the loader's mapped gate, JvmMappedFile, and BufferPackedTensorData; the planner pin for "unmappable stays heap-charged" moves from Q8_0 to TQ2_0 (ternary repack cache is the remaining #1192 follow-up). - Kernel matrix: mapped-serving table now shows all seven formats on both Android (native-jni-direct) and JVM (ffm-rowmajor). Verified on the Pixel 8a (cooled runs): SmolLM2 mapped no-prepack 33 ms/step, 0 faults, 0 fallbacks; Qwen2.5-1.5B mapped 66 ms/step, ✔ fits, 0 fallbacks. JVM: 149 native-cpu tests + FFM dispatch end-to-end green; apiCheck green. Co-authored-by: Claude Fable 5 --- .../reference/kernel-support-matrix.adoc | 9 +- .../backend/api/kernel/KernelDispatch.kt | 2 +- .../ainet/backend/api/kernel/MatmulKernels.kt | 9 + .../api/kernel/PackedViewMatmulKernel.kt | 37 +++- .../native/skainet_jni.c | 176 +++++++++++++++++ .../harness/android/M2A5DeviceMeasurement.kt | 6 + .../exec/kernel/jni/JniBufferPackedKernels.kt | 130 ++++++++----- .../sk/ainet/exec/kernel/jni/JniKernels.kt | 100 ++++++++++ .../native/include/skainet_kernels.h | 94 +++++++++ .../native/src/q4_0_matmul.c | 126 +++++++++--- .../native/src/q5_0_matmul.c | 143 ++++++++++---- .../native/src/q5_1_matmul.c | 143 ++++++++++---- .../native/src/q5k_matmul.c | 125 ++++++++---- .../native/src/q8_0_matmul.c | 184 ++++++++++++------ .../ainet/exec/kernel/FfmRowMajorKernels.kt | 171 ++++++++++++++++ .../exec/kernel/FfmRowMajorDispatchTest.kt | 133 +++++++++++++ .../exec/kernel/KernelSupportMatrixTest.kt | 8 +- .../exec/kernel/RowMajorMatmulParityTest.kt | 66 ++++++- .../kotlin/sk/ainet/io/JvmMappedFile.kt | 7 +- .../io/gguf/StreamingGgufParametersLoader.kt | 7 +- .../ainet/io/gguf/MappedPackedStagingTest.kt | 6 +- .../ainet/io/gguf/StagingPolicyParityTest.kt | 4 +- .../lang/memory/plan/AllocationResolver.kt | 13 +- .../lang/memory/plan/MappedBudgetPlanTest.kt | 21 +- .../tensor/data/BufferPackedTensorData.kt | 44 ++++- 25 files changed, 1492 insertions(+), 272 deletions(-) create mode 100644 skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/FfmRowMajorKernels.kt create mode 100644 skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/FfmRowMajorDispatchTest.kt diff --git a/docs/modules/ROOT/pages/reference/kernel-support-matrix.adoc b/docs/modules/ROOT/pages/reference/kernel-support-matrix.adoc index b6d4f1f1f..2f48adfc2 100644 --- a/docs/modules/ROOT/pages/reference/kernel-support-matrix.adoc +++ b/docs/modules/ROOT/pages/reference/kernel-support-matrix.adoc @@ -28,8 +28,13 @@ Kernels that read the weight in canonical row-major GGUF file order straight fro |=== | Weight format | JVM | Android | Native·linux | Native·apple | JS/WASM -| `Q4_K` | — | native-jni-direct | — | — | — -| `Q6_K` | — | native-jni-direct | — | — | — +| `Q8_0` | ffm-rowmajor | native-jni-direct | — | — | — +| `Q4_0` | ffm-rowmajor | native-jni-direct | — | — | — +| `Q4_K` | ffm-rowmajor | native-jni-direct | — | — | — +| `Q6_K` | ffm-rowmajor | native-jni-direct | — | — | — +| `Q5_K` | ffm-rowmajor | native-jni-direct | — | — | — +| `Q5_1` | ffm-rowmajor | native-jni-direct | — | — | — +| `Q5_0` | ffm-rowmajor | native-jni-direct | — | — | — |=== See also the eager backends & kernels mindmap (xref:explanation/eager-execution.adoc[]) for the narrative overview and gaps. diff --git a/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/KernelDispatch.kt b/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/KernelDispatch.kt index 8d5f8a9ef..1f24ce0d3 100644 --- a/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/KernelDispatch.kt +++ b/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/KernelDispatch.kt @@ -155,7 +155,7 @@ public object KernelDispatch { output = out.id, bytesRead = inputs.sumOf { it.elementCount * it.format.dtype.sizeInBytes }, bytesWritten = out.elementCount * out.format.dtype.sizeInBytes, - ) { kernel.run(inputs, out) } + ) { kernel.run(inputs, out, sink) } } } diff --git a/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/MatmulKernels.kt b/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/MatmulKernels.kt index 5aae52106..47d12ebbc 100644 --- a/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/MatmulKernels.kt +++ b/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/MatmulKernels.kt @@ -19,6 +19,15 @@ public interface ViewKernel { /** Run the kernel: [inputs] as described by [key], result written into [out]. */ public fun run(inputs: List, out: TensorView) + + /** + * Run with a [sink] for events the kernel itself must report — above all an internal + * fallback to the decoding reference (#1193: a silent 1000× slowdown deserves a trace + * line the way an adapter gets one, SKEEP-003 §5.1). Default delegates to [run]; kernels + * with an internal fallback override this and emit before punting. + */ + public fun run(inputs: List, out: TensorView, sink: sk.ainet.lang.memory.trace.TraceSink): Unit = + run(inputs, out) } /** diff --git a/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/PackedViewMatmulKernel.kt b/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/PackedViewMatmulKernel.kt index c4b968072..dcdcb725f 100644 --- a/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/PackedViewMatmulKernel.kt +++ b/skainet-backends/skainet-backend-api/src/commonMain/kotlin/sk/ainet/backend/api/kernel/PackedViewMatmulKernel.kt @@ -35,7 +35,10 @@ public class PackedViewMatmulKernel( override val name: String = "$providerName-$encodingName" - override fun run(inputs: List, out: TensorView) { + override fun run(inputs: List, out: TensorView): Unit = + run(inputs, out, sk.ainet.lang.memory.trace.NoopTraceSink) + + override fun run(inputs: List, out: TensorView, sink: sk.ainet.lang.memory.trace.TraceSink) { require(inputs.size == 2) { "matmul takes two operands" } val a = inputs[0] val w = inputs[1] @@ -48,15 +51,15 @@ public class PackedViewMatmulKernel( "should have prepacked it (#973)" } - val aHeap = a.storage as? Storage.Heap ?: return fallback(inputs, out) - val wHeap = w.storage as? Storage.Heap ?: return fallback(inputs, out) - val oHeap = out.storage as? Storage.Heap ?: return fallback(inputs, out) - val activation = aHeap.floats ?: return fallback(inputs, out) - val weight = wHeap.bytes ?: return fallback(inputs, out) - val output = oHeap.floats ?: return fallback(inputs, out) + val aHeap = a.storage as? Storage.Heap ?: return fallback(inputs, out, sink, "activation storage ${a.storage::class.simpleName}") + val wHeap = w.storage as? Storage.Heap ?: return fallback(inputs, out, sink, "weight storage ${w.storage::class.simpleName} — feed-order kernels take heap bytes") + val oHeap = out.storage as? Storage.Heap ?: return fallback(inputs, out, sink, "output storage ${out.storage::class.simpleName}") + val activation = aHeap.floats ?: return fallback(inputs, out, sink, "activation is not a FloatArray") + val weight = wHeap.bytes ?: return fallback(inputs, out, sink, "weight is not a ByteArray") + val output = oHeap.floats ?: return fallback(inputs, out, sink, "output is not a FloatArray") // The SPI takes a contiguous activation row; a strided one would be mis-indexed, so it goes // to the reference kernel rather than silently reading the wrong floats. - if (!a.isContiguous) return fallback(inputs, out) + if (!a.isContiguous) return fallback(inputs, out, sink, "strided activation") val weightOffset = wHeap.arrayOffset + (w.layout.offsetElements * w.layout.elementBytes).toInt() for (r in 0 until rows) { @@ -69,7 +72,23 @@ public class PackedViewMatmulKernel( } } - private fun fallback(inputs: List, out: TensorView) { + private fun fallback( + inputs: List, + out: TensorView, + sink: sk.ainet.lang.memory.trace.TraceSink, + reason: String, + ) { + // #1193: the decoding reference is ~1000× slower — a trace line, never silent. + if (sink.isEnabled) { + sink.emit( + sk.ainet.lang.memory.trace.TraceEvent.KernelRun( + op = key.op, + kernel = "reference-fallback from $name: $reason", + inputs = inputs.map { it.id }, + output = out.id, + ), + ) + } ReferenceMatmulKernel(key).run(inputs, out) } diff --git a/skainet-backends/skainet-backend-jni-cpu/native/skainet_jni.c b/skainet-backends/skainet-backend-jni-cpu/native/skainet_jni.c index 2b9c58e44..11a8bdc5b 100644 --- a/skainet-backends/skainet-backend-jni-cpu/native/skainet_jni.c +++ b/skainet-backends/skainet-backend-jni-cpu/native/skainet_jni.c @@ -194,6 +194,182 @@ Java_sk_ainet_exec_kernel_jni_JniKernels_q6kMatmulRmDirect( inputDim, outputDim, out, outputOffset)) } +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q80MatmulRmDirect( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jobject weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_RM_DIRECT_BODY( + skainet_q8_0_matmul_rm(in, inputOffset, w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +/* + * Heap-array row-major matmuls (#1193): the same _rm kernels over a pinned + * ByteArray weight in canonical file order. This is what serves a HEAP-staged + * canonical weight without a prepack — before these existed, a heap tensor + * whose format was not mapped-servable fell to the decoding reference kernel + * silently (measured 48,771 ms/step vs 65 on a mixed-quant model). + */ +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q4kMatmulRm( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jbyteArray weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_BODY( + skainet_q4k_matmul_rm(in, inputOffset, (const uint8_t*) w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q6kMatmulRm( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jbyteArray weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_BODY( + skainet_q6k_matmul_rm(in, inputOffset, (const uint8_t*) w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q80MatmulRm( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jbyteArray weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_BODY( + skainet_q8_0_matmul_rm(in, inputOffset, (const uint8_t*) w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q40MatmulRmDirect( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jobject weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_RM_DIRECT_BODY( + skainet_q4_0_matmul_rm(in, inputOffset, w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q40MatmulRm( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jbyteArray weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_BODY( + skainet_q4_0_matmul_rm(in, inputOffset, (const uint8_t*) w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q50MatmulRmDirect( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jobject weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_RM_DIRECT_BODY( + skainet_q5_0_matmul_rm(in, inputOffset, w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q50MatmulRm( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jbyteArray weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_BODY( + skainet_q5_0_matmul_rm(in, inputOffset, (const uint8_t*) w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q51MatmulRmDirect( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jobject weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_RM_DIRECT_BODY( + skainet_q5_1_matmul_rm(in, inputOffset, w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q51MatmulRm( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jbyteArray weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_BODY( + skainet_q5_1_matmul_rm(in, inputOffset, (const uint8_t*) w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q5kMatmulRmDirect( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jobject weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_RM_DIRECT_BODY( + skainet_q5k_matmul_rm(in, inputOffset, w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + +JNIEXPORT void JNICALL +Java_sk_ainet_exec_kernel_jni_JniKernels_q5kMatmulRm( + JNIEnv* env, jobject thiz, + jfloatArray input, jint inputOffset, + jbyteArray weight, jint weightByteOffset, + jint inputDim, jint outputDim, + jfloatArray output, jint outputOffset +) { + (void) thiz; + SKAINET_JNI_MATMUL_BODY( + skainet_q5k_matmul_rm(in, inputOffset, (const uint8_t*) w, weightByteOffset, + inputDim, outputDim, out, outputOffset)) +} + /* * bitnet_gemv (SKEEP-003 §5.3, #1041): int8 activations against ternary TQ2_0 * weights. Its activation is a *byte* array, not floats, so it does not fit diff --git a/skainet-backends/skainet-backend-jni-cpu/src/androidTest/kotlin/sk/ainet/exec/harness/android/M2A5DeviceMeasurement.kt b/skainet-backends/skainet-backend-jni-cpu/src/androidTest/kotlin/sk/ainet/exec/harness/android/M2A5DeviceMeasurement.kt index 44145bf76..766e1d2be 100644 --- a/skainet-backends/skainet-backend-jni-cpu/src/androidTest/kotlin/sk/ainet/exec/harness/android/M2A5DeviceMeasurement.kt +++ b/skainet-backends/skainet-backend-jni-cpu/src/androidTest/kotlin/sk/ainet/exec/harness/android/M2A5DeviceMeasurement.kt @@ -258,6 +258,12 @@ class M2A5DeviceMeasurement { line() line("- steady-state (after $warmup warm-up steps): ${steady.map { it.ms }.average().toInt()} ms/step, " + "major faults total ${steady.mapNotNull { it.majDelta }.sum()}") + // #1193: reference fallbacks are trace events now — a nonzero count here is the + // 48 s/step trap announcing itself instead of hiding in the timings. + val fallbacks = sink.events().filterIsInstance() + .filter { "reference-fallback" in it.kernel } + line("- packed-kernel reference fallbacks: ${fallbacks.size}" + + if (fallbacks.isEmpty()) "" else " — e.g. ${fallbacks.first().kernel}") line("- forward-scope allocations after warm-up should be 0; " + "whole-run forward allocations: ${actual.allocationsByScope[ScopeKind.FORWARD] ?: 0}") } diff --git a/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniBufferPackedKernels.kt b/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniBufferPackedKernels.kt index ecefa3e6d..5cf9bce57 100644 --- a/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniBufferPackedKernels.kt +++ b/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniBufferPackedKernels.kt @@ -14,35 +14,50 @@ import sk.ainet.lang.memory.Format import sk.ainet.lang.memory.MappedBufferStorage import sk.ainet.lang.memory.Storage import sk.ainet.lang.memory.TensorView +import sk.ainet.lang.memory.trace.NoopTraceSink +import sk.ainet.lang.memory.trace.TraceEvent +import sk.ainet.lang.memory.trace.TraceSink import sk.ainet.lang.tensor.storage.TensorEncoding import sk.ainet.lang.types.FP32 /** - * A [ViewKernel] over the JNI **direct-buffer row-major** matmuls (#1189): the weight is a - * `BLOCKED_ROW_MAJOR` packed view whose storage is off-heap ([MappedBufferStorage] — mmap'd GGUF - * pages — or [DirectBufferStorage]), read by the native kernel straight at its address. No heap - * `ByteArray`, no relayout: this is what lets a model bigger than the ART cap decode on Android. + * A [ViewKernel] over the JNI **row-major** matmuls, polymorphic in the weight's storage + * (#1189/#1193): a `BLOCKED_ROW_MAJOR` packed weight is served wherever its bytes live — + * off-heap ([MappedBufferStorage] mmap'd GGUF pages, [DirectBufferStorage]) through the + * direct-buffer entries, or a heap `ByteArray` through the pinned-array entries. Same C kernels + * either way; the pointer's provenance is a bridge detail, not a kernel property. + * + * The heap arm is what makes canonical (un-prepacked) heap weights fast: before it existed, a + * heap tensor in file order fell to the decoding reference silently — measured 48,771 ms/step + * vs 65 on a mixed-quant model (#1193). The remaining fallbacks (non-contiguous activation, + * exotic storage) now announce themselves to the trace sink instead of hiding. * * Contrast with [sk.ainet.backend.api.kernel.PackedViewMatmulKernel], which serves the same - * encodings from heap bytes in `BLOCKED_INPUT_MAJOR` (prepacked feed) order. Registering both - * lets the dispatcher pick by what the weight actually is — canonical mapped weights match this - * key exactly and are served with zero copies. + * encodings from heap bytes in `BLOCKED_INPUT_MAJOR` (prepacked feed) order. */ @ExperimentalMemoryApi -public class JniBufferPackedMatmulKernel( +public class JniRowMajorMatmulKernel( encodingName: String, override val key: KernelKey, - private val matmul: ( + private val bufferMatmul: ( input: FloatArray, inputOffset: Int, weight: ByteBuffer, weightByteOffset: Int, inputDim: Int, outputDim: Int, output: FloatArray, outputOffset: Int, ) -> Unit, + private val arrayMatmul: ( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) -> Unit, ) : ViewKernel { - override val name: String = "jni-buffer-$encodingName" + override val name: String = "jni-rowmajor-$encodingName" + + override fun run(inputs: List, out: TensorView): Unit = run(inputs, out, NoopTraceSink) - override fun run(inputs: List, out: TensorView) { + override fun run(inputs: List, out: TensorView, sink: TraceSink) { require(inputs.size == 2) { "matmul takes two operands" } val a = inputs[0] val w = inputs[1] @@ -54,30 +69,54 @@ public class JniBufferPackedMatmulKernel( "$name reads canonical row-major weights; this one is ${w.layout.blockOrder}" } - val buffer = when (val s = w.storage) { - is MappedBufferStorage -> s.buffer() - is DirectBufferStorage -> s.buffer() - else -> return fallback(inputs, out) - } - val aHeap = a.storage as? Storage.Heap ?: return fallback(inputs, out) - val oHeap = out.storage as? Storage.Heap ?: return fallback(inputs, out) - val activation = aHeap.floats ?: return fallback(inputs, out) - val output = oHeap.floats ?: return fallback(inputs, out) + val aHeap = a.storage as? Storage.Heap ?: return fallback(inputs, out, sink, "activation storage ${a.storage::class.simpleName}") + val oHeap = out.storage as? Storage.Heap ?: return fallback(inputs, out, sink, "output storage ${out.storage::class.simpleName}") + val activation = aHeap.floats ?: return fallback(inputs, out, sink, "activation is not a FloatArray") + val output = oHeap.floats ?: return fallback(inputs, out, sink, "output is not a FloatArray") // The JNI kernel takes a contiguous activation row; a strided one would be mis-indexed. - if (!a.isContiguous) return fallback(inputs, out) + if (!a.isContiguous) return fallback(inputs, out, sink, "strided activation") val weightOffset = (w.layout.offsetElements * w.layout.elementBytes).toInt() - for (r in 0 until rows) { - matmul( - activation, aHeap.arrayOffset + (a.layout.offsetElements + r.toLong() * k).toInt(), - buffer, weightOffset, - k, n, - output, oHeap.arrayOffset + (out.layout.offsetElements + r.toLong() * n).toInt(), - ) + when (val s = w.storage) { + is MappedBufferStorage, is DirectBufferStorage -> { + val buffer = if (s is MappedBufferStorage) s.buffer() else (s as DirectBufferStorage).buffer() + for (r in 0 until rows) { + bufferMatmul( + activation, aHeap.arrayOffset + (a.layout.offsetElements + r.toLong() * k).toInt(), + buffer, weightOffset, + k, n, + output, oHeap.arrayOffset + (out.layout.offsetElements + r.toLong() * n).toInt(), + ) + } + } + is Storage.Heap -> { + val bytes = s.bytes ?: return fallback(inputs, out, sink, "heap weight is not a ByteArray") + for (r in 0 until rows) { + arrayMatmul( + activation, aHeap.arrayOffset + (a.layout.offsetElements + r.toLong() * k).toInt(), + bytes, s.arrayOffset + weightOffset, + k, n, + output, oHeap.arrayOffset + (out.layout.offsetElements + r.toLong() * n).toInt(), + ) + } + } + else -> return fallback(inputs, out, sink, "weight storage ${s::class.simpleName}") } } - private fun fallback(inputs: List, out: TensorView) { + private fun fallback(inputs: List, out: TensorView, sink: TraceSink, reason: String) { + // #1193: never silent — the decoding reference is ~1000× the packed kernel, and that has + // to be a line in the trace, exactly the way an inserted adapter is (SKEEP-003 §5.1). + if (sink.isEnabled) { + sink.emit( + TraceEvent.KernelRun( + op = key.op, + kernel = "reference-fallback from $name: $reason", + inputs = inputs.map { it.id }, + output = out.id, + ), + ) + } ReferenceMatmulKernel(key).run(inputs, out) } @@ -94,28 +133,27 @@ public class JniBufferPackedMatmulKernel( } /** - * Registers the direct-buffer row-major kernels (#1189) into [KernelDispatch]. Call next to - * `KernelPacks.install(JniKernelProvider)` on Android when weights are loaded with - * `WeightResidency.MAPPED` — without it a mapped packed weight falls back to the decoding - * reference kernel (correct, hours-scale slow). + * Registers the row-major kernels (#1189/#1192/#1193) into [KernelDispatch]: all seven GGML block formats, each serving mapped, direct and heap weights in canonical file order. Call next to + * `KernelPacks.install(JniKernelProvider)` on Android — without it a canonical packed weight + * (mapped OR un-prepacked heap) falls back to the decoding reference. */ @ExperimentalMemoryApi public object JniMappedKernelPack { public fun install() { if (!JniKernelProvider.isAvailable()) return - KernelDispatch.register( - JniBufferPackedMatmulKernel( - TensorEncoding.Q4_K.name, - JniBufferPackedMatmulKernel.keyFor(TensorEncoding.Q4_K), - JniKernels::q4kMatmulRmDirect, - ), - ) - KernelDispatch.register( - JniBufferPackedMatmulKernel( - TensorEncoding.Q6_K.name, - JniBufferPackedMatmulKernel.keyFor(TensorEncoding.Q6_K), - JniKernels::q6kMatmulRmDirect, - ), + fun register( + encoding: TensorEncoding, + buffer: (FloatArray, Int, ByteBuffer, Int, Int, Int, FloatArray, Int) -> Unit, + array: (FloatArray, Int, ByteArray, Int, Int, Int, FloatArray, Int) -> Unit, + ) = KernelDispatch.register( + JniRowMajorMatmulKernel(encoding.name, JniRowMajorMatmulKernel.keyFor(encoding), buffer, array), ) + register(TensorEncoding.Q4_K, JniKernels::q4kMatmulRmDirect, JniKernels::q4kMatmulRm) + register(TensorEncoding.Q6_K, JniKernels::q6kMatmulRmDirect, JniKernels::q6kMatmulRm) + register(TensorEncoding.Q5_K, JniKernels::q5kMatmulRmDirect, JniKernels::q5kMatmulRm) + register(TensorEncoding.Q8_0, JniKernels::q80MatmulRmDirect, JniKernels::q80MatmulRm) + register(TensorEncoding.Q4_0, JniKernels::q40MatmulRmDirect, JniKernels::q40MatmulRm) + register(TensorEncoding.Q5_0, JniKernels::q50MatmulRmDirect, JniKernels::q50MatmulRm) + register(TensorEncoding.Q5_1, JniKernels::q51MatmulRmDirect, JniKernels::q51MatmulRm) } } diff --git a/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniKernels.kt b/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniKernels.kt index 9101fc36c..60d4019b5 100644 --- a/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniKernels.kt +++ b/skainet-backends/skainet-backend-jni-cpu/src/main/kotlin/sk/ainet/exec/kernel/jni/JniKernels.kt @@ -192,4 +192,104 @@ public object JniKernels { inputDim: Int, outputDim: Int, output: FloatArray, outputOffset: Int, ) + + /** Row-major Q8_0 matmul over a direct ByteBuffer weight (#1192); see [q4kMatmulRmDirect]. */ + public external fun q80MatmulRmDirect( + input: FloatArray, inputOffset: Int, + weight: java.nio.ByteBuffer, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** + * Row-major matmuls over a **heap** `ByteArray` weight in canonical GGUF file order (#1193): + * the same `_rm` kernels, pinned-array entry. These serve heap-staged canonical weights + * without a prepack — the case that used to fall silently to the decoding reference. + */ + public external fun q4kMatmulRm( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** See [q4kMatmulRm]. */ + public external fun q6kMatmulRm( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** See [q4kMatmulRm]. */ + public external fun q80MatmulRm( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** Row-major q40 matmul over a direct ByteBuffer weight (#1192); see [q4kMatmulRmDirect]. */ + public external fun q40MatmulRmDirect( + input: FloatArray, inputOffset: Int, + weight: java.nio.ByteBuffer, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** See [q4kMatmulRm]. */ + public external fun q40MatmulRm( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** Row-major q50 matmul over a direct ByteBuffer weight (#1192); see [q4kMatmulRmDirect]. */ + public external fun q50MatmulRmDirect( + input: FloatArray, inputOffset: Int, + weight: java.nio.ByteBuffer, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** See [q4kMatmulRm]. */ + public external fun q50MatmulRm( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** Row-major q51 matmul over a direct ByteBuffer weight (#1192); see [q4kMatmulRmDirect]. */ + public external fun q51MatmulRmDirect( + input: FloatArray, inputOffset: Int, + weight: java.nio.ByteBuffer, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** See [q4kMatmulRm]. */ + public external fun q51MatmulRm( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** Row-major q5k matmul over a direct ByteBuffer weight (#1192); see [q4kMatmulRmDirect]. */ + public external fun q5kMatmulRmDirect( + input: FloatArray, inputOffset: Int, + weight: java.nio.ByteBuffer, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) + + /** See [q4kMatmulRm]. */ + public external fun q5kMatmulRm( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) } diff --git a/skainet-backends/skainet-backend-native-cpu/native/include/skainet_kernels.h b/skainet-backends/skainet-backend-native-cpu/native/include/skainet_kernels.h index 02ac35694..dbe7e0adb 100644 --- a/skainet-backends/skainet-backend-native-cpu/native/include/skainet_kernels.h +++ b/skainet-backends/skainet-backend-native-cpu/native/include/skainet_kernels.h @@ -165,6 +165,100 @@ SKAINET_API void skainet_q6k_matmul_rm( int32_t output_offset ); +/* + * Q8_0 matrix-vector multiply over a ROW-MAJOR (canonical GGUF file order) + * weight (#1192). Same math and block format as skainet_q8_0_matmul; the + * weight is addressed + * weight + weight_byte_offset + (o * blocks_per_row + block_idx) * 34 + * — the bytes exactly as they sit in a .gguf file, so an mmap'd weight + * needs no relayout copy. input_dim must be a multiple of 32. + * + * Threads over output rows when output_dim >= 512 (#1195) — see + * skainet_q4k_matmul; bit-identical to the single-threaded result. + */ +SKAINET_API void skainet_q8_0_matmul_rm( + const float* input, + int32_t input_offset, + const uint8_t* weight, + int32_t weight_byte_offset, + int32_t input_dim, + int32_t output_dim, + float* output, + int32_t output_offset +); + +/* + * Q40 matrix-vector multiply over a ROW-MAJOR (canonical GGUF file + * order) weight (#1192): weight + weight_byte_offset + (o * blocks_per_row + + * block_idx) * 18 — mmap-fed as-is, no relayout copy. Same math as the + * feed-order kernel, bit-identical; threads over output rows >= 512 (#1195). + * input_dim must be a multiple of 32. + */ +SKAINET_API void skainet_q4_0_matmul_rm( + const float* input, + int32_t input_offset, + const uint8_t* weight, + int32_t weight_byte_offset, + int32_t input_dim, + int32_t output_dim, + float* output, + int32_t output_offset +); + +/* + * Q50 matrix-vector multiply over a ROW-MAJOR (canonical GGUF file + * order) weight (#1192): weight + weight_byte_offset + (o * blocks_per_row + + * block_idx) * 22 — mmap-fed as-is, no relayout copy. Same math as the + * feed-order kernel, bit-identical; threads over output rows >= 512 (#1195). + * input_dim must be a multiple of 32. + */ +SKAINET_API void skainet_q5_0_matmul_rm( + const float* input, + int32_t input_offset, + const uint8_t* weight, + int32_t weight_byte_offset, + int32_t input_dim, + int32_t output_dim, + float* output, + int32_t output_offset +); + +/* + * Q51 matrix-vector multiply over a ROW-MAJOR (canonical GGUF file + * order) weight (#1192): weight + weight_byte_offset + (o * blocks_per_row + + * block_idx) * 24 — mmap-fed as-is, no relayout copy. Same math as the + * feed-order kernel, bit-identical; threads over output rows >= 512 (#1195). + * input_dim must be a multiple of 32. + */ +SKAINET_API void skainet_q5_1_matmul_rm( + const float* input, + int32_t input_offset, + const uint8_t* weight, + int32_t weight_byte_offset, + int32_t input_dim, + int32_t output_dim, + float* output, + int32_t output_offset +); + +/* + * Q5K matrix-vector multiply over a ROW-MAJOR (canonical GGUF file + * order) weight (#1192): weight + weight_byte_offset + (o * blocks_per_row + + * block_idx) * 176 — mmap-fed as-is, no relayout copy. Same math as the + * feed-order kernel, bit-identical; threads over output rows >= 512 (#1195). + * input_dim must be a multiple of 256. + */ +SKAINET_API void skainet_q5k_matmul_rm( + const float* input, + int32_t input_offset, + const uint8_t* weight, + int32_t weight_byte_offset, + int32_t input_dim, + int32_t output_dim, + float* output, + int32_t output_offset +); + /* * Row-major FP32 SGEMM: C(m, n) = A(m, k) * B(k, n). * diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/q4_0_matmul.c b/skainet-backends/skainet-backend-native-cpu/native/src/q4_0_matmul.c index f9830a96e..325c6a759 100644 --- a/skainet-backends/skainet-backend-native-cpu/native/src/q4_0_matmul.c +++ b/skainet-backends/skainet-backend-native-cpu/native/src/q4_0_matmul.c @@ -1,5 +1,6 @@ #include "skainet_kernels.h" #include "skainet_simd.h" +#include "skainet_row_threads.h" #include #include @@ -66,35 +67,18 @@ static inline float skainet_q4_0_fp16_to_fp32(uint16_t h) { return r; } -SKAINET_API void skainet_q4_0_matmul( - const float* SKAINET_RESTRICT input, int32_t input_offset, - const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, - int32_t input_dim, int32_t output_dim, - float* SKAINET_RESTRICT output, int32_t output_offset -) { - if (output_dim <= 0) return; - if (input_dim <= 0) { - for (int32_t o = 0; o < output_dim; ++o) { - output[output_offset + o] = 0.0f; - } - return; - } - - const int32_t BLOCK_SIZE = 32; - const int32_t BYTES_PER_BLOCK = 18; - const int32_t blocks_per_input_dim = input_dim / BLOCK_SIZE; - float* SKAINET_RESTRICT out_base = output + output_offset; - - for (int32_t o = 0; o < output_dim; ++o) out_base[o] = 0.0f; +#define Q40_BLOCK_SIZE 32 +#define Q40_BYTES_PER_BLOCK 18 - for (int32_t block_idx = 0; block_idx < blocks_per_input_dim; ++block_idx) { - const float* SKAINET_RESTRICT input_block = - input + input_offset + (size_t) block_idx * BLOCK_SIZE; - const uint8_t* SKAINET_RESTRICT block = - weight + weight_byte_offset + - (size_t)(block_idx * output_dim) * BYTES_PER_BLOCK; - - for (int32_t o = 0; o < output_dim; ++o, block += BYTES_PER_BLOCK) { +/* + * One block's contribution to out[o] — the loop body of the original kernel, + * shared by the feed-order and row-major entries so both orders (and any row + * partition, #1195) stay bit-identical per output row (#1192). + */ +static inline float q4_0_block_term( + const uint8_t* SKAINET_RESTRICT block, + const float* SKAINET_RESTRICT input_block +) { uint16_t d_bits = (uint16_t) block[0] | ((uint16_t) block[1] << 8); float d = skainet_q4_0_fp16_to_fp32(d_bits); const uint8_t* SKAINET_RESTRICT codes = block + 2; @@ -141,7 +125,91 @@ SKAINET_API void skainet_q4_0_matmul( block_sum += input_block[k + 16] * (float) hi; } #endif - out_base[o] += block_sum * d; + return block_sum * d; +} + +/* Everything a row-range worker needs; read-only during the parallel section. */ +typedef struct { + const uint8_t* weight_base; /* weight + weight_byte_offset */ + const float* in_base; /* input + input_offset */ + float* out_base; + int32_t blocks_per_input_dim; + int32_t output_dim; +} q4_0_ctx; + +/* Feed-order rows [o_start, o_end): block OUTER, row INNER — see q4k_matmul.c. */ +static void q4_0_rows_feed(void* vctx, int32_t o_start, int32_t o_end) { + const q4_0_ctx* c = (const q4_0_ctx*) vctx; + for (int32_t o = o_start; o < o_end; ++o) c->out_base[o] = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; ++block_idx) { + const float* input_block = c->in_base + (size_t) block_idx * Q40_BLOCK_SIZE; + const uint8_t* block = c->weight_base + + ((size_t) block_idx * c->output_dim + o_start) * Q40_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o, block += Q40_BYTES_PER_BLOCK) { + c->out_base[o] += q4_0_block_term(block, input_block); } } } + +/* Row-major rows (#1192): canonical GGUF file order, (o·bpr + b)·18 — mmap-fed as-is. */ +static void q4_0_rows_rm(void* vctx, int32_t o_start, int32_t o_end) { + const q4_0_ctx* c = (const q4_0_ctx*) vctx; + const uint8_t* block = c->weight_base + + (size_t) o_start * c->blocks_per_input_dim * Q40_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o) { + float acc_row = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; + ++block_idx, block += Q40_BYTES_PER_BLOCK) { + acc_row += q4_0_block_term(block, c->in_base + (size_t) block_idx * Q40_BLOCK_SIZE); + } + c->out_base[o] = acc_row; + } +} + +/* Shared entry — guards, ctx fill, threaded run (skainet_row_threads.h, #1195). */ +static void q4_0_matmul_run( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset, + skainet_row_range_fn worker +) { + if (output_dim <= 0) return; + if (input_dim <= 0) { + for (int32_t o = 0; o < output_dim; ++o) output[output_offset + o] = 0.0f; + return; + } + q4_0_ctx ctx; + ctx.weight_base = weight + weight_byte_offset; + ctx.in_base = input + input_offset; + ctx.out_base = output + output_offset; + ctx.blocks_per_input_dim = input_dim / Q40_BLOCK_SIZE; + ctx.output_dim = output_dim; + skainet_run_rows(worker, &ctx, output_dim); +} + +SKAINET_API void skainet_q4_0_matmul( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q4_0_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q4_0_rows_feed); +} + +/* + * Row-major variant (#1192): canonical GGUF file order — see q4_0_rows_rm. + * Numerically equivalent to the feed-order kernel (same per-row block order; -ffast-math FMA/reassociation may differ at ULP scale); threads over rows >= 512 (#1195). + */ +SKAINET_API void skainet_q4_0_matmul_rm( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q4_0_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q4_0_rows_rm); +} diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/q5_0_matmul.c b/skainet-backends/skainet-backend-native-cpu/native/src/q5_0_matmul.c index 8851e904f..856369ffd 100644 --- a/skainet-backends/skainet-backend-native-cpu/native/src/q5_0_matmul.c +++ b/skainet-backends/skainet-backend-native-cpu/native/src/q5_0_matmul.c @@ -1,5 +1,7 @@ +#include #include "skainet_kernels.h" #include "skainet_simd.h" +#include "skainet_row_threads.h" #include #include @@ -73,39 +75,18 @@ static inline float skainet_q5_0_fp16_to_fp32(uint16_t h) { return r; } -SKAINET_API void skainet_q5_0_matmul( - const float* SKAINET_RESTRICT input, int32_t input_offset, - const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, - int32_t input_dim, int32_t output_dim, - float* SKAINET_RESTRICT output, int32_t output_offset -) { - if (output_dim <= 0) return; - if (input_dim <= 0) { - for (int32_t o = 0; o < output_dim; ++o) { - output[output_offset + o] = 0.0f; - } - return; - } - - const int32_t BLOCK_SIZE = 32; - const int32_t BYTES_PER_BLOCK = 22; - const int32_t blocks_per_input_dim = input_dim / BLOCK_SIZE; - float* SKAINET_RESTRICT out_base = output + output_offset; - - for (int32_t o = 0; o < output_dim; ++o) out_base[o] = 0.0f; - - for (int32_t block_idx = 0; block_idx < blocks_per_input_dim; ++block_idx) { - const float* SKAINET_RESTRICT input_block = - input + input_offset + (size_t) block_idx * BLOCK_SIZE; - const uint8_t* SKAINET_RESTRICT block = - weight + weight_byte_offset + - (size_t)(block_idx * output_dim) * BYTES_PER_BLOCK; +#define Q50_BLOCK_SIZE 32 +#define Q50_BYTES_PER_BLOCK 22 - /* Depends only on the activations — hoisted out of the o-loop. */ - float input_sum = 0.0f; - for (int32_t k = 0; k < BLOCK_SIZE; ++k) input_sum += input_block[k]; - - for (int32_t o = 0; o < output_dim; ++o, block += BYTES_PER_BLOCK) { +/* + * One block's contribution to out[o] — the loop body of the original kernel, + * shared by the feed-order and row-major entries so both orders (and any row + * partition, #1195) stay bit-identical per output row (#1192). + */ +static inline float q5_0_block_term( + const uint8_t* SKAINET_RESTRICT block, + const float* SKAINET_RESTRICT input_block, float input_sum +) { uint16_t d_bits = (uint16_t) block[0] | ((uint16_t) block[1] << 8); float d = skainet_q5_0_fp16_to_fp32(d_bits); const uint8_t* SKAINET_RESTRICT qh = block + 2; @@ -158,7 +139,103 @@ SKAINET_API void skainet_q5_0_matmul( code_dot += input_block[k + 16] * (float) hi; } #endif - out_base[o] += d * (code_dot - 16.0f * input_sum); + return d * (code_dot - 16.0f * input_sum); +} + +/* Everything a row-range worker needs; read-only during the parallel section. */ +typedef struct { + const uint8_t* weight_base; /* weight + weight_byte_offset */ + const float* in_base; /* input + input_offset */ + float* out_base; + const float* in_sums; /* per-block activation sums (hoisted; see term) */ + int32_t blocks_per_input_dim; + int32_t output_dim; +} q5_0_ctx; + +/* Feed-order rows [o_start, o_end): block OUTER, row INNER — see q4k_matmul.c. */ +static void q5_0_rows_feed(void* vctx, int32_t o_start, int32_t o_end) { + const q5_0_ctx* c = (const q5_0_ctx*) vctx; + for (int32_t o = o_start; o < o_end; ++o) c->out_base[o] = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; ++block_idx) { + const float* input_block = c->in_base + (size_t) block_idx * Q50_BLOCK_SIZE; + const uint8_t* block = c->weight_base + + ((size_t) block_idx * c->output_dim + o_start) * Q50_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o, block += Q50_BYTES_PER_BLOCK) { + c->out_base[o] += q5_0_block_term(block, input_block, c->in_sums[block_idx]); } } } + +/* Row-major rows (#1192): canonical GGUF file order, (o·bpr + b)·22 — mmap-fed as-is. */ +static void q5_0_rows_rm(void* vctx, int32_t o_start, int32_t o_end) { + const q5_0_ctx* c = (const q5_0_ctx*) vctx; + const uint8_t* block = c->weight_base + + (size_t) o_start * c->blocks_per_input_dim * Q50_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o) { + float acc_row = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; + ++block_idx, block += Q50_BYTES_PER_BLOCK) { + acc_row += q5_0_block_term(block, c->in_base + (size_t) block_idx * Q50_BLOCK_SIZE, c->in_sums[block_idx]); + } + c->out_base[o] = acc_row; + } +} + +/* Shared entry — guards, ctx fill, threaded run (skainet_row_threads.h, #1195). */ +static void q5_0_matmul_run( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset, + skainet_row_range_fn worker +) { + if (output_dim <= 0) return; + if (input_dim <= 0) { + for (int32_t o = 0; o < output_dim; ++o) output[output_offset + o] = 0.0f; + return; + } + q5_0_ctx ctx; + ctx.weight_base = weight + weight_byte_offset; + ctx.in_base = input + input_offset; + ctx.out_base = output + output_offset; + ctx.blocks_per_input_dim = input_dim / Q50_BLOCK_SIZE; + ctx.output_dim = output_dim; + + float* in_sums = (float*) malloc((size_t) ctx.blocks_per_input_dim * sizeof(float)); + if (in_sums == NULL) return; + for (int32_t b = 0; b < ctx.blocks_per_input_dim; ++b) { + const float* ib = ctx.in_base + (size_t) b * Q50_BLOCK_SIZE; + float sum = 0.0f; + for (int32_t k = 0; k < Q50_BLOCK_SIZE; ++k) sum += ib[k]; + in_sums[b] = sum; + } + ctx.in_sums = in_sums; + skainet_run_rows(worker, &ctx, output_dim); + free(in_sums); +} + +SKAINET_API void skainet_q5_0_matmul( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q5_0_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q5_0_rows_feed); +} + +/* + * Row-major variant (#1192): canonical GGUF file order — see q5_0_rows_rm. + * Numerically equivalent to the feed-order kernel (same per-row block order; -ffast-math FMA/reassociation may differ at ULP scale); threads over rows >= 512 (#1195). + */ +SKAINET_API void skainet_q5_0_matmul_rm( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q5_0_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q5_0_rows_rm); +} diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/q5_1_matmul.c b/skainet-backends/skainet-backend-native-cpu/native/src/q5_1_matmul.c index d47842c2d..6e6debaec 100644 --- a/skainet-backends/skainet-backend-native-cpu/native/src/q5_1_matmul.c +++ b/skainet-backends/skainet-backend-native-cpu/native/src/q5_1_matmul.c @@ -1,5 +1,7 @@ +#include #include "skainet_kernels.h" #include "skainet_simd.h" +#include "skainet_row_threads.h" #include #include @@ -71,39 +73,18 @@ static inline float skainet_q5_1_fp16_to_fp32(uint16_t h) { return r; } -SKAINET_API void skainet_q5_1_matmul( - const float* SKAINET_RESTRICT input, int32_t input_offset, - const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, - int32_t input_dim, int32_t output_dim, - float* SKAINET_RESTRICT output, int32_t output_offset -) { - if (output_dim <= 0) return; - if (input_dim <= 0) { - for (int32_t o = 0; o < output_dim; ++o) { - output[output_offset + o] = 0.0f; - } - return; - } - - const int32_t BLOCK_SIZE = 32; - const int32_t BYTES_PER_BLOCK = 24; - const int32_t blocks_per_input_dim = input_dim / BLOCK_SIZE; - float* SKAINET_RESTRICT out_base = output + output_offset; - - for (int32_t o = 0; o < output_dim; ++o) out_base[o] = 0.0f; - - for (int32_t block_idx = 0; block_idx < blocks_per_input_dim; ++block_idx) { - const float* SKAINET_RESTRICT input_block = - input + input_offset + (size_t) block_idx * BLOCK_SIZE; - const uint8_t* SKAINET_RESTRICT block = - weight + weight_byte_offset + - (size_t)(block_idx * output_dim) * BYTES_PER_BLOCK; +#define Q51_BLOCK_SIZE 32 +#define Q51_BYTES_PER_BLOCK 24 - /* Depends only on the activations — hoisted out of the o-loop. */ - float input_sum = 0.0f; - for (int32_t k = 0; k < BLOCK_SIZE; ++k) input_sum += input_block[k]; - - for (int32_t o = 0; o < output_dim; ++o, block += BYTES_PER_BLOCK) { +/* + * One block's contribution to out[o] — the loop body of the original kernel, + * shared by the feed-order and row-major entries so both orders (and any row + * partition, #1195) stay bit-identical per output row (#1192). + */ +static inline float q5_1_block_term( + const uint8_t* SKAINET_RESTRICT block, + const float* SKAINET_RESTRICT input_block, float input_sum +) { uint16_t d_bits = (uint16_t) block[0] | ((uint16_t) block[1] << 8); uint16_t m_bits = (uint16_t) block[2] | ((uint16_t) block[3] << 8); float d = skainet_q5_1_fp16_to_fp32(d_bits); @@ -158,7 +139,103 @@ SKAINET_API void skainet_q5_1_matmul( code_dot += input_block[k + 16] * (float) hi; } #endif - out_base[o] += d * code_dot + m * input_sum; + return d * code_dot + m * input_sum; +} + +/* Everything a row-range worker needs; read-only during the parallel section. */ +typedef struct { + const uint8_t* weight_base; /* weight + weight_byte_offset */ + const float* in_base; /* input + input_offset */ + float* out_base; + const float* in_sums; /* per-block activation sums (hoisted; see term) */ + int32_t blocks_per_input_dim; + int32_t output_dim; +} q5_1_ctx; + +/* Feed-order rows [o_start, o_end): block OUTER, row INNER — see q4k_matmul.c. */ +static void q5_1_rows_feed(void* vctx, int32_t o_start, int32_t o_end) { + const q5_1_ctx* c = (const q5_1_ctx*) vctx; + for (int32_t o = o_start; o < o_end; ++o) c->out_base[o] = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; ++block_idx) { + const float* input_block = c->in_base + (size_t) block_idx * Q51_BLOCK_SIZE; + const uint8_t* block = c->weight_base + + ((size_t) block_idx * c->output_dim + o_start) * Q51_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o, block += Q51_BYTES_PER_BLOCK) { + c->out_base[o] += q5_1_block_term(block, input_block, c->in_sums[block_idx]); } } } + +/* Row-major rows (#1192): canonical GGUF file order, (o·bpr + b)·24 — mmap-fed as-is. */ +static void q5_1_rows_rm(void* vctx, int32_t o_start, int32_t o_end) { + const q5_1_ctx* c = (const q5_1_ctx*) vctx; + const uint8_t* block = c->weight_base + + (size_t) o_start * c->blocks_per_input_dim * Q51_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o) { + float acc_row = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; + ++block_idx, block += Q51_BYTES_PER_BLOCK) { + acc_row += q5_1_block_term(block, c->in_base + (size_t) block_idx * Q51_BLOCK_SIZE, c->in_sums[block_idx]); + } + c->out_base[o] = acc_row; + } +} + +/* Shared entry — guards, ctx fill, threaded run (skainet_row_threads.h, #1195). */ +static void q5_1_matmul_run( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset, + skainet_row_range_fn worker +) { + if (output_dim <= 0) return; + if (input_dim <= 0) { + for (int32_t o = 0; o < output_dim; ++o) output[output_offset + o] = 0.0f; + return; + } + q5_1_ctx ctx; + ctx.weight_base = weight + weight_byte_offset; + ctx.in_base = input + input_offset; + ctx.out_base = output + output_offset; + ctx.blocks_per_input_dim = input_dim / Q51_BLOCK_SIZE; + ctx.output_dim = output_dim; + + float* in_sums = (float*) malloc((size_t) ctx.blocks_per_input_dim * sizeof(float)); + if (in_sums == NULL) return; + for (int32_t b = 0; b < ctx.blocks_per_input_dim; ++b) { + const float* ib = ctx.in_base + (size_t) b * Q51_BLOCK_SIZE; + float sum = 0.0f; + for (int32_t k = 0; k < Q51_BLOCK_SIZE; ++k) sum += ib[k]; + in_sums[b] = sum; + } + ctx.in_sums = in_sums; + skainet_run_rows(worker, &ctx, output_dim); + free(in_sums); +} + +SKAINET_API void skainet_q5_1_matmul( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q5_1_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q5_1_rows_feed); +} + +/* + * Row-major variant (#1192): canonical GGUF file order — see q5_1_rows_rm. + * Numerically equivalent to the feed-order kernel (same per-row block order; -ffast-math FMA/reassociation may differ at ULP scale); threads over rows >= 512 (#1195). + */ +SKAINET_API void skainet_q5_1_matmul_rm( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q5_1_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q5_1_rows_rm); +} diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/q5k_matmul.c b/skainet-backends/skainet-backend-native-cpu/native/src/q5k_matmul.c index 9edb22673..a5909be39 100644 --- a/skainet-backends/skainet-backend-native-cpu/native/src/q5k_matmul.c +++ b/skainet-backends/skainet-backend-native-cpu/native/src/q5k_matmul.c @@ -1,5 +1,6 @@ #include "skainet_kernels.h" #include "skainet_simd.h" +#include "skainet_row_threads.h" #include #include @@ -84,43 +85,17 @@ static inline void skainet_q5k_decode_scales( * arithmetic so -O3 auto-vectorizes on AVX2/NEON. A hand-written NEON * path is layered on behind __ARM_NEON in a later PR. */ -SKAINET_API void skainet_q5k_matmul( - const float* SKAINET_RESTRICT input, - int32_t input_offset, - const uint8_t* SKAINET_RESTRICT weight, - int32_t weight_byte_offset, - int32_t input_dim, - int32_t output_dim, - float* SKAINET_RESTRICT output, - int32_t output_offset +/* + * One block's contribution to out[o] — the loop body of the original kernel, + * shared by the feed-order and row-major entries so both orders (and any row + * partition, #1195) stay bit-identical per output row (#1192). + */ +static inline float q5k_block_term( + const uint8_t* SKAINET_RESTRICT block, + const float* SKAINET_RESTRICT in_block ) { - if (output_dim <= 0 || input_dim <= 0) return; - - const int32_t blocks_per_input_dim = input_dim / Q5K_BLOCK_SIZE; - const float* in_base = input + input_offset; - float* out_base = output + output_offset; - int scale_idx[Q5K_SUB_BLOCKS]; int min_idx[Q5K_SUB_BLOCKS]; - - /* - * Loop order: block OUTER, output row INNER — see q4k_matmul.c for the - * rationale. The weight is block-major (blockIdx*output_dim + o)*176, so for - * a fixed block consecutive `o` are 176 bytes apart: weight bytes are read - * sequentially (cache/prefetch friendly) instead of striding output_dim*176 - * per step, which on the in-order A55 makes every read a cold miss. - * out_base[o] accumulates across blocks (a per-o register `acc` holds the - * inner sum); accumulation order over blocks is unchanged ⇒ numerically - * identical to the o-outer form. - */ - for (int32_t o = 0; o < output_dim; ++o) out_base[o] = 0.0f; - - for (int32_t block_idx = 0; block_idx < blocks_per_input_dim; ++block_idx) { - const float* in_block = in_base + (size_t) block_idx * Q5K_BLOCK_SIZE; - const uint8_t* block = weight + weight_byte_offset - + (size_t)(block_idx * output_dim) * Q5K_BYTES_PER_BLOCK; - - for (int32_t o = 0; o < output_dim; ++o, block += Q5K_BYTES_PER_BLOCK) { float acc = 0.0f; /* d, dMin (FP16 LE -> FP32). */ @@ -208,7 +183,87 @@ SKAINET_API void skainet_q5k_matmul( acc += code_sum_hi * scale_hi - input_sum_hi * offset_hi; } - out_base[o] += acc; + return acc; +} + +/* Everything a row-range worker needs; read-only during the parallel section. */ +typedef struct { + const uint8_t* weight_base; /* weight + weight_byte_offset */ + const float* in_base; /* input + input_offset */ + float* out_base; + int32_t blocks_per_input_dim; + int32_t output_dim; +} q5k_ctx; + +/* Feed-order rows [o_start, o_end): block OUTER, row INNER — see q4k_matmul.c. */ +static void q5k_rows_feed(void* vctx, int32_t o_start, int32_t o_end) { + const q5k_ctx* c = (const q5k_ctx*) vctx; + for (int32_t o = o_start; o < o_end; ++o) c->out_base[o] = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; ++block_idx) { + const float* in_block = c->in_base + (size_t) block_idx * Q5K_BLOCK_SIZE; + const uint8_t* block = c->weight_base + + ((size_t) block_idx * c->output_dim + o_start) * Q5K_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o, block += Q5K_BYTES_PER_BLOCK) { + c->out_base[o] += q5k_block_term(block, in_block); + } + } +} + +/* Row-major rows (#1192): canonical GGUF file order, (o·bpr + b)·176 — mmap-fed as-is. */ +static void q5k_rows_rm(void* vctx, int32_t o_start, int32_t o_end) { + const q5k_ctx* c = (const q5k_ctx*) vctx; + const uint8_t* block = c->weight_base + + (size_t) o_start * c->blocks_per_input_dim * Q5K_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o) { + float acc_row = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; + ++block_idx, block += Q5K_BYTES_PER_BLOCK) { + acc_row += q5k_block_term(block, c->in_base + (size_t) block_idx * Q5K_BLOCK_SIZE); } + c->out_base[o] = acc_row; } } + +/* Shared entry — guards, ctx fill, threaded run (skainet_row_threads.h, #1195). */ +static void q5k_matmul_run( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset, + skainet_row_range_fn worker +) { + if (output_dim <= 0 || input_dim <= 0) return; + q5k_ctx ctx; + ctx.weight_base = weight + weight_byte_offset; + ctx.in_base = input + input_offset; + ctx.out_base = output + output_offset; + ctx.blocks_per_input_dim = input_dim / Q5K_BLOCK_SIZE; + ctx.output_dim = output_dim; + skainet_run_rows(worker, &ctx, output_dim); +} + +SKAINET_API void skainet_q5k_matmul( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q5k_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q5k_rows_feed); +} + +/* + * Row-major variant (#1192): canonical GGUF file order — see q5k_rows_rm. + * Numerically equivalent to the feed-order kernel (same per-row block order; -ffast-math FMA/reassociation may differ at ULP scale); threads over rows >= 512 (#1195). + */ +SKAINET_API void skainet_q5k_matmul_rm( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + q5k_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + q5k_rows_rm); +} diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/q8_0_matmul.c b/skainet-backends/skainet-backend-native-cpu/native/src/q8_0_matmul.c index 780e0d7d7..8e4f2e600 100644 --- a/skainet-backends/skainet-backend-native-cpu/native/src/q8_0_matmul.c +++ b/skainet-backends/skainet-backend-native-cpu/native/src/q8_0_matmul.c @@ -1,5 +1,6 @@ #include "skainet_kernels.h" #include "skainet_simd.h" +#include "skainet_row_threads.h" #include #include @@ -54,71 +55,140 @@ static inline float skainet_fp16_to_fp32(uint16_t h) { return r; } -SKAINET_API void skainet_q8_0_matmul( +#define Q80_BLOCK_SIZE 32 +#define Q80_BYTES_PER_BLOCK 34 + +/* + * One block's contribution to out[o]: `d * Σ_k input[k] * code[k]`. + * Activations are FP32, so int8 codes widen to float and FMA (int8 dotprod + * would need int8 activations). Shared by the feed-order and row-major + * entries, so both orders (and any row partition, #1195) stay bit-identical + * per output row. + */ +static inline float skainet_q8_0_block_term( + const uint8_t* SKAINET_RESTRICT block, + const float* SKAINET_RESTRICT input_block +) { + const uint16_t d_bits = (uint16_t) block[0] | ((uint16_t) block[1] << 8); + const float d = skainet_fp16_to_fp32(d_bits); + const int8_t* SKAINET_RESTRICT codes = (const int8_t*) (block + 2); + float block_sum = 0.0f; +#ifdef SKAINET_HAVE_NEON + float32x4_t accv = vdupq_n_f32(0.0f); + for (int32_t k = 0; k < Q80_BLOCK_SIZE; k += 16) { + const int8x16_t c8 = vld1q_s8(codes + k); + const int16x8_t lo16 = vmovl_s8(vget_low_s8(c8)); + const int16x8_t hi16 = vmovl_s8(vget_high_s8(c8)); + const float32x4_t cf0 = vcvtq_f32_s32(vmovl_s16(vget_low_s16(lo16))); + const float32x4_t cf1 = vcvtq_f32_s32(vmovl_s16(vget_high_s16(lo16))); + const float32x4_t cf2 = vcvtq_f32_s32(vmovl_s16(vget_low_s16(hi16))); + const float32x4_t cf3 = vcvtq_f32_s32(vmovl_s16(vget_high_s16(hi16))); + accv = vfmaq_f32(accv, vld1q_f32(input_block + k), cf0); + accv = vfmaq_f32(accv, vld1q_f32(input_block + k + 4), cf1); + accv = vfmaq_f32(accv, vld1q_f32(input_block + k + 8), cf2); + accv = vfmaq_f32(accv, vld1q_f32(input_block + k + 12), cf3); + } + block_sum = skainet_neon_hadd_f32(accv); +#else + for (int32_t k = 0; k < Q80_BLOCK_SIZE; ++k) { + block_sum += input_block[k] * (float) codes[k]; + } +#endif + return block_sum * d; +} + +/* Everything a row-range worker needs; read-only during the parallel section. */ +typedef struct { + const uint8_t* weight_base; /* weight + weight_byte_offset */ + const float* in_base; /* input + input_offset */ + float* out_base; + int32_t blocks_per_input_dim; + int32_t output_dim; +} skainet_q8_0_ctx; + +/* + * Feed-order rows [o_start, o_end): block OUTER, row INNER — see q4k_matmul.c + * for the cache rationale. out[o] accumulates across blocks in unchanged + * order under any partition. + */ +static void skainet_q8_0_rows_feed(void* vctx, int32_t o_start, int32_t o_end) { + const skainet_q8_0_ctx* c = (const skainet_q8_0_ctx*) vctx; + for (int32_t o = o_start; o < o_end; ++o) c->out_base[o] = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; ++block_idx) { + const float* input_block = c->in_base + (size_t) block_idx * Q80_BLOCK_SIZE; + const uint8_t* block = c->weight_base + + ((size_t) block_idx * c->output_dim + o_start) * Q80_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o, block += Q80_BYTES_PER_BLOCK) { + c->out_base[o] += skainet_q8_0_block_term(block, input_block); + } + } +} + +/* + * Row-major rows [o_start, o_end) (#1189/#1192): canonical GGUF file order — + * (o * blocks_per_row + b) * 34 — an mmap'd tensor is fed as-is, no relayout + * copy; each row's blocks are contiguous on disk. Per-row accumulation order + * matches the feed-order worker's. + */ +static void skainet_q8_0_rows_rm(void* vctx, int32_t o_start, int32_t o_end) { + const skainet_q8_0_ctx* c = (const skainet_q8_0_ctx*) vctx; + const uint8_t* block = c->weight_base + + (size_t) o_start * c->blocks_per_input_dim * Q80_BYTES_PER_BLOCK; + for (int32_t o = o_start; o < o_end; ++o) { + float acc = 0.0f; + for (int32_t block_idx = 0; block_idx < c->blocks_per_input_dim; + ++block_idx, block += Q80_BYTES_PER_BLOCK) { + acc += skainet_q8_0_block_term(block, c->in_base + (size_t) block_idx * Q80_BLOCK_SIZE); + } + c->out_base[o] = acc; + } +} + +/* Shared entry — guards, ctx fill, threaded run (skainet_row_threads.h, #1195). */ +static void skainet_q8_0_matmul_run( const float* SKAINET_RESTRICT input, int32_t input_offset, const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, int32_t input_dim, int32_t output_dim, - float* SKAINET_RESTRICT output, int32_t output_offset + float* SKAINET_RESTRICT output, int32_t output_offset, + skainet_row_range_fn worker ) { if (output_dim <= 0) return; if (input_dim <= 0) { - for (int32_t o = 0; o < output_dim; ++o) { - output[output_offset + o] = 0.0f; - } + for (int32_t o = 0; o < output_dim; ++o) output[output_offset + o] = 0.0f; return; } + skainet_q8_0_ctx ctx; + ctx.weight_base = weight + weight_byte_offset; + ctx.in_base = input + input_offset; + ctx.out_base = output + output_offset; + ctx.blocks_per_input_dim = input_dim / Q80_BLOCK_SIZE; + ctx.output_dim = output_dim; + skainet_run_rows(worker, &ctx, output_dim); +} - const int32_t BLOCK_SIZE = 32; - const int32_t BYTES_PER_BLOCK = 34; - const int32_t blocks_per_input_dim = input_dim / BLOCK_SIZE; - float* SKAINET_RESTRICT out_base = output + output_offset; - - /* - * Loop order: block OUTER, output row INNER — see q4k_matmul.c for the - * rationale. The weight is block-major (blockIdx*output_dim + o)*34, so for - * a fixed block consecutive `o` are 34 bytes apart: weight bytes are read - * sequentially instead of striding output_dim*34 per step, which on the - * in-order A55 makes every read a cold cache miss. out_base[o] accumulates - * across blocks; accumulation order is unchanged ⇒ numerically identical. - */ - for (int32_t o = 0; o < output_dim; ++o) out_base[o] = 0.0f; - - for (int32_t block_idx = 0; block_idx < blocks_per_input_dim; ++block_idx) { - const float* SKAINET_RESTRICT input_block = - input + input_offset + (size_t) block_idx * BLOCK_SIZE; - const uint8_t* SKAINET_RESTRICT block = - weight + weight_byte_offset + - (size_t)(block_idx * output_dim) * BYTES_PER_BLOCK; +SKAINET_API void skainet_q8_0_matmul( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + skainet_q8_0_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + skainet_q8_0_rows_feed); +} - for (int32_t o = 0; o < output_dim; ++o, block += BYTES_PER_BLOCK) { - uint16_t d_bits = (uint16_t) block[0] | ((uint16_t) block[1] << 8); - float d = skainet_fp16_to_fp32(d_bits); - const int8_t* SKAINET_RESTRICT codes = (const int8_t*) (block + 2); - float block_sum = 0.0f; -#ifdef SKAINET_HAVE_NEON - /* Activations are FP32, so widen int8 codes to float and FMA - * (int8 dotprod would need int8 activations — see plan note). */ - float32x4_t accv = vdupq_n_f32(0.0f); - for (int32_t k = 0; k < BLOCK_SIZE; k += 16) { - const int8x16_t c8 = vld1q_s8(codes + k); - const int16x8_t lo16 = vmovl_s8(vget_low_s8(c8)); - const int16x8_t hi16 = vmovl_s8(vget_high_s8(c8)); - const float32x4_t cf0 = vcvtq_f32_s32(vmovl_s16(vget_low_s16(lo16))); - const float32x4_t cf1 = vcvtq_f32_s32(vmovl_s16(vget_high_s16(lo16))); - const float32x4_t cf2 = vcvtq_f32_s32(vmovl_s16(vget_low_s16(hi16))); - const float32x4_t cf3 = vcvtq_f32_s32(vmovl_s16(vget_high_s16(hi16))); - accv = vfmaq_f32(accv, vld1q_f32(input_block + k), cf0); - accv = vfmaq_f32(accv, vld1q_f32(input_block + k + 4), cf1); - accv = vfmaq_f32(accv, vld1q_f32(input_block + k + 8), cf2); - accv = vfmaq_f32(accv, vld1q_f32(input_block + k + 12), cf3); - } - block_sum = skainet_neon_hadd_f32(accv); -#else - for (int32_t k = 0; k < BLOCK_SIZE; ++k) { - block_sum += input_block[k] * (float) codes[k]; - } -#endif - out_base[o] += block_sum * d; - } - } +/* + * Row-major variant (#1192): the weight stays in canonical GGUF file order — + * see skainet_q8_0_rows_rm. Numerically equivalent to the feed-order kernel (same per-row block order; -ffast-math FMA/reassociation may differ at ULP scale); threads + * over output rows when outputDim >= 512 (#1195). + */ +SKAINET_API void skainet_q8_0_matmul_rm( + const float* SKAINET_RESTRICT input, int32_t input_offset, + const uint8_t* SKAINET_RESTRICT weight, int32_t weight_byte_offset, + int32_t input_dim, int32_t output_dim, + float* SKAINET_RESTRICT output, int32_t output_offset +) { + skainet_q8_0_matmul_run(input, input_offset, weight, weight_byte_offset, + input_dim, output_dim, output, output_offset, + skainet_q8_0_rows_rm); } diff --git a/skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/FfmRowMajorKernels.kt b/skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/FfmRowMajorKernels.kt new file mode 100644 index 000000000..31d276117 --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/FfmRowMajorKernels.kt @@ -0,0 +1,171 @@ +package sk.ainet.exec.kernel + +import java.lang.foreign.Arena +import java.lang.foreign.FunctionDescriptor +import java.lang.foreign.Linker +import java.lang.foreign.MemorySegment +import java.lang.foreign.ValueLayout +import java.lang.invoke.MethodHandle +import sk.ainet.backend.api.kernel.KernelDispatch +import sk.ainet.backend.api.kernel.KernelKey +import sk.ainet.backend.api.kernel.LayoutClass +import sk.ainet.backend.api.kernel.OperandKey +import sk.ainet.backend.api.kernel.ReferenceMatmulKernel +import sk.ainet.backend.api.kernel.ViewKernel +import sk.ainet.lang.memory.BlockOrder +import sk.ainet.lang.memory.DirectBufferStorage +import sk.ainet.lang.memory.ExperimentalMemoryApi +import sk.ainet.lang.memory.Format +import sk.ainet.lang.memory.MappedBufferStorage +import sk.ainet.lang.memory.Storage +import sk.ainet.lang.memory.TensorView +import sk.ainet.lang.memory.trace.NoopTraceSink +import sk.ainet.lang.memory.trace.TraceEvent +import sk.ainet.lang.memory.trace.TraceSink +import sk.ainet.lang.tensor.storage.TensorEncoding +import sk.ainet.lang.types.FP32 + +/** + * The JVM (FFM) row-major tier (#1191): serves `BLOCKED_ROW_MAJOR` packed weights through the + * same `_rm` C symbols the Android JNI tier uses (#1189/#1192), polymorphic in the weight's + * storage: + * + * - [MappedBufferStorage]/[DirectBufferStorage] → `MemorySegment.ofBuffer` — **zero-copy**: the + * kernel reads the mmap'd GGUF pages directly, which is the point of `WeightResidency.MAPPED`. + * Before this tier, a mapped packed weight on the JVM fell to the decoding reference kernel. + * - Heap `ByteArray` → staged into a confined arena per call, the same cost class as the + * existing heap FFM kernels ([NativeQ4KMatmulKernel] et al. stage their weight too). + * + * Activations and outputs are heap `FloatArray`s staged per call (small: `k` and `n` floats). + * Fallbacks to the decoding reference are announced to the trace sink (#1193) — never silent. + */ +@ExperimentalMemoryApi +public class FfmRowMajorMatmulKernel internal constructor( + encodingName: String, + override val key: KernelKey, + private val handle: MethodHandle, +) : ViewKernel { + + override val name: String = "ffm-rowmajor-$encodingName" + + override fun run(inputs: List, out: TensorView): Unit = run(inputs, out, NoopTraceSink) + + override fun run(inputs: List, out: TensorView, sink: TraceSink) { + require(inputs.size == 2) { "matmul takes two operands" } + val a = inputs[0] + val w = inputs[1] + val rows = a.shape[0] + val k = a.shape[1] + val n = w.shape[0] + require(w.shape[1] == k) { "inner dimensions disagree: [$rows, $k] × [$n, ${w.shape[1]}]" } + check(w.layout.blockOrder == BlockOrder.ROW_MAJOR) { + "$name reads canonical row-major weights; this one is ${w.layout.blockOrder}" + } + + val aHeap = a.storage as? Storage.Heap ?: return fallback(inputs, out, sink, "activation storage ${a.storage::class.simpleName}") + val oHeap = out.storage as? Storage.Heap ?: return fallback(inputs, out, sink, "output storage ${out.storage::class.simpleName}") + val activation = aHeap.floats ?: return fallback(inputs, out, sink, "activation is not a FloatArray") + val output = oHeap.floats ?: return fallback(inputs, out, sink, "output is not a FloatArray") + if (!a.isContiguous) return fallback(inputs, out, sink, "strided activation") + + val weightOffset = (w.layout.offsetElements * w.layout.elementBytes).toInt() + + Arena.ofConfined().use { arena -> + val weightSeg: MemorySegment + val weightSegOffset: Int + when (val s = w.storage) { + is MappedBufferStorage, is DirectBufferStorage -> { + val buffer = if (s is MappedBufferStorage) s.buffer() else (s as DirectBufferStorage).buffer() + weightSeg = MemorySegment.ofBuffer(buffer) // zero-copy over the mapping + weightSegOffset = weightOffset + } + is Storage.Heap -> { + val bytes = s.bytes ?: return fallback(inputs, out, sink, "heap weight is not a ByteArray") + val length = s.sizeBytes - weightOffset + val seg = arena.allocate(length, 1L) + MemorySegment.copy(bytes, s.arrayOffset + weightOffset, seg, ValueLayout.JAVA_BYTE, 0L, length.toInt()) + weightSeg = seg + weightSegOffset = 0 + } + else -> return fallback(inputs, out, sink, "weight storage ${s::class.simpleName}") + } + + val inSeg = arena.allocate(k.toLong() * java.lang.Float.BYTES, ValueLayout.JAVA_FLOAT.byteAlignment()) + val outSeg = arena.allocate(n.toLong() * java.lang.Float.BYTES, ValueLayout.JAVA_FLOAT.byteAlignment()) + for (r in 0 until rows) { + MemorySegment.copy( + activation, aHeap.arrayOffset + (a.layout.offsetElements + r.toLong() * k).toInt(), + inSeg, ValueLayout.JAVA_FLOAT, 0L, k, + ) + handle.invoke(inSeg, 0, weightSeg, weightSegOffset, k, n, outSeg, 0) + MemorySegment.copy( + outSeg, ValueLayout.JAVA_FLOAT, 0L, + output, oHeap.arrayOffset + (out.layout.offsetElements + r.toLong() * n).toInt(), n, + ) + } + } + } + + private fun fallback(inputs: List, out: TensorView, sink: TraceSink, reason: String) { + // #1193: the decoding reference is ~1000× slower — that must be a trace line, never silent. + if (sink.isEnabled) { + sink.emit( + TraceEvent.KernelRun( + op = key.op, + kernel = "reference-fallback from $name: $reason", + inputs = inputs.map { it.id }, + output = out.id, + ), + ) + } + ReferenceMatmulKernel(key).run(inputs, out) + } + + public companion object { + /** The key this kernel serves: dense contiguous FP32 activation × row-major packed weight. */ + public fun keyFor(encoding: TensorEncoding): KernelKey = KernelKey( + op = "matmul", + operands = listOf( + OperandKey.contiguous(Format.dense(FP32)), + OperandKey(Format(FP32, encoding), LayoutClass.BLOCKED_ROW_MAJOR), + ), + ) + } +} + +/** + * Registers the FFM row-major kernels (#1191) into [KernelDispatch] — the JVM counterpart of + * Android's `JniMappedKernelPack.install()`. No-op when the native library is unavailable; + * a symbol that fails to bind simply leaves that encoding to the reference kernel. + */ +@ExperimentalMemoryApi +public object FfmRowMajorKernelPack { + + private val DESCRIPTOR = FunctionDescriptor.ofVoid( + ValueLayout.ADDRESS, ValueLayout.JAVA_INT, // input, input_offset + ValueLayout.ADDRESS, ValueLayout.JAVA_INT, // weight, weight_byte_offset + ValueLayout.JAVA_INT, ValueLayout.JAVA_INT, // input_dim, output_dim + ValueLayout.ADDRESS, ValueLayout.JAVA_INT, // output, output_offset + ) + + private val SYMBOLS: List> = listOf( + TensorEncoding.Q4_K to "skainet_q4k_matmul_rm", + TensorEncoding.Q6_K to "skainet_q6k_matmul_rm", + TensorEncoding.Q5_K to "skainet_q5k_matmul_rm", + TensorEncoding.Q8_0 to "skainet_q8_0_matmul_rm", + TensorEncoding.Q4_0 to "skainet_q4_0_matmul_rm", + TensorEncoding.Q5_0 to "skainet_q5_0_matmul_rm", + TensorEncoding.Q5_1 to "skainet_q5_1_matmul_rm", + ) + + public fun install() { + val lookup = NativeLibraryLoader.lookup() ?: return + for ((encoding, symbol) in SYMBOLS) { + val sym = lookup.find(symbol).orElse(null) ?: continue + val handle = runCatching { Linker.nativeLinker().downcallHandle(sym, DESCRIPTOR) }.getOrNull() ?: continue + KernelDispatch.register( + FfmRowMajorMatmulKernel(encoding.name, FfmRowMajorMatmulKernel.keyFor(encoding), handle), + ) + } + } +} diff --git a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/FfmRowMajorDispatchTest.kt b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/FfmRowMajorDispatchTest.kt new file mode 100644 index 000000000..a074fcf06 --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/FfmRowMajorDispatchTest.kt @@ -0,0 +1,133 @@ +package sk.ainet.exec.kernel + +import java.nio.file.Files +import kotlin.random.Random +import kotlin.test.AfterTest +import kotlin.test.Test +import kotlin.test.assertTrue +import sk.ainet.backend.api.kernel.KernelDispatch +import sk.ainet.lang.memory.ExperimentalMemoryApi +import sk.ainet.lang.memory.MappedBufferStorage +import sk.ainet.lang.memory.Storage +import sk.ainet.lang.memory.TensorView +import sk.ainet.lang.memory.trace.RecordingTraceSink +import sk.ainet.lang.memory.trace.TraceEvent +import sk.ainet.lang.tensor.Shape +import sk.ainet.lang.tensor.data.BufferPackedTensorData +import sk.ainet.lang.tensor.data.Q4_KBlockTensorData +import sk.ainet.lang.tensor.storage.TensorEncoding +import sk.ainet.lang.types.FP32 + +/** + * #1191 end-to-end on the JVM: a mapped Q4_K weight dispatches to the FFM row-major kernel + * (zero-copy over the mmap), produces the same numbers as the decoded reference, and the trace + * names the kernel that ran. Also pins #1193's visibility contract: when the fast path cannot + * serve, the fallback announces itself in the trace instead of silently decoding. + */ +@OptIn(ExperimentalMemoryApi::class) +class FfmRowMajorDispatchTest { + + @AfterTest + fun reset() = KernelDispatch.clearForTesting() + + private fun q4kPayload(rows: Int, k: Int, seed: Int): ByteArray { + val bpr = k / 256 + val rng = Random(seed) + val bytes = ByteArray(rows * bpr * 144) + rng.nextBytes(bytes) + for (b in 0 until rows * bpr) { + bytes[b * 144] = 0x00; bytes[b * 144 + 1] = 0x3C + bytes[b * 144 + 2] = 0x00; bytes[b * 144 + 3] = 0x3C + } + return bytes + } + + @Test + fun mapped_q4k_dispatches_to_the_ffm_row_major_kernel() { + KernelDispatch.clearForTesting() + FfmRowMajorKernelPack.install() + + val rows = 64 + val k = 512 + val payload = q4kPayload(rows, k, seed = 6) + val file = Files.createTempFile("ffm-rm", ".bin") + try { + Files.write(file, payload) + val storage = MappedBufferStorage.map(file, 0, payload.size.toLong()) + val weight = BufferPackedTensorData(Shape(rows, k), storage, TensorEncoding.Q4_K) + + val act = FloatArray(k) { Random(7 + it).nextFloat() - 0.5f } + val actStorage = Storage.Heap.wrap(act, mutable = false) + val actView = TensorView.dense(actStorage, Shape(1, k), FP32) + val out = FloatArray(rows) + val outView = TensorView.dense(Storage.Heap.wrap(out, mutable = true), Shape(1, rows), FP32) + + val sink = RecordingTraceSink() + KernelDispatch.matmul(actView, weight.packedView, outView, sink = sink) + + // Trace names the FFM kernel — proof it served, not the reference. + val runs = sink.events().filterIsInstance() + assertTrue(runs.any { "ffm-rowmajor-Q4_K" in it.kernel }, "expected ffm kernel in trace, got: ${runs.map { it.kernel }}") + assertTrue(runs.none { "reference-fallback" in it.kernel }, "no fallback expected: ${runs.map { it.kernel }}") + + // Numbers equal the feed-order kernel on permuted bytes — the same W4A8 family + // (Q4_K quantizes the activation to int8; the exact-float decode is NOT the oracle, + // see #944), so only -ffast-math reassociation separates them. + val bpr = k / 256 + val feed = ByteArray(payload.size) + for (o in 0 until rows) for (b in 0 until bpr) { + payload.copyInto(feed, (b * rows + o) * 144, (o * bpr + b) * 144, (o * bpr + b + 1) * 144) + } + val expected = FloatArray(rows) + NativeQ4KMatmulKernel.matmul(act, 0, feed, 0, k, rows, expected, 0) + for (o in 0 until rows) { + assertTrue(kotlin.math.abs(expected[o] - out[o]) <= maxOf(1e-5f, 1e-5f * kotlin.math.abs(expected[o])), + "row $o: dispatch ${out[o]} vs feed-order ${expected[o]}") + } + } finally { + Files.deleteIfExists(file) + } + } + + @Test + fun a_fallback_announces_itself_in_the_trace() { + // #1193: a packed kernel that cannot serve its operands must say so. Byte-backed output + // storage is un-servable for every packed matmul bridge — the kernel punts to the + // decoding reference AND leaves a trace line naming itself and the reason. + val kernel = sk.ainet.backend.api.kernel.PackedViewMatmulKernel( + "test", "Q4_K", + sk.ainet.backend.api.kernel.PackedViewMatmulKernel.keyFor(TensorEncoding.Q4_K), + ) { _, _, _, _, _, _, _, _ -> } + + val rows = 2 + val k = 256 + val payload = q4kPayload(rows, k, seed = 8) + val weightView = TensorView.packed( + storage = Storage.Heap.wrap(payload, mutable = false), + shape = Shape(rows, k), + encoding = TensorEncoding.Q4_K, + decoder = sk.ainet.lang.memory.PackedBlockDecoder(Q4_KBlockTensorData(Shape(rows, k), payload)), + blockOrder = sk.ainet.lang.memory.BlockOrder.INPUT_BLOCK_MAJOR, + ) + // A STRIDED activation (every second float of a doubled backing row): the C kernel + // would mis-index it, so the bridge must punt — and the reference reads any layout. + val actBacking = FloatArray(k * 2) { Random(9 + it).nextFloat() - 0.5f } + val actView = TensorView( + Shape(1, k), + sk.ainet.lang.memory.Format.dense(FP32), + sk.ainet.lang.memory.Layout(Shape(1, k), intArrayOf(k * 2, 2)), + Storage.Heap.wrap(actBacking, mutable = false), + null, + ) + val outView = TensorView.dense(Storage.Heap.wrap(FloatArray(rows), mutable = true), Shape(1, rows), FP32) + + val sink = RecordingTraceSink() + kernel.run(listOf(actView, weightView), outView, sink) + + val runs = sink.events().filterIsInstance() + assertTrue( + runs.any { "reference-fallback" in it.kernel && "strided" in it.kernel }, + "fallback must announce itself; trace: ${runs.map { it.kernel }}", + ) + } +} diff --git a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/KernelSupportMatrixTest.kt b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/KernelSupportMatrixTest.kt index d4cbc7cd6..5b6c5fd96 100644 --- a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/KernelSupportMatrixTest.kt +++ b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/KernelSupportMatrixTest.kt @@ -59,13 +59,15 @@ class KernelSupportMatrixTest { /** * Mapped serving (#1189): kernels that read the weight in canonical row-major GGUF file * order straight from off-heap bytes (mmap/direct buffer) — the `_rm` symbols behind - * `JniBufferPackedMatmulKernel` on Android. The JVM/FFM tier is #1191; other platforms - * and formats are #1192. Must stay in lockstep with + * `JniBufferPackedMatmulKernel` on Android. K/N and remaining formats: #1192 follow-ups. Must stay in lockstep with * `StorageCapabilities.MAPPED_SERVABLE_ENCODINGS` (dense F32 is mapped there too, but as * element-view serving, not a matmul kernel — it has no row here on purpose). */ private fun mappedTiers(): List = listOf( - Tier("native-jni-direct", 100, setOf("Android"), setOf("Q4_K", "Q6_K")), + Tier("native-jni-direct", 100, setOf("Android"), + setOf("Q4_K", "Q6_K", "Q5_K", "Q8_0", "Q4_0", "Q5_0", "Q5_1")), + Tier("ffm-rowmajor", 100, setOf("JVM"), + setOf("Q4_K", "Q6_K", "Q5_K", "Q8_0", "Q4_0", "Q5_0", "Q5_1")), ) private fun best(fmt: String, platform: String, tiers: List): String? = diff --git a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt index 5ec8da6d6..9d01dc615 100644 --- a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt +++ b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt @@ -28,6 +28,12 @@ class RowMajorMatmulParityTest { const val BLOCK = 256 const val Q4K_BPB = 144 const val Q6K_BPB = 210 + const val Q5K_BPB = 176 + const val SMALL = 32 + const val Q80_BPB = 34 + const val Q40_BPB = 18 + const val Q50_BPB = 22 + const val Q51_BPB = 24 } private fun bindRm(symbol: String): MethodHandle? { @@ -105,6 +111,18 @@ class RowMajorMatmulParityTest { return out } + + /** + * Cross-ORDER comparisons are numerically-equivalent, not bit-exact: under -O3 -ffast-math + * the compiler may contract/re-associate float accumulation differently between the two loop + * shapes (q5_1 measured 2 ULP apart; the integer-dot formats happen to match exactly). + * Threaded-vs-solo comparisons stay bit-exact — same function, same per-row order. + */ + private fun assertClose(expected: Float, got: Float, label: String) { + val tol = maxOf(1e-5f, 1e-5f * kotlin.math.abs(expected)) + assertTrue(kotlin.math.abs(expected - got) <= tol, "$label: $expected vs $got (tol $tol)") + } + private fun assertQ4kParity(inputDim: Int, outputDim: Int, seed: Int, pad: Int = 0) { val bpr = inputDim / BLOCK val rowMajor = randomBlocks(bpr * outputDim, Q4K_BPB, intArrayOf(0, 2), seed) @@ -117,7 +135,7 @@ class RowMajorMatmulParityTest { val padded = if (pad == 0) rowMajor else ByteArray(pad) + rowMajor val got = callRm(q4kRm!!, input, padded, pad, inputDim, outputDim) for (o in 0 until outputDim) { - assertEquals(expected[o].toRawBits(), got[o].toRawBits(), "Q4_K row $o diverged: ${expected[o]} vs ${got[o]}") + assertClose(expected[o], got[o], "Q4_K row $o") } } @@ -134,7 +152,7 @@ class RowMajorMatmulParityTest { val padded = if (pad == 0) rowMajor else ByteArray(pad) + rowMajor val got = callRm(q6kRm!!, input, padded, pad, inputDim, outputDim) for (o in 0 until outputDim) { - assertEquals(expected[o].toRawBits(), got[o].toRawBits(), "Q6_K row $o diverged: ${expected[o]} vs ${got[o]}") + assertClose(expected[o], got[o], "Q6_K row $o") } } @@ -189,4 +207,48 @@ class RowMajorMatmulParityTest { assertEquals(single[0].toRawBits(), full[o].toRawBits(), "Q6_K row $o: threaded diverged from solo") } } + + // ---- #1192: the four remaining GGML block formats, same two-oracle scheme ---- + + private fun assertParity( + rmSymbol: String, + feed: (FloatArray, Int, ByteArray, Int, Int, Int, FloatArray, Int) -> Unit, + blockSize: Int, bpb: Int, fp16At: IntArray, + inputDim: Int, outputDim: Int, seed: Int, + ) { + val handle = bindRm(rmSymbol) + assertNotNull(handle, "$rmSymbol must be bindable") + val bpr = inputDim / blockSize + val rowMajor = randomBlocks(bpr * outputDim, bpb, fp16At, seed) + val feedBytes = toFeedOrder(rowMajor, outputDim, bpr, bpb) + val input = FloatArray(inputDim) { Random(seed + it).nextFloat() - 0.5f } + + val expected = FloatArray(outputDim) + feed(input, 0, feedBytes, 0, inputDim, outputDim, expected, 0) + val got = callRm(handle, input, rowMajor, 0, inputDim, outputDim) + for (o in 0 until outputDim) { + assertClose(expected[o], got[o], "$rmSymbol row $o") + } + } + + @Test fun q5k_rm_parity() = assertParity( + "skainet_q5k_matmul_rm", NativeQ5KMatmulKernel::matmul, BLOCK, Q5K_BPB, intArrayOf(0, 2), 512, 48, seed = 21) + @Test fun q5k_rm_parity_threaded() = assertParity( + "skainet_q5k_matmul_rm", NativeQ5KMatmulKernel::matmul, BLOCK, Q5K_BPB, intArrayOf(0, 2), 512, 1024, seed = 22) + @Test fun q80_rm_parity() = assertParity( + "skainet_q8_0_matmul_rm", NativeQ8_0MatmulKernel::matmul, SMALL, Q80_BPB, intArrayOf(0), 256, 48, seed = 23) + @Test fun q80_rm_parity_threaded() = assertParity( + "skainet_q8_0_matmul_rm", NativeQ8_0MatmulKernel::matmul, SMALL, Q80_BPB, intArrayOf(0), 256, 1024, seed = 24) + @Test fun q40_rm_parity() = assertParity( + "skainet_q4_0_matmul_rm", NativeQ4_0MatmulKernel::matmul, SMALL, Q40_BPB, intArrayOf(0), 256, 48, seed = 25) + @Test fun q40_rm_parity_threaded() = assertParity( + "skainet_q4_0_matmul_rm", NativeQ4_0MatmulKernel::matmul, SMALL, Q40_BPB, intArrayOf(0), 256, 1024, seed = 26) + @Test fun q50_rm_parity() = assertParity( + "skainet_q5_0_matmul_rm", NativeQ5_0MatmulKernel::matmul, SMALL, Q50_BPB, intArrayOf(0), 256, 48, seed = 27) + @Test fun q50_rm_parity_threaded() = assertParity( + "skainet_q5_0_matmul_rm", NativeQ5_0MatmulKernel::matmul, SMALL, Q50_BPB, intArrayOf(0), 256, 1024, seed = 28) + @Test fun q51_rm_parity() = assertParity( + "skainet_q5_1_matmul_rm", NativeQ5_1MatmulKernel::matmul, SMALL, Q51_BPB, intArrayOf(0, 2), 256, 48, seed = 29) + @Test fun q51_rm_parity_threaded() = assertParity( + "skainet_q5_1_matmul_rm", NativeQ5_1MatmulKernel::matmul, SMALL, Q51_BPB, intArrayOf(0, 2), 256, 1024, seed = 30) } diff --git a/skainet-io/skainet-io-core/src/jvmAndroidMain/kotlin/sk/ainet/io/JvmMappedFile.kt b/skainet-io/skainet-io-core/src/jvmAndroidMain/kotlin/sk/ainet/io/JvmMappedFile.kt index 39ba010b4..86c8ed149 100644 --- a/skainet-io/skainet-io-core/src/jvmAndroidMain/kotlin/sk/ainet/io/JvmMappedFile.kt +++ b/skainet-io/skainet-io-core/src/jvmAndroidMain/kotlin/sk/ainet/io/JvmMappedFile.kt @@ -33,7 +33,7 @@ public class JvmMappedFile private constructor( } /** - * Q4_K/Q6_K payloads as [sk.ainet.lang.tensor.data.BufferPackedTensorData] borrowing a slice + * GGML block-format payloads as [sk.ainet.lang.tensor.data.BufferPackedTensorData] borrowing a slice * of the one file mapping (#1189) — the packed counterpart of [denseFloats]: zero heap bytes, * blocks left in canonical row-major file order for the buffer-reading kernels. */ @@ -45,6 +45,11 @@ public class JvmMappedFile private constructor( ): TensorData<*, *>? = when (encoding) { sk.ainet.lang.tensor.storage.TensorEncoding.Q4_K, sk.ainet.lang.tensor.storage.TensorEncoding.Q6_K, + sk.ainet.lang.tensor.storage.TensorEncoding.Q5_K, + sk.ainet.lang.tensor.storage.TensorEncoding.Q8_0, + sk.ainet.lang.tensor.storage.TensorEncoding.Q4_0, + sk.ainet.lang.tensor.storage.TensorEncoding.Q5_0, + sk.ainet.lang.tensor.storage.TensorEncoding.Q5_1, -> { val length = checkNotNull(encoding.physicalBytes(shape.volume.toLong())) { "${encoding.name} has size-determinate blocks" diff --git a/skainet-io/skainet-io-gguf/src/commonMain/kotlin/sk/ainet/io/gguf/StreamingGgufParametersLoader.kt b/skainet-io/skainet-io-gguf/src/commonMain/kotlin/sk/ainet/io/gguf/StreamingGgufParametersLoader.kt index f6641edff..a8c70c641 100644 --- a/skainet-io/skainet-io-gguf/src/commonMain/kotlin/sk/ainet/io/gguf/StreamingGgufParametersLoader.kt +++ b/skainet-io/skainet-io-gguf/src/commonMain/kotlin/sk/ainet/io/gguf/StreamingGgufParametersLoader.kt @@ -291,7 +291,7 @@ public class StreamingGgufParametersLoader( onProgress(current, total, tensorInfo.name) continue } - // #1189: a Q4_K/Q6_K tensor under MAPPED staging is the packed counterpart of the + // #1189/#1192: a GGML block-format tensor under MAPPED staging is the packed counterpart of the // dense case above — served as a view over the file-backed pages, blocks left in // canonical row-major order, never a heap ByteArray. Only when nothing asks for a // different form: a dequantize/planes request or KERNEL_FEED order needs the bytes @@ -307,6 +307,11 @@ public class StreamingGgufParametersLoader( val enc = when (tensorInfo.tensorType) { GGMLQuantizationType.Q4_K -> sk.ainet.lang.tensor.storage.TensorEncoding.Q4_K GGMLQuantizationType.Q6_K -> sk.ainet.lang.tensor.storage.TensorEncoding.Q6_K + GGMLQuantizationType.Q5_K -> sk.ainet.lang.tensor.storage.TensorEncoding.Q5_K + GGMLQuantizationType.Q8_0 -> sk.ainet.lang.tensor.storage.TensorEncoding.Q8_0 + GGMLQuantizationType.Q4_0 -> sk.ainet.lang.tensor.storage.TensorEncoding.Q4_0 + GGMLQuantizationType.Q5_0 -> sk.ainet.lang.tensor.storage.TensorEncoding.Q5_0 + GGMLQuantizationType.Q5_1 -> sk.ainet.lang.tensor.storage.TensorEncoding.Q5_1 else -> null } enc?.let { encoding -> diff --git a/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/MappedPackedStagingTest.kt b/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/MappedPackedStagingTest.kt index c5b733d28..a393806b1 100644 --- a/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/MappedPackedStagingTest.kt +++ b/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/MappedPackedStagingTest.kt @@ -62,13 +62,13 @@ class MappedPackedStagingTest { val q6k = mapped.getValue("w_q6k").data assertTrue(q6k is BufferPackedTensorData, "Q6_K under MAPPED must be buffer-packed, got ${q6k::class.simpleName}") - // Unsupported-by-#1189 types keep their heap staging. + // #1192: Q8_0 (the k-quant fallback format) is mapped-servable too. val q80 = mapped.getValue("w_q80").data - assertTrue(q80 !is BufferPackedTensorData, "Q8_0 has no buffer kernel yet; stays heap-staged") + assertTrue(q80 is BufferPackedTensorData, "Q8_0 under MAPPED is buffer-packed since #1192, got ${q80::class.simpleName}") // Values equal the heap load, element for element. val heap = load(f, WeightForm(residency = WeightResidency.HEAP)) - for (name in listOf("w_q4k", "w_q6k")) { + for (name in listOf("w_q4k", "w_q6k", "w_q80")) { assertContentEquals( heap.getValue(name).data.copyToFloatArray(), mapped.getValue(name).data.copyToFloatArray(), diff --git a/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/StagingPolicyParityTest.kt b/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/StagingPolicyParityTest.kt index 6e44eae5d..543df1335 100644 --- a/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/StagingPolicyParityTest.kt +++ b/skainet-io/skainet-io-gguf/src/jvmTest/kotlin/sk/ainet/io/gguf/StagingPolicyParityTest.kt @@ -97,9 +97,9 @@ class StagingPolicyParityTest { "Q4_K under MAPPED stays off-heap (#1189)", ) assertEquals( - heap.getValue("w_q80").data::class.simpleName, + "BufferPackedTensorData", mapped.getValue("w_q80").data::class.simpleName, - "Q8_0 staging is unchanged by the mapping", + "Q8_0 under MAPPED stays off-heap (#1192)", ) } finally { f.delete() diff --git a/skainet-lang/skainet-lang-core/src/commonMain/kotlin/sk/ainet/lang/memory/plan/AllocationResolver.kt b/skainet-lang/skainet-lang-core/src/commonMain/kotlin/sk/ainet/lang/memory/plan/AllocationResolver.kt index f6393f377..21b506b75 100644 --- a/skainet-lang/skainet-lang-core/src/commonMain/kotlin/sk/ainet/lang/memory/plan/AllocationResolver.kt +++ b/skainet-lang/skainet-lang-core/src/commonMain/kotlin/sk/ainet/lang/memory/plan/AllocationResolver.kt @@ -27,9 +27,14 @@ public data class StorageCapabilities( val mappedServableEncodings: Set = MAPPED_SERVABLE_DEFAULT, ) { public companion object { - /** What the loaders can serve from mapped pages today: dense FP32, Q4_K, Q6_K (#1189). */ - public val MAPPED_SERVABLE_DEFAULT: Set = - setOf(TensorEncoding.Dense(4), TensorEncoding.Q4_K, TensorEncoding.Q6_K) + /** What the loaders can serve from mapped pages today: dense FP32 (#921) and the GGML + * block formats (#1189 Q4_K/Q6_K, #1192 the rest). Ternary formats need a load-time + * repack and stay heap until the repack cache lands. */ + public val MAPPED_SERVABLE_DEFAULT: Set = setOf( + TensorEncoding.Dense(4), + TensorEncoding.Q4_K, TensorEncoding.Q6_K, TensorEncoding.Q5_K, + TensorEncoding.Q8_0, TensorEncoding.Q4_0, TensorEncoding.Q5_0, TensorEncoding.Q5_1, + ) /** The platform this code is running on. */ public fun current(): StorageCapabilities = StorageCapabilities( @@ -150,7 +155,7 @@ public object AllocationResolver { form?.residency == WeightResidency.MAPPED && weight.format.encoding !in platform.mappedServableEncodings -> "form asks MAPPED but no loader/kernel serves ${weight.format.encoding.name} " + - "from a mapping yet (#1189 covers dense F32, Q4_K, Q6_K) — heap staging" + "from a mapping yet (see StorageCapabilities.mappedServableEncodings) — heap staging" else -> "resident ${MemoryPlans.formatBytes(weight.residentBytes)} vs off-heap threshold " + MemoryPlans.formatBytes(profile.offHeapThresholdBytes) + diff --git a/skainet-lang/skainet-lang-core/src/commonTest/kotlin/sk/ainet/lang/memory/plan/MappedBudgetPlanTest.kt b/skainet-lang/skainet-lang-core/src/commonTest/kotlin/sk/ainet/lang/memory/plan/MappedBudgetPlanTest.kt index 129a058c0..d66a102b0 100644 --- a/skainet-lang/skainet-lang-core/src/commonTest/kotlin/sk/ainet/lang/memory/plan/MappedBudgetPlanTest.kt +++ b/skainet-lang/skainet-lang-core/src/commonTest/kotlin/sk/ainet/lang/memory/plan/MappedBudgetPlanTest.kt @@ -67,28 +67,31 @@ class MappedBudgetPlanTest { @Test fun encodings_nobody_serves_from_a_mapping_stay_heap_charged_even_under_MAPPED() { + // was Q8_0 until #1192 made it servable — the pin moves to the ternary family val bytes = 100L * 1024 * 1024 val mappedForm = WeightForm(residency = WeightResidency.MAPPED) - val q80 = PlanTensor("w", null, Format(FP32, TensorEncoding.Q8_0), bytes / 34 * 32, bytes, mappedForm) - val plan = MemoryPlans.plan(input(listOf(q80)), Budget.of(256L * 1024 * 1024), platform) + // TQ2_0 needs a load-time repack — a copy cannot page from the file it no longer + // matches — so it stays heap-charged until the repack cache lands (#1192 follow-up). + val tq = PlanTensor("w", null, Format(FP32, TensorEncoding.TQ2_0), bytes * 4, bytes, mappedForm) + val plan = MemoryPlans.plan(input(listOf(tq)), Budget.of(256L * 1024 * 1024), platform) - assertEquals(0L, plan.weightsMappedBytes, "Q8_0 has no mapped-servable representation (#1189 covers Q4_K/Q6_K)") + assertEquals(0L, plan.weightsMappedBytes, "TQ2_0 has no mapped-servable representation") // and the resolver agrees — the plan and the load must tell the same story - val spec = AllocationResolver.resolve(q80, PlannerProfile.MOBILE_2GB, platform) + val spec = AllocationResolver.resolve(tq, PlannerProfile.MOBILE_2GB, platform) assertNotEquals(MemoryDomain.MMAP_FILE, spec.domain) - assertTrue("no loader/kernel serves Q8_0" in AllocationResolver.explain(q80, PlannerProfile.MOBILE_2GB, platform)) + assertTrue("no loader/kernel serves TQ2_0" in AllocationResolver.explain(tq, PlannerProfile.MOBILE_2GB, platform)) } @Test fun a_mix_splits_into_a_mapped_and_a_heap_line() { val mappedForm = WeightForm(residency = WeightResidency.MAPPED) val big = q4k("big", 512L * 1024 * 1024, mappedForm) - val q80 = PlanTensor( - "small", null, Format(FP32, TensorEncoding.Q8_0), - (10L * 1024 * 1024) / 34 * 32, 10L * 1024 * 1024, mappedForm, + val tq = PlanTensor( + "small", null, Format(FP32, TensorEncoding.TQ2_0), + (10L * 1024 * 1024) * 4, 10L * 1024 * 1024, mappedForm, ) - val plan = MemoryPlans.plan(input(listOf(big, q80)), Budget.of(256L * 1024 * 1024), platform) + val plan = MemoryPlans.plan(input(listOf(big, tq)), Budget.of(256L * 1024 * 1024), platform) assertEquals(512L * 1024 * 1024, plan.weightsMappedBytes) assertEquals(10L * 1024 * 1024, plan.weightsHeapBytes) diff --git a/skainet-lang/skainet-lang-core/src/jvmAndroidMain/kotlin/sk/ainet/lang/tensor/data/BufferPackedTensorData.kt b/skainet-lang/skainet-lang-core/src/jvmAndroidMain/kotlin/sk/ainet/lang/tensor/data/BufferPackedTensorData.kt index df332a72a..a0248e53a 100644 --- a/skainet-lang/skainet-lang-core/src/jvmAndroidMain/kotlin/sk/ainet/lang/tensor/data/BufferPackedTensorData.kt +++ b/skainet-lang/skainet-lang-core/src/jvmAndroidMain/kotlin/sk/ainet/lang/tensor/data/BufferPackedTensorData.kt @@ -84,9 +84,49 @@ public class BufferPackedTensorData( scratchDecoder = d scratchData = d } + TensorEncoding.Q5_K -> { + blockSize = TensorEncoding.Q5_K.BLOCK_SIZE + bytesPerBlock = TensorEncoding.Q5_K.BYTES_PER_BLOCK + scratch = ByteArray(bytesPerBlock) + val d = Q5_KBlockTensorData(Shape(blockSize), scratch) + scratchDecoder = d + scratchData = d + } + TensorEncoding.Q8_0 -> { + blockSize = TensorEncoding.Q8_0.BLOCK_SIZE + bytesPerBlock = TensorEncoding.Q8_0.BYTES_PER_BLOCK + scratch = ByteArray(bytesPerBlock) + val d = Q8_0BlockTensorData(Shape(blockSize), scratch) + scratchDecoder = d + scratchData = d + } + TensorEncoding.Q4_0 -> { + blockSize = TensorEncoding.Q4_0.BLOCK_SIZE + bytesPerBlock = TensorEncoding.Q4_0.BYTES_PER_BLOCK + scratch = ByteArray(bytesPerBlock) + val d = Q4_0BlockTensorData(Shape(blockSize), scratch) + scratchDecoder = d + scratchData = d + } + TensorEncoding.Q5_0 -> { + blockSize = TensorEncoding.Q5_0.BLOCK_SIZE + bytesPerBlock = TensorEncoding.Q5_0.BYTES_PER_BLOCK + scratch = ByteArray(bytesPerBlock) + val d = Q5_0BlockTensorData(Shape(blockSize), scratch) + scratchDecoder = d + scratchData = d + } + TensorEncoding.Q5_1 -> { + blockSize = TensorEncoding.Q5_1.BLOCK_SIZE + bytesPerBlock = TensorEncoding.Q5_1.BYTES_PER_BLOCK + scratch = ByteArray(bytesPerBlock) + val d = Q5_1BlockTensorData(Shape(blockSize), scratch) + scratchDecoder = d + scratchData = d + } else -> throw IllegalArgumentException( - "BufferPackedTensorData supports Q4_K and Q6_K (#1189); got ${encoding.name}. " + - "Other block formats keep their heap TensorData until a buffer kernel exists." + "BufferPackedTensorData supports the GGML block formats " + + "(Q4_K/Q6_K, #1189; Q5_K/Q8_0/Q4_0/Q5_0/Q5_1, #1192); got ${encoding.name}." ) } require(shape.volume % blockSize == 0) { From b4a4577e36c9604d0cbde745ca38e5f5a1c258e9 Mon Sep 17 00:00:00 2001 From: Michal Harakal Date: Thu, 27 Aug 2026 23:04:46 +0200 Subject: [PATCH 2/3] =?UTF-8?q?test(#1192):=20cross-order=20parity=20toler?= =?UTF-8?q?ance=201e-4=20=E2=80=94=20MSVC=20x86=20auto-vectorization=20dri?= =?UTF-8?q?fts=20past=201e-5?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit q80_rm_parity_threaded failed on windows-x86_64: the same shared block-term inlined into two loop shapes auto-vectorizes with different partial-sum splits under MSVC, exceeding 1e-5 relative. 1e-4 remains six orders below any addressing bug (which produces O(1) garbage); threaded-vs-solo comparisons stay bit-exact. Co-authored-by: Claude Fable 5 --- .../kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt index 9d01dc615..13613d3b9 100644 --- a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt +++ b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/RowMajorMatmulParityTest.kt @@ -119,7 +119,7 @@ class RowMajorMatmulParityTest { * Threaded-vs-solo comparisons stay bit-exact — same function, same per-row order. */ private fun assertClose(expected: Float, got: Float, label: String) { - val tol = maxOf(1e-5f, 1e-5f * kotlin.math.abs(expected)) + val tol = maxOf(1e-4f, 1e-4f * kotlin.math.abs(expected)) assertTrue(kotlin.math.abs(expected - got) <= tol, "$label: $expected vs $got (tol $tol)") } From 75f0d4d4b953265f9a359d0563fda596688f3478 Mon Sep 17 00:00:00 2001 From: Michal Harakal Date: Thu, 27 Aug 2026 23:06:22 +0200 Subject: [PATCH 3/3] =?UTF-8?q?docs(#1166):=200.50.0=20changelog=20?= =?UTF-8?q?=E2=80=94=20mapped=20staging,=20threaded=20kernels,=20truthful?= =?UTF-8?q?=20plan,=20third-party=20notice?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The release entry for the Android memory arc, including the explicit Third-party section for the vendored NeoGPU kernel (#1166's last in-repo deliverable — 0.49.0's entry mentioned it inline; this makes the attribution a named section with upstream commit and license). Co-authored-by: Claude Fable 5 --- CHANGELOG.md | 75 ++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 75 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 78fdfca62..7d10c4afd 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -2,6 +2,81 @@ ## [Unreleased] +## [0.50.0] - 2026-08-28 + +Headline: **model size on Android is now a page-cache question, not a heap question — and decode is +2.5× faster.** Quantized weights are served straight from the memory-mapped GGUF file (zero copies, +zero relayout), the packed matmul kernels thread across cores, and a 1.0 GB Qwen2.5-1.5B Q4_K_M — +which OOM'd at load under the 256 MB ART cap in 0.49.0 — loads in ~0.4 s with 566 KB of weight heap +and decodes at 61–66 ms/step with zero steady-state page faults (Pixel 8a, measured by the M2-A5 +harness). The same file-order kernels serve heap-staged weights too, closing a measured 48,771 ms/step +silent-fallback trap for mixed-quant models — and fallbacks are never silent again. + +### Added + +- **Packed-tensor mapped staging** ([#1189](https://github.com/SKaiNET-developers/SKaiNET/issues/1189), + [#1190](https://github.com/SKaiNET-developers/SKaiNET/pull/1190)): under + `WeightForm(residency = MAPPED)` every GGML block format (Q4_K, Q6_K, Q5_K, Q8_0, Q4_0, Q5_0, Q5_1 — + [#1192](https://github.com/SKaiNET-developers/SKaiNET/issues/1192)) is served from file-backed pages: + `BufferPackedTensorData` over `MappedBufferStorage`/`DirectBufferStorage`, row-major (`_rm`) C kernels + that read canonical GGUF file order (no prepack, no relayout copy), reached via JNI direct-buffer + entries on Android and FFM `MemorySegment.ofBuffer` on the JVM + ([#1191](https://github.com/SKaiNET-developers/SKaiNET/issues/1191)). Ternary (`BITNET_B1_58`) still + heap-stages — its load-time repack needs the sidecar cache tracked in + [#1198](https://github.com/SKaiNET-developers/SKaiNET/issues/1198). +- **Threaded packed matmuls** ([#1195](https://github.com/SKaiNET-developers/SKaiNET/issues/1195), + [#1196](https://github.com/SKaiNET-developers/SKaiNET/pull/1196)): a spin-then-park worker pool with + guided row-grains, engaged at `outputDim ≥ 512`. The design is measurement-driven and the failed + variants are documented in the PR: per-call `pthread_create` cost 994 ms/step against deep-idle + cores, and every *sleeping* pool lost to one pegged big core because sub-millisecond bursts never + build scheduler utilization — the ~1 ms spin before parking (the same trick llama.cpp uses) is what + unlocks big cores at full clocks. 153 → 61 ms/step on the 1.5B; results are bit-identical to + single-threaded (disjoint row ranges, unchanged per-row accumulation order). +- **Storage-polymorphic row-major dispatch, and no silent fallbacks** + ([#1193](https://github.com/SKaiNET-developers/SKaiNET/issues/1193), + [#1197](https://github.com/SKaiNET-developers/SKaiNET/pull/1197)): one kernel per format serves + `BLOCKED_ROW_MAJOR` weights from mapped, direct **or heap** storage — an un-prepacked heap canonical + weight used to fall to the decoding reference kernel silently (measured: 48,771 ms/step on + SmolLM2-135M, whose non-256-multiple dims made llama.cpp's quantizer emit mostly Q8_0; now 33 ms/step, + the fastest configuration measured). `ViewKernel` gained a sink-aware `run` overload and every packed + bridge announces a reference fallback as a `KernelRun` trace event with the reason; the M2-A5 harness + prints the count. +- **The plan tells the truth about mapped weights** + ([#1190](https://github.com/SKaiNET-developers/SKaiNET/pull/1190)): `MemoryPlan.budgetedBytes` + charges mapped-servable weights against device RAM/page cache instead of the heap budget + (`fits`, suggestions and `PlannerProfile`'s KV auto-quantization follow), rendered as its own + `mapped (page cache, evictable — not heap)` line. `AllocationResolver.servesFromMapping` is the one + predicate the resolver, the plan and the loader share, gated by + `StorageCapabilities.mappedServableEncodings`, so they cannot tell different stories; `planInput` + takes the `WeightForm` the load will use. +- **Kernel-support matrix: mapped serving section** — the generated matrix now shows, per platform, + which formats serve from a mapping (all seven on Android `native-jni-direct` and JVM `ffm-rowmajor`; + empty cells are the documented gaps). +- **"The DSL is compute"** ([#1194](https://github.com/SKaiNET-developers/SKaiNET/pull/1194)): + the architecture principle behind all of the above as a teachable explanation page — network + definitions describe computation only; memory intent lives in `WeightForm` at the load boundary, + priced by the plan, honored-or-visibly-rejected by resolvers, with runtime dispatch holding no + memory policy. + +### Changed + +- Cross-order kernel results (row-major vs feed-order) are **numerically equivalent, not bit-exact**: + under `-O3 -ffast-math` compilers may contract float accumulation differently per loop shape + (measured 2 ULP on clang/arm64, more under MSVC auto-vectorization). Integer-dot formats (Q4_K, Q6_K) + currently match exactly, but only threaded-vs-single-threaded identity is contractual. The parity + suites encode this. +- The M2-A5 measurement harness gained `residency=heap|mapped` and reports mapped vs heap weight bytes + and the packed-kernel fallback count. + +### Third-party + +- The vendored NeoGPU ternary NEON kernel + (`skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c`, + © 2024 NeoGPU Contributors, MIT, byte-identical to upstream + [anjaustin/neogpu](https://github.com/anjaustin/neogpu) @ `0846b24`) ships unchanged in this release; + REUSE metadata and `META-INF/THIRD-PARTY-NOTICES.md` in the published artifacts carry the attribution + ([#1166](https://github.com/SKaiNET-developers/SKaiNET/issues/1166)). + ## [0.49.0] - 2026-08-26 Headline: **the SKEEP-003 memory & storage architecture, complete — from accepted proposal to shipped system.**