qsa_real_shapes() takes the compiled geometry; gdn_parity's prefill cases and native_grouped_parity follow the build's shape

qsa_real_shapes() (24 / 2 / 256 / 64, indexer 4 x 128) is the shape every QSA parity program builds from; it now reads kCompiledGeometry (the same values for the default build), so the
q4_* builds' kv_stream / qsa_prompt_attn / snapshot checks run at their own head counts instead of being refused.  gdn_parity's prefill cases used 48 value heads; native_grouped_parity's
compiled-shape loop no longer asks the K-quant checker for an IQ4_NL down.
This commit is contained in:
Niko1221 committed 2026-11-08 07:00:42 +01:00
1 parent 8893ba5eec
commit 41717f29c1
3 files changed
+8 -4

No files matched your search

+6 -2
View File
@@ -87,6 +87,8 @@
// The RAW tail is fp32 for the same reason and because it is only `idx_block - 1` rows.
#pragma once
#include "strata/core/geometry.hpp"
#include <cstdint>
namespace strata::kernels {
@@ -108,8 +110,10 @@ struct QsaShapes {
/// The real artifact's geometry, so no caller transcribes it by hand.
inline QsaShapes qsa_real_shapes() {
QsaShapes s;
s.n_head = 24; s.n_head_kv = 2; s.head_dim = 256; s.n_rot = 64;
s.idx_n_head = 4; s.idx_dim = 128; s.idx_block = 4; s.idx_top_k = 2048;
// the shape the engine is compiled for (core/geometry.hpp: 24 / 2 / 256 / 64, indexer 4 x 128, block 4, top-k 2048 for Flash-Next)
const auto& g = strata::core::kCompiledGeometry;
s.n_head = g.n_head; s.n_head_kv = g.n_head_kv; s.head_dim = g.head_dim; s.n_rot = g.n_rot;
s.idx_n_head = g.idx_q_heads; s.idx_dim = g.idx_key_dim; s.idx_block = g.idx_block; s.idx_top_k = g.idx_top_k;
// one page = one indexer block (4 cells): the granule KV streaming keeps resident (kv_stream.hpp). The readers
// resolve a row per cell anyway, so the page size costs them nothing (measured: identical output and speed).
s.page_size = 4;
+1 -1
View File
@@ -100,7 +100,7 @@ double rel_l1(const std::vector<float>& a, const std::vector<float>& b) {
// Fixtures have nonzero initial state, modulo-paired heads, signed q/k/v, and
// both shallow and deep decay. A tolerance is diagnostic only, never a pass.
int prefill_case(int T, int seed, int repeats) {
constexpr int S = 128, HK = 16, HV = 48, C = 10240;
constexpr int S = strata::core::kCompiledGeometry.ssm_state_size, HK = strata::core::kCompiledGeometry.ssm_k_heads, HV = strata::core::kCompiledGeometry.ssm_v_heads, C = strata::core::kCompiledGeometry.ssm_conv_channels; // the compiled shape
sycl::queue q(dpct::get_in_order_queue().get_device(),
sycl::property_list{sycl::property::queue::in_order{}});
// Variant 6 runs on Intel GPUs only (the engine throws elsewhere): skip, not fail, on any other device.
+1 -1
View File
@@ -764,7 +764,7 @@ int main(int argc, char** argv) {
#ifdef STRATA_Q6K_EXPERTS
14,
#endif
7, 20, 8})
7, 8}) // (check_kq's reference has no IQ4_NL down: those pairs are in the loop below)
check_kq(gu, dt, KH, KF, s, rng);
for (int gu : {29, 16, 17, 22, 42, 18, 23, 21})
for (int dt : {42, 20, 6, 8}) check(gu, dt, KH, KF, s, rng);