From 4815fc6169bc9145797694c4a3ff103ef5d0874f Mon Sep 17 00:00:00 2001 From: Michal Harakal Date: Wed, 26 Aug 2026 10:04:09 +0200 Subject: [PATCH 1/2] feat(native): vendor the NeoGPU ternary LUT kernel, verbatim under MIT MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Byte-identical copy of src/hs_ml_ternary_neon.c from anjaustin/neogpu at 0846b24 (sha256 recorded in the vendor README), with the REUSE .license sidecar carrying the upstream copyright. Exact f32-activation ternary matmul on baseline NEON — no dotprod, the fast path for Cortex-A72/Pi-4. Agreed with upstream in anjaustin/neogpu#1. Refs #1137, #1136 Co-Authored-By: Claude Fable 5 --- .../native/src/vendor/neogpu/README.md | 53 +++ .../src/vendor/neogpu/hs_ml_ternary_neon.c | 338 ++++++++++++++++++ .../neogpu/hs_ml_ternary_neon.c.license | 2 + 3 files changed, 393 insertions(+) create mode 100644 skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/README.md create mode 100644 skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c create mode 100644 skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c.license diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/README.md b/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/README.md new file mode 100644 index 000000000..457b70aaa --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/README.md @@ -0,0 +1,53 @@ +# Vendored: NeoGPU ternary LUT kernel + +Byte-identical copy of one file from the NeoGPU project, vendored under its +MIT license (see the REUSE `.license` sidecar; `LICENSES/MIT.txt` at the repo +root carries the license text). Agreed with upstream in +[anjaustin/neogpu#1](https://github.com/anjaustin/neogpu/issues/1); SKaiNET +tracking issue: [#1136](https://github.com/SKaiNET-developers/SKaiNET/issues/1136). + +| | | +|---|---| +| Upstream | | +| File | `src/hs_ml_ternary_neon.c` | +| Vendored at commit | `0846b24ceb9f76a610b3efe2967ff4ead2ef10e6` | +| SHA-256 | `a560ffcf4d5a2e2f600d715b8193da999a92982160c625e5da48076d2257d587` | +| Local modifications | **none** | + +## What it is + +An f32-activation × ternary-weight ({-1, 0, +1}) matmul for AArch64 using a +256-entry × 4-float decode LUT (4 KB, L1-resident) — `vld1q_f32(lut[w[b]])` + +`vfmaq_f32` inner loop. Needs only baseline NEON (`armv8-a+simd`), no +`FEAT_DotProd`: it is the fast path for Cortex-A72 / Raspberry Pi 4-class +cores. Carries its own exact scalar fallback for non-ARM builds. The 2-bit +payload rule (4 codes per byte, low bit-pair first, code = value + 1) is +byte-identical to SKaiNET's `BITNET_B1_58` / `TernaryPacked` encoding. + +Exports three symbols; SKaiNET wraps the first two via +`src/skainet_ternary_f32.c` and never calls the third +(`hs_ml_lmhead_stage1_i8` uses a thread-unsafe static scratch buffer): + +- `hs_ml_ternary_f32_proj` — single-plane projection GEMV +- `hs_ml_lmhead_stage1` — fused 4-plane lm_head with FP16 row scales +- `hs_ml_lmhead_stage1_i8` — ABI-compat stub, **do not use** + +## Constraints the adapter enforces / works around + +- `build_lut()` uses a non-atomic init guard → the adapter warms it up under + `pthread_once` before any concurrent use. +- pthreads: threading is NeoGPU's own (4 threads once `N >= 512` for the proj; + the lm_head always threads). No MSVC build — the adapter compiles a portable + scalar fallback there instead (`SKAINET_HAVE_NEOGPU_TERNARY` unset). +- Byte code 3 decodes to +2.0 (matches `TernaryCodec.decodeBitNet`); loaders + must reject code 3 at import, the kernel never validates. +- Must be compiled at `-march=armv8-a` on non-Apple AArch64 — the library + default `-march=armv8.2-a+fp16+dotprod` would defeat its purpose + (dotprod-less targets). + +## Re-vendoring + +Copy the file byte-identical from upstream, update the commit + SHA-256 here, +keep "Local modifications: none" true (fixes belong in the adapter or +upstream), and re-run the ternary goldens in +`src/jvmTest/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernelTest.kt`. diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c b/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c new file mode 100644 index 000000000..967ca3331 --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c @@ -0,0 +1,338 @@ +/* + * NeoGPU ML — Ternary Projection Kernel (NEON + pthreads, Cortex-A72) + * + * Pure ternary arithmetic: weights in {-1, 0, +1}, activations in float32. + * No integer quantization of activations. No approximate path. + * + * I2_S packing: byte bi holds 4 codes for act[bi*4 .. bi*4+3] + * bits[1:0] -> code for act[bi*4+0] + * bits[3:2] -> code for act[bi*4+1] + * bits[5:4] -> code for act[bi*4+2] + * bits[7:6] -> code for act[bi*4+3] + * Code mapping: 0->-1, 1->0, 2->+1 + * + * Decode strategy: 256-entry LUT + * lut[byte] = {float(-1/0/+1), float(-1/0/+1), float(-1/0/+1), float(-1/0/+1)} + * Inner loop: vld1q_f32(lut[w[bi]]) + vld1q_f32(in+bi*4) + vfmaq_f32 + * 4KB table fits in L1 cache. No scalar decode arithmetic. + * Measured: 96ms / 6.78 GOPS (4-thread, full vocab) + * vs int8 path: 51ms / 3.20 GOPS (has 1.5% quantization error) + * vs F16 NEON: 175ms / 0.93 GOPS + * + * Two public functions: + * hs_ml_ternary_f32_proj -- single plane, N rows (projection layers) + * hs_ml_lmhead_stage1 -- fused 4 planes, all V rows (lm_head Stage 1) + */ + +#include +#include +#include + +#ifdef __ARM_NEON +#include +#endif + +#define N_THREADS 4 +#define THREAD_THRESHOLD 512 + +/* ============================================================ + * 256-entry decode LUT + * lut[b][lane] = (float)((b >> (lane*2)) & 3) - 1 ∈ {-1, 0, +1} + * 256 * 4 * 4 bytes = 4096 bytes = 4KB + * Aligned to 64 bytes (cache line) for prefetch efficiency. + * ============================================================ */ + +static float g_lut[256][4] __attribute__((aligned(64))); +static int g_lut_built = 0; + +static void build_lut(void) { + if (g_lut_built) return; + for (int b = 0; b < 256; b++) { + g_lut[b][0] = (float)((int)((b >> 0) & 3) - 1); + g_lut[b][1] = (float)((int)((b >> 2) & 3) - 1); + g_lut[b][2] = (float)((int)((b >> 4) & 3) - 1); + g_lut[b][3] = (float)((int)((b >> 6) & 3) - 1); + } + g_lut_built = 1; +} + +/* ============================================================ + * Single-plane inner kernel + * out[n] = dot(in[0..K-1], decode(W[n])) + * LUT inner loop: 2 loads + 1 FMA per byte = minimal decode overhead + * ============================================================ */ + +static void proj_rows(float *out, const float *in, + const uint8_t *W, + uint32_t row_start, uint32_t row_end, + uint32_t K) { + const uint32_t row_bytes = K / 4; + +#ifdef __ARM_NEON + for (uint32_t n = row_start; n < row_end; n++) { + const uint8_t *wrow = W + (size_t)n * row_bytes; + + float32x4_t acc0 = vdupq_n_f32(0.0f); + float32x4_t acc1 = vdupq_n_f32(0.0f); + float32x4_t acc2 = vdupq_n_f32(0.0f); + float32x4_t acc3 = vdupq_n_f32(0.0f); + + uint32_t b4 = row_bytes & ~3u; + for (uint32_t bi = 0; bi < b4; bi += 4) { + __builtin_prefetch(wrow + bi + 64, 0, 3); + __builtin_prefetch(in + (bi + 8) * 4, 0, 3); + + acc0 = vfmaq_f32(acc0, + vld1q_f32(g_lut[wrow[bi ]]), + vld1q_f32(in + (bi ) * 4)); + acc1 = vfmaq_f32(acc1, + vld1q_f32(g_lut[wrow[bi+1]]), + vld1q_f32(in + (bi+1) * 4)); + acc2 = vfmaq_f32(acc2, + vld1q_f32(g_lut[wrow[bi+2]]), + vld1q_f32(in + (bi+2) * 4)); + acc3 = vfmaq_f32(acc3, + vld1q_f32(g_lut[wrow[bi+3]]), + vld1q_f32(in + (bi+3) * 4)); + } + /* Tail */ + for (uint32_t bi = b4; bi < row_bytes; bi++) { + acc0 = vfmaq_f32(acc0, + vld1q_f32(g_lut[wrow[bi]]), + vld1q_f32(in + bi * 4)); + } + + out[n] = vaddvq_f32(vaddq_f32(vaddq_f32(acc0, acc1), + vaddq_f32(acc2, acc3))); + } +#else + for (uint32_t n = row_start; n < row_end; n++) { + const uint8_t *wrow = W + (size_t)n * row_bytes; + float acc = 0.0f; + for (uint32_t bi = 0; bi < row_bytes; bi++) { + uint8_t b = wrow[bi]; + acc += g_lut[b][0] * in[bi*4 ]; + acc += g_lut[b][1] * in[bi*4+1]; + acc += g_lut[b][2] * in[bi*4+2]; + acc += g_lut[b][3] * in[bi*4+3]; + } + out[n] = acc; + } +#endif +} + +/* ============================================================ + * Fused 4-plane lm_head Stage 1 inner kernel + * + * Single pass over N rows, reading 4 I2_S planes simultaneously. + * Plane weights: w0=1, w1=1/3, w2=1/9, w3=1/27. + * Applies F16 row_scale inline. + * + * out[n] = row_scale[n] * sum_{k=0}^{3} w_k * dot(in, plane_k[n]) + * + * Per byte-position: 4 LUT lookups (one per plane) + 4 FMAs. + * Same 4KB LUT — no extra memory needed. + * ============================================================ */ + +static void lmhead_stage1_rows(float *out, const float *in, + const uint8_t *P0, const uint8_t *P1, + const uint8_t *P2, const uint8_t *P3, + const uint16_t *row_scale, + uint32_t row_start, uint32_t row_end, + uint32_t K) { + const uint32_t row_bytes = K / 4; + const float W1 = 1.0f / 3.0f; + const float W2 = 1.0f / 9.0f; + const float W3 = 1.0f / 27.0f; + +#ifdef __ARM_NEON + for (uint32_t n = row_start; n < row_end; n++) { + const uint8_t *r0 = P0 + (size_t)n * row_bytes; + const uint8_t *r1 = P1 + (size_t)n * row_bytes; + const uint8_t *r2 = P2 + (size_t)n * row_bytes; + const uint8_t *r3 = P3 + (size_t)n * row_bytes; + + float32x4_t acc0 = vdupq_n_f32(0.0f); + float32x4_t acc1 = vdupq_n_f32(0.0f); + float32x4_t acc2 = vdupq_n_f32(0.0f); + float32x4_t acc3 = vdupq_n_f32(0.0f); + + uint32_t b4 = row_bytes & ~3u; + for (uint32_t bi = 0; bi < b4; bi += 4) { + __builtin_prefetch(r0 + bi + 64, 0, 1); + __builtin_prefetch(r1 + bi + 64, 0, 1); + __builtin_prefetch(r2 + bi + 64, 0, 1); + __builtin_prefetch(r3 + bi + 64, 0, 1); + __builtin_prefetch(in + (bi + 8) * 4, 0, 3); + + for (int s = 0; s < 4; s++) { + float32x4_t act = vld1q_f32(in + (bi + s) * 4); + + /* Combine 4 plane contributions into one weighted vector */ + float32x4_t v = vld1q_f32(g_lut[r0[bi+s]]); + float32x4_t u = vfmaq_n_f32(v, vld1q_f32(g_lut[r1[bi+s]]), W1); + u = vfmaq_n_f32(u, vld1q_f32(g_lut[r2[bi+s]]), W2); + u = vfmaq_n_f32(u, vld1q_f32(g_lut[r3[bi+s]]), W3); + + switch (s) { + case 0: acc0 = vfmaq_f32(acc0, u, act); break; + case 1: acc1 = vfmaq_f32(acc1, u, act); break; + case 2: acc2 = vfmaq_f32(acc2, u, act); break; + case 3: acc3 = vfmaq_f32(acc3, u, act); break; + } + } + } + /* Tail */ + for (uint32_t bi = b4; bi < row_bytes; bi++) { + float32x4_t act = vld1q_f32(in + bi * 4); + float32x4_t v = vld1q_f32(g_lut[r0[bi]]); + float32x4_t u = vfmaq_n_f32(v, vld1q_f32(g_lut[r1[bi]]), W1); + u = vfmaq_n_f32(u, vld1q_f32(g_lut[r2[bi]]), W2); + u = vfmaq_n_f32(u, vld1q_f32(g_lut[r3[bi]]), W3); + acc0 = vfmaq_f32(acc0, u, act); + } + + float s = vaddvq_f32(vaddq_f32(vaddq_f32(acc0, acc1), + vaddq_f32(acc2, acc3))); + + /* Inline F16 row_scale decode */ + uint32_t h16 = row_scale[n]; + uint32_t eu = (h16 >> 10) & 0x1Fu; + uint32_t mu = h16 & 0x03FFu; + uint32_t fu = ((eu + 112u) << 23) | (mu << 13); + float rsc; __builtin_memcpy(&rsc, &fu, 4); + out[n] = s * rsc; + } +#else + for (uint32_t n = row_start; n < row_end; n++) { + const uint8_t *r0=P0+(size_t)n*row_bytes, *r1=P1+(size_t)n*row_bytes; + const uint8_t *r2=P2+(size_t)n*row_bytes, *r3=P3+(size_t)n*row_bytes; + float acc = 0.0f; + for (uint32_t bi = 0; bi < row_bytes; bi++) { + for (int s = 0; s < 4; s++) { + float act = in[bi*4+s]; + acc += (g_lut[r0[bi]][s] + + g_lut[r1[bi]][s] * W1 + + g_lut[r2[bi]][s] * W2 + + g_lut[r3[bi]][s] * W3) * act; + } + } + uint32_t h16=row_scale[n]; + uint32_t eu=(h16>>10)&0x1Fu, mu=h16&0x03FFu; + uint32_t fu=((eu+112u)<<23)|(mu<<13); + float rsc; __builtin_memcpy(&rsc,&fu,4); + out[n] = acc * rsc; + } +#endif +} + +/* ============================================================ + * Thread pool + * ============================================================ */ + +typedef struct { + float *out; + const float *in; + const uint8_t *W; + const uint8_t *P0, *P1, *P2, *P3; + const uint16_t *row_scale; + uint32_t row_start, row_end, K; + int fused; +} WorkItem; + +static void *thread_fn(void *arg) { + WorkItem *w = (WorkItem *)arg; + if (w->fused) + lmhead_stage1_rows(w->out, w->in, + w->P0, w->P1, w->P2, w->P3, w->row_scale, + w->row_start, w->row_end, w->K); + else + proj_rows(w->out, w->in, w->W, w->row_start, w->row_end, w->K); + return NULL; +} + +static void dispatch(WorkItem *work, int n_active) { + pthread_t threads[N_THREADS]; + for (int t = 0; t < n_active; t++) + pthread_create(&threads[t], NULL, thread_fn, &work[t]); + for (int t = 0; t < n_active; t++) + pthread_join(threads[t], NULL); +} + +/* ============================================================ + * Public API — single-plane projection + * out[n] = dot(in, W[n]) for n in [0, N) + * ============================================================ */ + +void hs_ml_ternary_f32_proj(float *out, const float *in, + const uint8_t *W, uint32_t N, uint32_t K) { + build_lut(); + + if (N < THREAD_THRESHOLD) { + proj_rows(out, in, W, 0, N, K); + return; + } + + WorkItem work[N_THREADS]; + uint32_t chunk = (N + N_THREADS - 1) / N_THREADS; + int active = 0; + + for (int t = 0; t < N_THREADS; t++) { + uint32_t start = (uint32_t)t * chunk; + uint32_t end = start + chunk; + if (start >= N) break; + if (end > N) end = N; + work[t] = (WorkItem){ out, in, W, NULL, NULL, NULL, NULL, NULL, + start, end, K, 0 }; + active++; + } + dispatch(work, active); +} + +/* ============================================================ + * Public API — fused 4-plane lm_head Stage 1 (LUT, float32) + * out[n] = row_scale[n] * (dot(in,P0[n]) + dot(in,P1[n])/3 + * + dot(in,P2[n])/9 + dot(in,P3[n])/27) + * ============================================================ */ + +void hs_ml_lmhead_stage1(float *out, const float *in, + const uint8_t *P0, const uint8_t *P1, + const uint8_t *P2, const uint8_t *P3, + const uint16_t *row_scale, + uint32_t N, uint32_t K) { + build_lut(); + + WorkItem work[N_THREADS]; + uint32_t chunk = (N + N_THREADS - 1) / N_THREADS; + int active = 0; + + for (int t = 0; t < N_THREADS; t++) { + uint32_t start = (uint32_t)t * chunk; + uint32_t end = start + chunk; + if (start >= N) break; + if (end > N) end = N; + work[t] = (WorkItem){ out, in, NULL, P0, P1, P2, P3, row_scale, + start, end, K, 1 }; + active++; + } + dispatch(work, active); +} + +/* ============================================================ + * Public API — int8 Stage 1 (kept for ABI compatibility, delegates to float) + * The int8 path is no longer the preferred path; use hs_ml_lmhead_stage1. + * This stub dequantizes in_i8 back to float and calls the LUT kernel. + * ============================================================ */ + +void hs_ml_lmhead_stage1_i8(float *out, const int8_t *in_i8, float act_scale, + const uint8_t *P0, const uint8_t *P1, + const uint8_t *P2, const uint8_t *P3, + const uint16_t *row_scale, + uint32_t N, uint32_t K) { + /* Dequantize to float, then use LUT path */ + static float scratch[4096]; /* H=2560, stack-safe */ + float inv = 1.0f / act_scale; + for (uint32_t i = 0; i < K * 4; i++) + scratch[i] = (float)in_i8[i] * inv; + hs_ml_lmhead_stage1(out, scratch, P0, P1, P2, P3, row_scale, N, K); +} diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c.license b/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c.license new file mode 100644 index 000000000..460bbca9b --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/native/src/vendor/neogpu/hs_ml_ternary_neon.c.license @@ -0,0 +1,2 @@ +SPDX-FileCopyrightText: 2024 NeoGPU Contributors +SPDX-License-Identifier: MIT From 814b2fd73ef436996f2374f14a1ecc6388f84017 Mon Sep 17 00:00:00 2001 From: Michal Harakal Date: Wed, 26 Aug 2026 10:04:09 +0200 Subject: [PATCH 2/2] feat(native): expose the vendored ternary f32 kernel through the FFM pipeline MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit skainet_ternary_f32.c adapts the vendored file to the skainet_kernels ABI: skainet_ternary_f32_gemv (single-plane proj, no scale applied — the caller owns it) and skainet_ternary_lmhead_stage1 (fused 4-plane + FP16 row scales, stride-addressed so one concatenated buffer works). A pthread_once warm-up closes the vendored non-atomic LUT-init guard; MSVC builds compile a portable scalar mirror instead (SKAINET_HAVE_NEOGPU_TERNARY unset). Build: both native trees compile the two files; on non-Apple aarch64 they are pinned back to -march=armv8-a — the library-wide v8.2+dotprod flags would defeat a kernel that exists for dotprod-less cores. NativeTernaryF32GemvKernel downcalls the gemv via FFM and round-trips the output reach so row-looping callers keep previously written rows. Goldens pin the full 256-byte decode table (code 3 → +2), the pthread partitioning above the 512-row threshold, offsets, and edge cases — the same goldens cover the NEON branch on arm64 and the scalar branch on x86. The TernaryF32GemvNative SPI + KernelDispatch wiring follows in #1138; the kernel-support matrix is untouched until the kernel joins a provider. Refs #1137, #1136 Co-Authored-By: Claude Fable 5 --- .../native/CMakeLists.txt | 5 + .../native/CMakeLists.txt | 25 +++ .../native/include/skainet_kernels.h | 48 +++++ .../native/src/skainet_ternary_f32.c | 152 +++++++++++++++ .../exec/kernel/NativeTernaryF32GemvKernel.kt | 113 ++++++++++++ .../kernel/NativeTernaryF32GemvKernelTest.kt | 173 ++++++++++++++++++ 6 files changed, 516 insertions(+) create mode 100644 skainet-backends/skainet-backend-native-cpu/native/src/skainet_ternary_f32.c create mode 100644 skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernel.kt create mode 100644 skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernelTest.kt diff --git a/skainet-backends/skainet-backend-jni-cpu/native/CMakeLists.txt b/skainet-backends/skainet-backend-jni-cpu/native/CMakeLists.txt index 05a8eef4d..bcb2be72f 100644 --- a/skainet-backends/skainet-backend-jni-cpu/native/CMakeLists.txt +++ b/skainet-backends/skainet-backend-jni-cpu/native/CMakeLists.txt @@ -18,6 +18,10 @@ set(SKAINET_KERNEL_SOURCES ${SKAINET_KERNELS_ROOT}/src/q5k_matmul.c ${SKAINET_KERNELS_ROOT}/src/q6k_matmul.c ${SKAINET_KERNELS_ROOT}/src/bitnet_gemv.c + ${SKAINET_KERNELS_ROOT}/src/skainet_ternary_f32.c + # Vendored NeoGPU ternary LUT kernel (MIT, verbatim). NDK toolchains are + # clang, and bionic bundles pthreads — no Threads link needed. + ${SKAINET_KERNELS_ROOT}/src/vendor/neogpu/hs_ml_ternary_neon.c ) set(SKAINET_JNI_SHIM ${CMAKE_CURRENT_SOURCE_DIR}/skainet_jni.c) @@ -34,6 +38,7 @@ function(skainet_add_jni_lib target) add_library(${target} SHARED ${SKAINET_JNI_SHIM} ${SKAINET_KERNEL_SOURCES}) target_include_directories(${target} PRIVATE ${SKAINET_KERNELS_ROOT}/include) target_compile_options(${target} PRIVATE ${SKAINET_C_FLAGS}) + target_compile_definitions(${target} PRIVATE SKAINET_HAVE_NEOGPU_TERNARY) target_link_options(${target} PRIVATE ${SKAINET_LINK_FLAGS}) endfunction() diff --git a/skainet-backends/skainet-backend-native-cpu/native/CMakeLists.txt b/skainet-backends/skainet-backend-native-cpu/native/CMakeLists.txt index 5525e0a4f..e97b9d0b0 100644 --- a/skainet-backends/skainet-backend-native-cpu/native/CMakeLists.txt +++ b/skainet-backends/skainet-backend-native-cpu/native/CMakeLists.txt @@ -23,8 +23,17 @@ set(SKAINET_KERNEL_SOURCES src/q4_0_matmul.c src/q5_0_matmul.c src/q5_1_matmul.c + src/skainet_ternary_f32.c ) +# Vendored NeoGPU ternary LUT kernel (MIT, verbatim — src/vendor/neogpu/). +# Needs pthreads + GNU/Clang builtins, so MSVC builds get the portable +# scalar fallback inside skainet_ternary_f32.c instead. +if(NOT MSVC) + list(APPEND SKAINET_KERNEL_SOURCES src/vendor/neogpu/hs_ml_ternary_neon.c) + find_package(Threads REQUIRED) +endif() + # SHARED: consumed by the JVM via java.lang.foreign (FFM), bundled as a JAR # resource (libskainet_kernels.{so,dylib,dll}). # STATIC: consumed by Kotlin/Native via cinterop, linked into the K/N binary @@ -49,6 +58,12 @@ endif() foreach(tgt IN LISTS SKAINET_KERNEL_TARGETS) target_include_directories(${tgt} PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/include) + if(NOT MSVC) + # The adapter forwards to the vendored NeoGPU kernel (pthreads). + target_compile_definitions(${tgt} PRIVATE SKAINET_HAVE_NEOGPU_TERNARY) + target_link_libraries(${tgt} PRIVATE Threads::Threads) + endif() + # Strip the "lib" prefix on Windows so the shared artifact name matches # the resource-bundle path skainet_kernels.{dll,so,dylib}. if(WIN32) @@ -83,6 +98,16 @@ foreach(tgt IN LISTS SKAINET_KERNEL_TARGETS) target_compile_options(${tgt} PRIVATE -march=armv8.2-a+fp16+dotprod ) + # The vendored NeoGPU ternary LUT kernel exists FOR dotprod-less + # cores (Cortex-A72/Pi-4): pin it and its adapter back to the + # armv8-a baseline so the library-wide v8.2 flags above cannot + # emit instructions those cores lack. Source-level COMPILE_OPTIONS + # append after target options, so the later -march wins. + set_source_files_properties( + src/vendor/neogpu/hs_ml_ternary_neon.c + src/skainet_ternary_f32.c + PROPERTIES COMPILE_OPTIONS "-march=armv8-a" + ) endif() set_target_properties(${tgt} PROPERTIES C_VISIBILITY_PRESET hidden) elseif(CMAKE_C_COMPILER_ID MATCHES "MSVC") 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 e3f8f0897..43637691c 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 @@ -291,6 +291,54 @@ SKAINET_API void skainet_bitnet_gemv_tq2_0( int32_t input_dim, int32_t output_dim, float* output, int32_t output_offset); +/* + * Ternary f32 GEMV — exact FP32 activations against sequentially-packed + * ternary weights (the BitNet b1.58 / `BITNET_B1_58` payload: 4 codes per + * byte, low bit-pair first, code {0,1,2} → {-1,0,+1}; byte code 3 decodes + * to +2, loaders reject it at import). + * + * output[output_offset + o] = sum_j input[input_offset + j] * + * decode(weight row o) + * + * NO scale is applied — the caller owns the per-tensor scale. Weights are + * row-major, input_dim/4 bytes per output row, at + * weight + weight_byte_offset + o * (input_dim / 4) + * + * input_dim must be a multiple of 4. Unlike the int8 `bitnet_gemv` path + * there is no activation quantization: results are exact. Backed by the + * vendored NeoGPU LUT kernel (baseline NEON, no dotprod — the fast path for + * Cortex-A72/Pi-4 class cores); it threads internally with pthreads once + * output_dim >= 512. A portable scalar build stands in where the vendored + * file cannot compile (MSVC). + */ +SKAINET_API void skainet_ternary_f32_gemv( + 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); + +/* + * Fused 4-plane ternary lm_head, Stage 1 (NeoGPU multi-plane trit format). + * + * output[output_offset + o] = f16(row_scale[row_scale_offset + o]) * + * sum_{p=0}^{3} (1/3^p) * sum_j input[..+j] * decode(plane p, row o) + * + * Each plane is a full sequentially-packed ternary matrix (same payload rule + * as skainet_ternary_f32_gemv); plane p starts at + * planes + planes_byte_offset + p * plane_stride_bytes + * so a single buffer of four concatenated planes uses + * plane_stride_bytes == output_dim * input_dim / 4. row_scale holds raw + * little-endian IEEE binary16 bit patterns, one per output row (sign + * ignored — encoders store max|row| >= 0). input_dim must be a multiple + * of 4. Threads internally with pthreads at any output_dim. + */ +SKAINET_API void skainet_ternary_lmhead_stage1( + const float* input, int32_t input_offset, + const uint8_t* planes, int32_t planes_byte_offset, int32_t plane_stride_bytes, + const uint16_t* row_scale, int32_t row_scale_offset, + int32_t input_dim, int32_t output_dim, + float* output, int32_t output_offset); + #ifdef __cplusplus } #endif diff --git a/skainet-backends/skainet-backend-native-cpu/native/src/skainet_ternary_f32.c b/skainet-backends/skainet-backend-native-cpu/native/src/skainet_ternary_f32.c new file mode 100644 index 000000000..a145b5677 --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/native/src/skainet_ternary_f32.c @@ -0,0 +1,152 @@ +/* + * Adapter over the vendored NeoGPU ternary LUT kernel + * (src/vendor/neogpu/hs_ml_ternary_neon.c, MIT, vendored verbatim — see the + * vendor README). Exposes the two useful symbols under the skainet_kernels + * ABI and closes the vendored file's non-atomic LUT-init guard with a + * pthread_once warm-up, so first use is race-free under any threading. + * + * The vendored file needs pthreads and GNU/Clang builtins, so it is only in + * the build where SKAINET_HAVE_NEOGPU_TERNARY is defined (non-MSVC). The + * fallback branch below is a portable scalar mirror with identical semantics, + * including the byte-code-3 → +2.0 decode (loaders reject code 3 at import; + * the kernel contract is "garbage in, garbage out", never a crash). + */ + +#include "skainet_kernels.h" + +#include + +#ifdef SKAINET_HAVE_NEOGPU_TERNARY + +#include + +/* The vendored file ships no header; these mirror its public definitions. */ +extern void hs_ml_ternary_f32_proj(float *out, const float *in, + const uint8_t *W, uint32_t N, uint32_t K); +extern void hs_ml_lmhead_stage1(float *out, const float *in, + const uint8_t *P0, const uint8_t *P1, + const uint8_t *P2, const uint8_t *P3, + const uint16_t *row_scale, + uint32_t N, uint32_t K); + +static pthread_once_t skainet_ternary_lut_once = PTHREAD_ONCE_INIT; + +static void skainet_ternary_lut_warmup(void) { + /* Any call builds the 256-entry LUT; N=1 stays on the calling thread + * (far below the vendored THREAD_THRESHOLD of 512). */ + static const float in[4] = { 0.0f, 0.0f, 0.0f, 0.0f }; + static const uint8_t w[1] = { 0 }; + float out; + hs_ml_ternary_f32_proj(&out, in, w, 1u, 4u); +} + +void skainet_ternary_f32_gemv( + 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) +{ + if (output_dim <= 0) return; + pthread_once(&skainet_ternary_lut_once, skainet_ternary_lut_warmup); + /* input_dim == 0 falls through: the vendored kernel writes 0.0f per row. */ + hs_ml_ternary_f32_proj( + output + output_offset, + input + input_offset, + weight + weight_byte_offset, + (uint32_t)output_dim, + (uint32_t)input_dim); +} + +void skainet_ternary_lmhead_stage1( + const float* input, int32_t input_offset, + const uint8_t* planes, int32_t planes_byte_offset, int32_t plane_stride_bytes, + const uint16_t* row_scale, int32_t row_scale_offset, + int32_t input_dim, int32_t output_dim, + float* output, int32_t output_offset) +{ + if (output_dim <= 0) return; + pthread_once(&skainet_ternary_lut_once, skainet_ternary_lut_warmup); + const uint8_t* base = planes + planes_byte_offset; + hs_ml_lmhead_stage1( + output + output_offset, + input + input_offset, + base, + base + (size_t)plane_stride_bytes, + base + (size_t)plane_stride_bytes * 2, + base + (size_t)plane_stride_bytes * 3, + row_scale + row_scale_offset, + (uint32_t)output_dim, + (uint32_t)input_dim); +} + +#else /* !SKAINET_HAVE_NEOGPU_TERNARY — portable scalar mirror (MSVC etc.) */ + +/* decode(byte, lane) = ((byte >> (lane*2)) & 3) - 1, exactly the vendored LUT + * (so byte code 3 yields +2, matching TernaryCodec.decodeBitNet). */ +static float skainet_ternary_decode(uint8_t b, int lane) { + return (float)((int)((b >> (lane * 2)) & 3) - 1); +} + +void skainet_ternary_f32_gemv( + 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) +{ + if (output_dim <= 0) return; + const float* in = input + input_offset; + const uint8_t* w = weight + weight_byte_offset; + const int32_t row_bytes = input_dim / 4; + for (int32_t n = 0; n < output_dim; n++) { + const uint8_t* wrow = w + (size_t)n * row_bytes; + float acc = 0.0f; + for (int32_t bi = 0; bi < row_bytes; bi++) { + const uint8_t b = wrow[bi]; + acc += skainet_ternary_decode(b, 0) * in[bi * 4 ]; + acc += skainet_ternary_decode(b, 1) * in[bi * 4 + 1]; + acc += skainet_ternary_decode(b, 2) * in[bi * 4 + 2]; + acc += skainet_ternary_decode(b, 3) * in[bi * 4 + 3]; + } + output[output_offset + n] = acc; + } +} + +void skainet_ternary_lmhead_stage1( + const float* input, int32_t input_offset, + const uint8_t* planes, int32_t planes_byte_offset, int32_t plane_stride_bytes, + const uint16_t* row_scale, int32_t row_scale_offset, + int32_t input_dim, int32_t output_dim, + float* output, int32_t output_offset) +{ + if (output_dim <= 0) return; + const float* in = input + input_offset; + const uint8_t* base = planes + planes_byte_offset; + const uint16_t* rs = row_scale + row_scale_offset; + const int32_t row_bytes = input_dim / 4; + const float pw[4] = { 1.0f, 1.0f / 3.0f, 1.0f / 9.0f, 1.0f / 27.0f }; + for (int32_t n = 0; n < output_dim; n++) { + float acc = 0.0f; + for (int32_t bi = 0; bi < row_bytes; bi++) { + for (int lane = 0; lane < 4; lane++) { + float u = 0.0f; + for (int p = 0; p < 4; p++) { + const uint8_t b = + base[(size_t)plane_stride_bytes * p + (size_t)n * row_bytes + bi]; + u += skainet_ternary_decode(b, lane) * pw[p]; + } + acc += u * in[bi * 4 + lane]; + } + } + /* FP16 row scale decode, same bit arithmetic as the vendored file + * (sign ignored — encoders store max|row| >= 0). */ + const uint32_t h16 = rs[n]; + const uint32_t eu = (h16 >> 10) & 0x1Fu; + const uint32_t mu = h16 & 0x03FFu; + const uint32_t fu = ((eu + 112u) << 23) | (mu << 13); + float rsc; + memcpy(&rsc, &fu, 4); + output[output_offset + n] = acc * rsc; + } +} + +#endif /* SKAINET_HAVE_NEOGPU_TERNARY */ diff --git a/skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernel.kt b/skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernel.kt new file mode 100644 index 000000000..89b142c69 --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/src/jvmMain/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernel.kt @@ -0,0 +1,113 @@ +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 + +/** + * Native (FFM) downcall to the vendored NeoGPU ternary LUT kernel. + * + * Wraps the bundled C symbol + * + * void skainet_ternary_f32_gemv( + * 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); + * + * Exact FP32 activations against sequentially-packed ternary weights (the + * `BITNET_B1_58` payload: 4 codes per byte, low bit-pair first, + * code {0,1,2} → {-1,0,+1}; byte code 3 decodes to +2 — loaders reject it + * at import). NO scale is applied — the caller owns the per-tensor scale. + * Unlike the int8 `bitnet_gemv` path there is no activation quantization, + * so results are exact. + * + * The C side needs only baseline NEON (no dotprod) and threads internally + * with pthreads once outputDim >= 512. Goldens + threading parity are + * pinned by `NativeTernaryF32GemvKernelTest`. + * + * Refs SKaiNET issue #1137 (vendored from anjaustin/neogpu, MIT — see + * native/src/vendor/neogpu/README.md). The `TernaryF32GemvNative` SPI + * wiring into KernelDispatch follows in #1138. + */ +internal object NativeTernaryF32GemvKernel { + + fun isAvailable(): Boolean = handle != null + + fun gemv( + input: FloatArray, inputOffset: Int, + weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + output: FloatArray, outputOffset: Int, + ) { + require(inputDim % 4 == 0) { + "NativeTernaryF32GemvKernel: inputDim must be a multiple of 4; got $inputDim" + } + if (outputDim == 0) return + + val mh = handle + ?: error("NativeTernaryF32GemvKernel.gemv invoked while native library unavailable") + + // Reach calculations: copy offset + payload into off-heap arenas, + // same convention as the other Native*MatmulKernel objects. + val rowBytes = inputDim / 4 + val inputReachFloats = if (inputDim == 0) 0 else inputOffset + inputDim + val weightReachBytes = if (inputDim == 0 || outputDim == 0) 0 + else weightByteOffset + rowBytes * outputDim + val outputReachFloats = outputOffset + outputDim + + Arena.ofConfined().use { arena -> + val fAlign = ValueLayout.JAVA_FLOAT.byteAlignment() + val bAlign = ValueLayout.JAVA_BYTE.byteAlignment() + + val inputSeg: MemorySegment = if (inputReachFloats > 0) + arena.allocate(inputReachFloats.toLong() * java.lang.Float.BYTES, fAlign) + else MemorySegment.NULL + val weightSeg: MemorySegment = if (weightReachBytes > 0) + arena.allocate(weightReachBytes.toLong(), bAlign) + else MemorySegment.NULL + val outputSeg: MemorySegment = + arena.allocate(outputReachFloats.toLong() * java.lang.Float.BYTES, fAlign) + + if (inputReachFloats > 0) { + MemorySegment.copy(input, 0, inputSeg, ValueLayout.JAVA_FLOAT, 0L, inputReachFloats) + } + if (weightReachBytes > 0) { + MemorySegment.copy(weight, 0, weightSeg, ValueLayout.JAVA_BYTE, 0L, weightReachBytes) + } + // Round-trip the existing output reach so array content before + // outputOffset survives the copy-back (callers loop rows into + // one array at increasing offsets — a zeroed segment would + // clobber the rows already written). + MemorySegment.copy(output, 0, outputSeg, ValueLayout.JAVA_FLOAT, 0L, outputReachFloats) + + mh.invoke( + inputSeg, inputOffset, + weightSeg, weightByteOffset, + inputDim, outputDim, + outputSeg, outputOffset, + ) + + MemorySegment.copy(outputSeg, ValueLayout.JAVA_FLOAT, 0L, output, 0, outputReachFloats) + } + } + + private val handle: MethodHandle? by lazy { + val lookup = NativeLibraryLoader.lookup() ?: return@lazy null + val symbol = lookup.find("skainet_ternary_f32_gemv").orElse(null) ?: return@lazy null + val descriptor = FunctionDescriptor.ofVoid( + ValueLayout.ADDRESS, // input + ValueLayout.JAVA_INT, // input_offset + ValueLayout.ADDRESS, // weight + ValueLayout.JAVA_INT, // weight_byte_offset + ValueLayout.JAVA_INT, // input_dim + ValueLayout.JAVA_INT, // output_dim + ValueLayout.ADDRESS, // output + ValueLayout.JAVA_INT, // output_offset + ) + runCatching { Linker.nativeLinker().downcallHandle(symbol, descriptor) }.getOrNull() + } +} diff --git a/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernelTest.kt b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernelTest.kt new file mode 100644 index 000000000..bda76a9cc --- /dev/null +++ b/skainet-backends/skainet-backend-native-cpu/src/jvmTest/kotlin/sk/ainet/exec/kernel/NativeTernaryF32GemvKernelTest.kt @@ -0,0 +1,173 @@ +package sk.ainet.exec.kernel + +import kotlin.math.abs +import kotlin.random.Random +import kotlin.test.BeforeTest +import kotlin.test.Test +import kotlin.test.assertEquals +import kotlin.test.assertFailsWith +import kotlin.test.assertTrue + +/** + * Goldens for [NativeTernaryF32GemvKernel] (the vendored NeoGPU LUT + * kernel, #1137). The reference decode lives in this file: + * `((byte >> (lane*2)) & 3) - 1` — the sequential `BITNET_B1_58` payload + * rule, including the byte-code-3 → +2 identity. + * + * The same goldens pin BOTH native branches: NEON on arm64 runners, + * the vendored scalar fallback elsewhere (x86 CI). + * + * Integer-valued activations keep every sum exact (well below 2^24), so + * the structural goldens assert bitwise equality; only the random float + * parity case allows a summation-order tolerance. + */ +class NativeTernaryF32GemvKernelTest { + + @BeforeTest + fun checkAvailable() { + assertTrue( + NativeTernaryF32GemvKernel.isAvailable(), + "Native ternary f32 kernel must be available — bundled libskainet_kernels " + + "missing or skainet_ternary_f32_gemv symbol unresolved", + ) + } + + /** Reference decode — deliberately independent of any production codec. */ + private fun decode(b: Byte, lane: Int): Float = + (((b.toInt() and 0xFF) shr (lane * 2)) and 3).toFloat() - 1f + + private fun referenceGemv( + input: FloatArray, weight: ByteArray, weightByteOffset: Int, + inputDim: Int, outputDim: Int, + ): FloatArray { + val rowBytes = inputDim / 4 + return FloatArray(outputDim) { n -> + var acc = 0.0 + for (bi in 0 until rowBytes) { + val b = weight[weightByteOffset + n * rowBytes + bi] + for (lane in 0 until 4) { + acc += decode(b, lane) * input[bi * 4 + lane] + } + } + acc.toFloat() + } + } + + @Test + fun all_256_byte_values_match_reference_exactly() { + // One row whose 256 weight bytes enumerate every possible byte — + // pins the full decode table (incl. code 3 → +2) on whichever + // branch (NEON or scalar) this runner compiled. + val inputDim = 1024 + val input = FloatArray(inputDim) { ((it % 7) - 3).toFloat() } + val weight = ByteArray(256) { it.toByte() } + val expected = referenceGemv(input, weight, 0, inputDim, 1) + val out = FloatArray(1) + NativeTernaryF32GemvKernel.gemv(input, 0, weight, 0, inputDim, 1, out, 0) + assertEquals(expected[0], out[0], "all-bytes golden must match bit-exactly") + } + + @Test + fun byte_code_3_decodes_to_plus_two() { + // 0xFF = four 2-bit codes of 3 → each lane decodes to +2.0. + val out = FloatArray(1) + NativeTernaryF32GemvKernel.gemv( + floatArrayOf(1f, 1f, 1f, 1f), 0, + byteArrayOf(0xFF.toByte()), 0, + 4, 1, out, 0, + ) + assertEquals(8f, out[0], "0xFF row against ones must sum to 4 * (+2)") + } + + @Test + fun random_floats_match_reference_within_summation_tolerance() { + val inputDim = 2560 // BitNet-2B hidden size + val outputDim = 64 + val rng = Random(42) + val input = FloatArray(inputDim) { rng.nextFloat() - 0.5f } + val weight = ByteArray(outputDim * inputDim / 4).also { rng.nextBytes(it) } + val expected = referenceGemv(input, weight, 0, inputDim, outputDim) + val out = FloatArray(outputDim) + NativeTernaryF32GemvKernel.gemv(input, 0, weight, 0, inputDim, outputDim, out, 0) + for (i in out.indices) { + val diff = abs(expected[i] - out[i]) + assertTrue( + diff <= 1e-3f, + "mismatch at $i: reference=${expected[i]} native=${out[i]} diff=$diff", + ) + } + } + + @Test + fun threaded_regime_matches_single_row_calls_exactly() { + // outputDim 2048 crosses the vendored THREAD_THRESHOLD (512): the C + // side fans out over 4 pthreads. Per-row math is independent of the + // partitioning, so the threaded result must equal 2048 single-row + // calls bit-for-bit. + val inputDim = 256 + val outputDim = 2048 + val rowBytes = inputDim / 4 + val rng = Random(7) + val input = FloatArray(inputDim) { rng.nextFloat() - 0.5f } + val weight = ByteArray(outputDim * rowBytes).also { rng.nextBytes(it) } + + val threaded = FloatArray(outputDim) + NativeTernaryF32GemvKernel.gemv(input, 0, weight, 0, inputDim, outputDim, threaded, 0) + + val perRow = FloatArray(outputDim) + for (n in 0 until outputDim) { + NativeTernaryF32GemvKernel.gemv( + input, 0, weight, n * rowBytes, inputDim, 1, perRow, n, + ) + } + for (n in 0 until outputDim) { + assertEquals( + perRow[n], threaded[n], + "thread partitioning changed row $n", + ) + } + } + + @Test + fun offsets_are_honoured() { + val inputDim = 8 + val pad = 3 + val input = FloatArray(pad + inputDim) { if (it < pad) 99f else (it - pad + 1).toFloat() } + // Rows start at byte offset 5. 0x22 = codes {2,0,2,0} → {+1,-1,+1,-1}; + // 0x55 = codes {1,1,1,1} → all zero. + val weight = ByteArray(5 + 2 * 2) { 0x55.toByte() } + weight[5] = 0x22 + weight[6] = 0x22 + val out = FloatArray(4) { -1f } + NativeTernaryF32GemvKernel.gemv(input, pad, weight, 5, inputDim, 2, out, 2) + // in (after offset) = 1..8; row0 decode = {+1,-1,+1,-1, +1,-1,+1,-1} + // → 1-2+3-4+5-6+7-8 = -4; row1 all zeros → 0. + assertEquals(-1f, out[0]); assertEquals(-1f, out[1]) + assertEquals(-4f, out[2]); assertEquals(0f, out[3]) + } + + @Test + fun rejects_non_multiple_of_4_input_dim() { + assertFailsWith { + NativeTernaryF32GemvKernel.gemv( + FloatArray(6), 0, ByteArray(2), 0, 6, 1, FloatArray(1), 0, + ) + } + } + + @Test + fun zero_output_dim_is_no_op() { + NativeTernaryF32GemvKernel.gemv( + FloatArray(4) { 1f }, 0, ByteArray(1), 0, 4, 0, FloatArray(0), 0, + ) + } + + @Test + fun zero_input_dim_zeros_output() { + val out = FloatArray(3) { 9f } + NativeTernaryF32GemvKernel.gemv( + FloatArray(0), 0, ByteArray(0), 0, 0, 3, out, 0, + ) + for (v in out) assertEquals(0f, v, "output should be zeroed for inputDim=0") + } +}