1057 lines
52 KiB
C
1057 lines
52 KiB
C
/* SPDX-License-Identifier: AGPL-3.0-or-later */
|
|
/* Copyright (C) 2026 David (Corey) Herrera */
|
|
/**
|
|
* @file test_attention_ops.c
|
|
* @brief OFF-HARDWARE gate for flash-decode attention: the wire format, the derived
|
|
* quantities, the refusals, the constants, the capacity wall — and the ALGORITHM, scored
|
|
* as a blocked online softmax against a two-pass reference.
|
|
*
|
|
* Built WITHOUT -DFT_TT_PHASE1B_NOC_ENABLED, so every device call in the objects it links
|
|
* would be refused before touching the NoC. It cannot reach the card even if it is wrong,
|
|
* which is what makes it safe to run while somebody else owns the silicon.
|
|
*
|
|
* ---------------------------------------------------------------------------
|
|
* WHAT THIS CAN PROVE AND WHAT IT CANNOT
|
|
* ---------------------------------------------------------------------------
|
|
* It proves the ALGORITHM: that an online softmax over padded, masked KV blocks reproduces
|
|
* a two-pass softmax over the live row, on shapes where the two are not trivially the same.
|
|
* It proves the CONTRACT: that the sixteen words mean on the host what they mean in the
|
|
* kernel, by PARSING the kernel header rather than retyping its constants — the same
|
|
* arrangement test_vector_ops.c uses, and the reason is dev notes §4's catalogue of
|
|
* confident sentences nobody re-derived.
|
|
*
|
|
* It cannot prove the KERNEL. Nothing here executes a tile, a face permutation or a
|
|
* circular buffer; PCC is blind to a consistent permutation, and there is no aperture
|
|
* without a card (dev notes §3). tests/test_phase6_attention.c is where those go red.
|
|
*
|
|
* ---------------------------------------------------------------------------
|
|
* THE INVARIANCE TESTS ARE THE POINT OF THE ALGORITHM SECTION
|
|
* ---------------------------------------------------------------------------
|
|
* softmax(x - c) = softmax(x), so a WRONG RUNNING MAX is structurally invisible to
|
|
* correlation — that exact mutation came back green on seven checks in the reduce suite
|
|
* (docs/tt_bringup_log.md §22.2). Two shapes here exist only to break that symmetry:
|
|
* an outlier planted in a LATE block, so a block-local max produces values outside
|
|
* representable range rather than a shifted-but-equal answer, and a dead-region poison, so
|
|
* a mask that does not mask changes the answer instead of being absorbed.
|
|
*
|
|
* Usage: ./build-attn/test_attention_ops
|
|
*
|
|
* @status See the Phase 6 attention report.
|
|
* @issues None
|
|
* @todo None
|
|
*/
|
|
|
|
#include <math.h>
|
|
#include <stdbool.h>
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
#include <string.h>
|
|
|
|
#include "../convert/ft_float_convert.h"
|
|
#include "../ft_tt_op_attention.h"
|
|
#include "../ft_tt_op_l1_layout.h"
|
|
#include "../ft_tt_op_memory.h"
|
|
#include "../ft_tt_queue.h"
|
|
#include "../ft_tt_tilize.h"
|
|
#include "ft_tt_attention_reference.h"
|
|
#include "ft_tt_test_report.h"
|
|
|
|
/* THE DEVICE HALF, INCLUDED RATHER THAN RETYPED. The chunk mode's mask selection and its
|
|
* geometry refusal are INLINE FUNCTIONS in the kernel header, so including it makes this
|
|
* suite exercise the code the card runs instead of a host copy of it — the same code the
|
|
* reader and the compute thread compile. The constants are still PARSED out of the file by
|
|
* kernel_constant(), because those really are two declarations that can drift.
|
|
* `FTQ_L1_PTR` is the SFPI GCC fork's own pointer attribute and means nothing here. */
|
|
#define FTQ_L1_PTR
|
|
#include "../kernels/resident_attention_protocol.h"
|
|
|
|
#define TILE_DIM 32u
|
|
#define TILE_BYTES FT_TT_OP_TILE_BYTES_BF16
|
|
|
|
/** MiniCPM5-1B, from the GGUF header (docs/tt_phase0_decisions.md §3): 16 Q heads, 2 KV
|
|
* heads, head_dim 128, 24 layers, context 131072. Every shape below is that model's. */
|
|
#define MODEL_Q_HEADS 16u
|
|
#define MODEL_KV_HEADS 2u
|
|
#define MODEL_HEAD_DIM 128u
|
|
#define MODEL_DIM_TILES (MODEL_HEAD_DIM / TILE_DIM)
|
|
#define MODEL_GROUP_ROWS (MODEL_Q_HEADS / MODEL_KV_HEADS)
|
|
|
|
/** Ornith-9B (`general.architecture = qwen35`), the SECOND model this op is asked to serve.
|
|
*
|
|
* **READ OUT OF THE GGUF HEADER, 2026-08-20, after this stage spent a round asserting a
|
|
* shape the model does not have.** `qwen35.attention.head_count = 16`,
|
|
* `head_count_kv = 4`, `key_length = value_length = 256`, `block_count = 32` with 8 of them
|
|
* full attention. The 32/16 this file used to carry is the GDN block's head count, not the
|
|
* attention layers' — and every number derived from it (602 L1 pages, 128 descriptors a
|
|
* token, a 4 GiB KV family) was wrong while the stage reported 0 failures. Pantheon prints
|
|
* the truth at arm time: "16 Q heads / 4 KV groups of 4 rows".
|
|
*
|
|
* It is still the near-inverse of MiniCPM5-1B — a WIDE head (dim_tiles 8 against 4) over a
|
|
* SHALLOWER group (q_rows 4 against 8) — which is why it is worth its own stage: every
|
|
* bound in this op is a product of `q_rows` and `dim_tiles`. */
|
|
#define ORNITH_Q_HEADS 16u
|
|
#define ORNITH_KV_HEADS 4u
|
|
#define ORNITH_HEAD_DIM 256u
|
|
#define ORNITH_ATTN_LAYERS 8u
|
|
#define ORNITH_DIM_TILES (ORNITH_HEAD_DIM / TILE_DIM)
|
|
#define ORNITH_GROUP_ROWS (ORNITH_Q_HEADS / ORNITH_KV_HEADS)
|
|
|
|
/** The residency every algorithm check runs at. */
|
|
#define TEST_BLOCK_TILES 8u
|
|
|
|
/** What `pn_tt_attention.c` ships as its flash-decode prefetch depth. */
|
|
#define SHIPPED_BLOCK_TILES 16u
|
|
|
|
/* ================================================================
|
|
* The kernel header, PARSED — not retyped
|
|
* ================================================================ */
|
|
|
|
/**
|
|
* Read `enum { <name> = <value> };` out of the kernel's protocol header -- the
|
|
* C/C++-shared declaration form the device headers use since plan Phase 6.6.
|
|
*
|
|
* The host and the device halves of this op are two files that must agree word for word.
|
|
* Retyping the device constants here would make the check a copy of a copy: it would pass
|
|
* against a header that had moved, because both would have been edited from the same wrong
|
|
* belief. Reading the actual file is the only version of this check that can fail.
|
|
*
|
|
* @return the value, or UINT32_MAX if the name is absent (which is itself a failure)
|
|
*/
|
|
static uint32_t kernel_constant(const char *name)
|
|
{
|
|
static char buffer[262144];
|
|
static size_t length = 0;
|
|
const char *directory = getenv("FT_TT_KERNEL_DIR");
|
|
char path[512];
|
|
char needle[128];
|
|
const char *found;
|
|
|
|
if (length == 0u) {
|
|
FILE *file;
|
|
|
|
snprintf(path, sizeof(path), "%s/resident_attention_protocol.h",
|
|
directory ? directory : "../kernels");
|
|
file = fopen(path, "rb");
|
|
if (!file) {
|
|
/* SAY WHY, ONCE. Without this the whole wire-contract stage reports UINT32_MAX
|
|
* against every host value and reads as 22 protocol mismatches, when the only
|
|
* fact is that the binary was run from a directory where `../kernels` is not the
|
|
* kernel directory. A harness that fails for a reason other than the one it
|
|
* prints costs a round. */
|
|
static bool said = false;
|
|
|
|
if (!said)
|
|
fprintf(stderr, "test_attention_ops: cannot open %s — set FT_TT_KERNEL_DIR, "
|
|
"or run from tt/tests. Every kernel constant below reads as "
|
|
"UINT32_MAX and NONE of it is a protocol finding.\n", path);
|
|
said = true;
|
|
return UINT32_MAX;
|
|
}
|
|
length = fread(buffer, 1u, sizeof(buffer) - 1u, file);
|
|
buffer[length] = '\0';
|
|
fclose(file);
|
|
}
|
|
|
|
snprintf(needle, sizeof(needle), "enum { %s = ", name);
|
|
found = strstr(buffer, needle);
|
|
if (!found)
|
|
return UINT32_MAX;
|
|
return (uint32_t)strtoul(found + strlen(needle), NULL, 0);
|
|
}
|
|
|
|
static void stage_wire_contract(void)
|
|
{
|
|
static const struct {
|
|
const char *kernel_name;
|
|
uint32_t host_value;
|
|
} pairs[] = {
|
|
{ "kAttnHeads", FT_TT_OP_ATTENTION_HEADS },
|
|
{ "kAttnDimTiles", FT_TT_OP_ATTENTION_DIM_TILES },
|
|
{ "kAttnKvPositions", FT_TT_OP_ATTENTION_KV_POSITIONS },
|
|
{ "kAttnBlockTiles", FT_TT_OP_ATTENTION_BLOCK_TILES },
|
|
{ "kAttnCapacityTiles", FT_TT_OP_ATTENTION_CAPACITY_TILES },
|
|
{ "kAttnReserved", FT_TT_OP_ATTENTION_RESERVED },
|
|
{ "kAttnQCoord", FT_TT_OP_ATTENTION_Q_COORD },
|
|
{ "kAttnQBase", FT_TT_OP_ATTENTION_Q_BASE },
|
|
{ "kAttnKCoord", FT_TT_OP_ATTENTION_K_COORD },
|
|
{ "kAttnKBase", FT_TT_OP_ATTENTION_K_BASE },
|
|
{ "kAttnVCoord", FT_TT_OP_ATTENTION_V_COORD },
|
|
{ "kAttnVBase", FT_TT_OP_ATTENTION_V_BASE },
|
|
{ "kAttnConstCoord", FT_TT_OP_ATTENTION_CONST_COORD },
|
|
{ "kAttnConstBase", FT_TT_OP_ATTENTION_CONST_BASE },
|
|
{ "kAttnOutCoord", FT_TT_OP_ATTENTION_OUT_COORD },
|
|
{ "kAttnOutBase", FT_TT_OP_ATTENTION_OUT_BASE },
|
|
{ "kAttnWords", FT_TT_OP_ATTENTION_WORDS },
|
|
{ "kAttnConstScalerTile", FT_TT_OP_ATTENTION_CONST_SCALER_TILE },
|
|
{ "kAttnConstScaleTile", FT_TT_OP_ATTENTION_CONST_SCALE_TILE },
|
|
{ "kAttnConstMaskTile", FT_TT_OP_ATTENTION_CONST_MASK_TILE },
|
|
/* kAttnConstTiles is DERIVED on both sides (`causal + 32`), so it is checked by the
|
|
* "66 tiles" arithmetic in stage_chunk_contract() rather than pinned here — this
|
|
* parser reads literals, and a pin that reads 0 for a derived name is a false
|
|
* finding, not a check. */
|
|
{ "kAttnConstCausalTile", FT_TT_OP_ATTENTION_CONST_CAUSAL_TILE },
|
|
{ "kAttnFlagChunk", FT_TT_OP_ATTENTION_FLAG_CHUNK },
|
|
{ "kAttnCbConstPages", FT_TT_OP_ATTENTION_CONST_PAGES },
|
|
{ "kAttnTileRows", TILE_DIM }
|
|
};
|
|
size_t index;
|
|
|
|
ft_tt_test_step("the host and the kernel agree word for word — by PARSING the kernel "
|
|
"header, not by retyping it");
|
|
|
|
for (index = 0; index < sizeof(pairs) / sizeof(pairs[0]); index++) {
|
|
const uint32_t device = kernel_constant(pairs[index].kernel_name);
|
|
char label[128];
|
|
|
|
snprintf(label, sizeof(label), "%s", pairs[index].kernel_name);
|
|
ft_tt_test_check_fmt(label, device == pairs[index].host_value, "kernel %u host %u",
|
|
device, pairs[index].host_value);
|
|
}
|
|
}
|
|
|
|
/* ================================================================
|
|
* The parameter block
|
|
* ================================================================ */
|
|
|
|
static void fill_operand(ft_tt_op_operand *operand, uint16_t x, uint16_t y, uint32_t address)
|
|
{
|
|
operand->x = x;
|
|
operand->y = y;
|
|
operand->address = address;
|
|
operand->tile_bytes = TILE_BYTES;
|
|
}
|
|
|
|
/** A request whose every field is a DIFFERENT recognisable value, so a word in the wrong
|
|
* slot cannot pass for its neighbour. */
|
|
static void fill_reference_request(ft_tt_op_attention_request *request)
|
|
{
|
|
memset(request, 0, sizeof(*request));
|
|
request->q_rows = MODEL_GROUP_ROWS;
|
|
request->dim_tiles = MODEL_DIM_TILES;
|
|
request->block_tiles = TEST_BLOCK_TILES;
|
|
request->kv_positions = 1000u;
|
|
request->capacity_tiles = 4096u;
|
|
request->query_row = 7u;
|
|
fill_operand(&request->query, 1u, 2u, 0x00010000u);
|
|
fill_operand(&request->keys, 3u, 4u, 0x00020000u);
|
|
fill_operand(&request->values, 5u, 6u, 0x00030000u);
|
|
fill_operand(&request->constants, 7u, 8u, 0x00040000u);
|
|
fill_operand(&request->output, 9u, 10u, 0x00050000u);
|
|
}
|
|
|
|
static void stage_parameters(void)
|
|
{
|
|
ft_tt_op_attention_request request;
|
|
uint32_t words[FT_TT_OP_ATTENTION_WORDS];
|
|
int rc;
|
|
|
|
ft_tt_test_step("the parameter block, pinned as LITERALS");
|
|
|
|
fill_reference_request(&request);
|
|
rc = ft_tt_op_attention_parameters(&request, words);
|
|
ft_tt_test_check_fmt("a MiniCPM5-1B group encodes", rc == FT_TT_OK, "rc=%d", rc);
|
|
if (rc != FT_TT_OK)
|
|
return;
|
|
|
|
ft_tt_test_check_fmt("heads word packs q_rows and the query row",
|
|
words[FT_TT_OP_ATTENTION_HEADS] == (8u | (7u << 8)), "0x%X",
|
|
words[FT_TT_OP_ATTENTION_HEADS]);
|
|
ft_tt_test_check_fmt("dim tiles", words[FT_TT_OP_ATTENTION_DIM_TILES] == 4u, "%u",
|
|
words[FT_TT_OP_ATTENTION_DIM_TILES]);
|
|
ft_tt_test_check_fmt("kv positions", words[FT_TT_OP_ATTENTION_KV_POSITIONS] == 1000u, "%u",
|
|
words[FT_TT_OP_ATTENTION_KV_POSITIONS]);
|
|
ft_tt_test_check_fmt("block tiles", words[FT_TT_OP_ATTENTION_BLOCK_TILES] == 8u, "%u",
|
|
words[FT_TT_OP_ATTENTION_BLOCK_TILES]);
|
|
ft_tt_test_check_fmt("capacity tiles", words[FT_TT_OP_ATTENTION_CAPACITY_TILES] == 4096u,
|
|
"%u", words[FT_TT_OP_ATTENTION_CAPACITY_TILES]);
|
|
ft_tt_test_check_fmt("reserved is zero", words[FT_TT_OP_ATTENTION_RESERVED] == 0u, "%u",
|
|
words[FT_TT_OP_ATTENTION_RESERVED]);
|
|
ft_tt_test_check_fmt("Q coordinate packs (y << 16) | x",
|
|
words[FT_TT_OP_ATTENTION_Q_COORD] == ((2u << 16) | 1u), "0x%X",
|
|
words[FT_TT_OP_ATTENTION_Q_COORD]);
|
|
ft_tt_test_check_fmt("K coordinate", words[FT_TT_OP_ATTENTION_K_COORD] == ((4u << 16) | 3u),
|
|
"0x%X", words[FT_TT_OP_ATTENTION_K_COORD]);
|
|
ft_tt_test_check_fmt("V coordinate", words[FT_TT_OP_ATTENTION_V_COORD] == ((6u << 16) | 5u),
|
|
"0x%X", words[FT_TT_OP_ATTENTION_V_COORD]);
|
|
ft_tt_test_check_fmt("constants coordinate",
|
|
words[FT_TT_OP_ATTENTION_CONST_COORD] == ((8u << 16) | 7u), "0x%X",
|
|
words[FT_TT_OP_ATTENTION_CONST_COORD]);
|
|
ft_tt_test_check_fmt("output coordinate",
|
|
words[FT_TT_OP_ATTENTION_OUT_COORD] == ((10u << 16) | 9u), "0x%X",
|
|
words[FT_TT_OP_ATTENTION_OUT_COORD]);
|
|
ft_tt_test_check_fmt("every base survives",
|
|
words[FT_TT_OP_ATTENTION_Q_BASE] == 0x00010000u &&
|
|
words[FT_TT_OP_ATTENTION_K_BASE] == 0x00020000u &&
|
|
words[FT_TT_OP_ATTENTION_V_BASE] == 0x00030000u &&
|
|
words[FT_TT_OP_ATTENTION_CONST_BASE] == 0x00040000u &&
|
|
words[FT_TT_OP_ATTENTION_OUT_BASE] == 0x00050000u,
|
|
"0x%X 0x%X 0x%X 0x%X 0x%X", words[FT_TT_OP_ATTENTION_Q_BASE],
|
|
words[FT_TT_OP_ATTENTION_K_BASE], words[FT_TT_OP_ATTENTION_V_BASE],
|
|
words[FT_TT_OP_ATTENTION_CONST_BASE],
|
|
words[FT_TT_OP_ATTENTION_OUT_BASE]);
|
|
ft_tt_test_check("the opcode is 0x11", FT_TT_QUEUE_OP_ATTENTION == 0x11u, "");
|
|
}
|
|
|
|
/** Each refusal names its REASON, not merely "it failed": a check with several sufficient
|
|
* grounds tests whichever is cheapest (dev notes §3). */
|
|
static void stage_refusals(void)
|
|
{
|
|
ft_tt_op_attention_request request;
|
|
uint32_t words[FT_TT_OP_ATTENTION_WORDS];
|
|
|
|
ft_tt_test_step("what the encoder refuses, and for which reason");
|
|
|
|
fill_reference_request(&request);
|
|
request.kv_positions = 0u;
|
|
ft_tt_test_check("a zero KV length is refused",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
|
|
fill_reference_request(&request);
|
|
request.q_rows = 33u;
|
|
ft_tt_test_check("a group wider than a tile is refused",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
|
|
fill_reference_request(&request);
|
|
request.query_row = 32u;
|
|
ft_tt_test_check("a query row past the tile is refused",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
|
|
fill_reference_request(&request);
|
|
request.block_tiles = 0u;
|
|
ft_tt_test_check("a zero block is refused (it is a zero-length NoC read)",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
|
|
/* THE ONE THAT MATTERS. The reader fetches the PADDED window, so a capacity that covers
|
|
* the live length but not the padding is a read past the end of the KV region. */
|
|
fill_reference_request(&request);
|
|
request.kv_positions = 33u * TILE_DIM;
|
|
request.capacity_tiles = 33u;
|
|
ft_tt_test_check("a capacity that covers the LIVE length but not the PADDED window is "
|
|
"refused",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
request.capacity_tiles = 48u;
|
|
ft_tt_test_check("...and the same shape passes once the region holds the padding",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_OK, "");
|
|
|
|
fill_reference_request(&request);
|
|
request.keys.tile_bytes = FT_TT_OP_TILE_BYTES_BFP8_B;
|
|
ft_tt_test_check("a block-float KV operand is refused (ft_tt_op_memory.h §1.6)",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
}
|
|
|
|
static void stage_derived_quantities(void)
|
|
{
|
|
static const struct {
|
|
uint32_t kv;
|
|
uint32_t live;
|
|
uint32_t padded;
|
|
} cases[] = {
|
|
{ 1u, 1u, 16u }, { 32u, 1u, 16u }, { 33u, 2u, 16u },
|
|
{ 512u, 16u, 16u }, { 513u, 17u, 32u }, { 1000u, 32u, 32u },
|
|
{ 8192u, 256u, 256u }, { 8193u, 257u, 272u }
|
|
};
|
|
size_t index;
|
|
|
|
ft_tt_test_step("the derived quantities — live tiles and the padded window at "
|
|
"block_tiles = 8");
|
|
|
|
for (index = 0; index < sizeof(cases) / sizeof(cases[0]); index++) {
|
|
const uint32_t live = ft_tt_op_attention_live_tiles(cases[index].kv);
|
|
const uint32_t padded = ft_tt_op_attention_padded_tiles(cases[index].kv, TEST_BLOCK_TILES);
|
|
char label[96];
|
|
|
|
snprintf(label, sizeof(label), "kv=%u", cases[index].kv);
|
|
ft_tt_test_check_fmt(label, live == cases[index].live && padded == cases[index].padded,
|
|
"live %u (want %u) padded %u (want %u)", live, cases[index].live,
|
|
padded, cases[index].padded);
|
|
}
|
|
|
|
/* The block-prefetch invariant's second condition, as arithmetic: the descriptor has to
|
|
* leave the write pointer where it found it, which means the padded tile count is a
|
|
* whole number of DOUBLE blocks. */
|
|
for (index = 0; index < sizeof(cases) / sizeof(cases[0]); index++) {
|
|
const uint32_t padded = ft_tt_op_attention_padded_tiles(cases[index].kv, TEST_BLOCK_TILES);
|
|
char label[128];
|
|
|
|
snprintf(label, sizeof(label), "kv=%u leaves the CB write pointer at the base",
|
|
cases[index].kv);
|
|
ft_tt_test_check_fmt(label, padded % (2u * TEST_BLOCK_TILES) == 0u, "padded %u", padded);
|
|
}
|
|
}
|
|
|
|
/* ================================================================
|
|
* The constants
|
|
* ================================================================ */
|
|
|
|
static void stage_constants(void)
|
|
{
|
|
static uint16_t constants[FT_TT_OP_ATTENTION_CONST_TILES * TILE_DIM * TILE_DIM];
|
|
const float expected_scale = 1.0f / sqrtf((float)MODEL_HEAD_DIM);
|
|
uint32_t tail;
|
|
int rc;
|
|
|
|
ft_tt_test_step("the constants operand — one scaler, one query scale, 32 masks");
|
|
|
|
rc = ft_tt_op_attention_build_constants(MODEL_HEAD_DIM, TILE_BYTES, constants,
|
|
sizeof(constants));
|
|
ft_tt_test_check_fmt("it builds", rc == FT_TT_OK, "rc=%d", rc);
|
|
if (rc != FT_TT_OK)
|
|
return;
|
|
|
|
ft_tt_test_check_fmt("the reduce scaler is 1.0 everywhere",
|
|
ft_float_from_bf16(constants[0]) == 1.0f &&
|
|
ft_float_from_bf16(constants[1023]) == 1.0f,
|
|
"%f", (double)ft_float_from_bf16(constants[0]));
|
|
{
|
|
const float scale = ft_float_from_bf16(constants[TILE_DIM * TILE_DIM]);
|
|
|
|
ft_tt_test_check_fmt("the query scale is 1/sqrt(head_dim)",
|
|
fabs((double)scale - (double)expected_scale) < 0.002,
|
|
"%f vs %f", (double)scale, (double)expected_scale);
|
|
}
|
|
|
|
/* A mask that is right in face (0,0) and wrong everywhere else is the exact failure a
|
|
* flat-index check cannot see (docs/tt_kernel_model.md §9), so every row of every mask
|
|
* is walked in FACE order. */
|
|
for (tail = 0; tail < TILE_DIM; tail++) {
|
|
const uint16_t *tile =
|
|
constants + (FT_TT_OP_ATTENTION_CONST_MASK_TILE + tail) * TILE_DIM * TILE_DIM;
|
|
bool correct = true;
|
|
uint32_t row;
|
|
|
|
for (row = 0; row < TILE_DIM && correct; row++) {
|
|
uint32_t column;
|
|
|
|
for (column = 0; column < TILE_DIM; column++) {
|
|
const float value = ft_float_from_bf16(tile[ft_tt_tile_face_index(row, column)]);
|
|
const bool live = column < tail;
|
|
|
|
if (live ? (value != 0.0f) : (value > -1.0e29f)) {
|
|
correct = false;
|
|
break;
|
|
}
|
|
}
|
|
}
|
|
if (tail == 0u || tail == 1u || tail == 17u || tail == 31u) {
|
|
char label[96];
|
|
|
|
snprintf(label, sizeof(label), "mask %u is dead at and above column %u, in FACE "
|
|
"order",
|
|
tail, tail);
|
|
ft_tt_test_check(label, correct, "");
|
|
} else {
|
|
ft_tt_test_check("a mask tile is correct", correct, "");
|
|
}
|
|
}
|
|
}
|
|
|
|
/* ================================================================
|
|
* CHUNK MODE — the flag, the geometry refusal, the mask selection, the causal tile
|
|
* ================================================================ */
|
|
|
|
/** A chunk descriptor for `rows` positions starting at `first`. */
|
|
static void fill_chunk_request(ft_tt_op_attention_request *request, uint32_t first,
|
|
uint32_t rows)
|
|
{
|
|
fill_reference_request(request);
|
|
request->q_rows = rows;
|
|
request->query_row = 0u;
|
|
request->flags = FT_TT_OP_ATTENTION_FLAG_CHUNK;
|
|
request->kv_positions = first + rows;
|
|
request->capacity_tiles = 4096u;
|
|
}
|
|
|
|
/** Is causal tile `constants[CAUSAL + live - 1]` live exactly where `c <= r && c < live`,
|
|
* walked in FACE order rather than flat? A mask right in face (0,0) and wrong elsewhere is
|
|
* the exact failure a flat-index check cannot see. */
|
|
static bool causal_tile_is_correct(const uint16_t *constants, uint32_t live)
|
|
{
|
|
const uint16_t *tile =
|
|
constants + (size_t)(FT_TT_OP_ATTENTION_CONST_CAUSAL_TILE + live - 1u) * TILE_DIM *
|
|
TILE_DIM;
|
|
uint32_t row;
|
|
|
|
for (row = 0; row < TILE_DIM; row++) {
|
|
uint32_t column;
|
|
|
|
for (column = 0; column < TILE_DIM; column++) {
|
|
const float value = ft_float_from_bf16(tile[ft_tt_tile_face_index(row, column)]);
|
|
const bool want_live = column <= row && column < live;
|
|
|
|
if (want_live ? (value != 0.0f) : (value > -1.0e29f))
|
|
return false;
|
|
}
|
|
}
|
|
return true;
|
|
}
|
|
|
|
static void stage_chunk_contract(void)
|
|
{
|
|
static uint16_t constants[FT_TT_OP_ATTENTION_CONST_TILES * TILE_DIM * TILE_DIM];
|
|
ft_tt_op_attention_request request;
|
|
uint32_t words[FT_TT_OP_ATTENTION_WORDS];
|
|
fta_shape shape;
|
|
uint32_t live;
|
|
int rc;
|
|
|
|
ft_tt_test_step("CHUNK MODE — one Q head over C positions, and causality as ONE tile");
|
|
|
|
/* ---- the flag rides in byte 2, and the DEVICE decoder reads it back ---- */
|
|
fill_chunk_request(&request, 992u, 32u);
|
|
rc = ft_tt_op_attention_parameters(&request, words);
|
|
ft_tt_test_check_fmt("a 32-position chunk at an aligned first position encodes",
|
|
rc == FT_TT_OK, "rc=%d", rc);
|
|
if (rc != FT_TT_OK)
|
|
return;
|
|
ft_tt_test_check_fmt("the flag rides in byte 2 and C in byte 0 of the heads word",
|
|
words[FT_TT_OP_ATTENTION_HEADS] == (32u | (1u << 16)), "0x%X",
|
|
words[FT_TT_OP_ATTENTION_HEADS]);
|
|
|
|
shape = fta_decode(words);
|
|
ft_tt_test_check_fmt("the DEVICE decoder reads C, row 0 and the chunk flag back",
|
|
shape.q_rows == 32u && shape.query_row == 0u &&
|
|
shape.flags == (uint32_t)kAttnFlagChunk,
|
|
"q_rows %u query_row %u flags %u", shape.q_rows, shape.query_row,
|
|
shape.flags);
|
|
ft_tt_test_check("the device accepts the decoded chunk shape", fta_shape_valid(&shape), "");
|
|
|
|
/* ---- THE MASK SELECTION, WHICH IS THE WHOLE OF CAUSALITY ----
|
|
* At first = 992, C = 32: kv_positions 1024, live_tiles 32, tail 0. DECODE calls the
|
|
* last live tile entirely live; CHUNK must mask it. That one difference is the bug this
|
|
* mode exists to avoid, and it is invisible to every length and address check. */
|
|
ft_tt_test_check_fmt("chunk: a tile before the last live one takes NO mask",
|
|
fta_mask_for_tile(&shape, 30u) == (uint32_t)kAttnCbConstPages, "%u",
|
|
fta_mask_for_tile(&shape, 30u));
|
|
ft_tt_test_check_fmt("chunk: the LAST LIVE tile takes the mask slot even at tail 0",
|
|
fta_mask_for_tile(&shape, 31u) == (uint32_t)kAttnCbConstTailMask,
|
|
"tail %u -> slot %u", shape.tail, fta_mask_for_tile(&shape, 31u));
|
|
ft_tt_test_check_fmt("chunk: a tile past the end takes the all-dead mask",
|
|
fta_mask_for_tile(&shape, 32u) == (uint32_t)kAttnCbConstDeadMask, "%u",
|
|
fta_mask_for_tile(&shape, 32u));
|
|
ft_tt_test_check_fmt("chunk: the tail slot is loaded with the CAUSAL tile for C",
|
|
fta_tail_slot_source_tile(&shape) ==
|
|
(uint32_t)kAttnConstCausalTile + 31u,
|
|
"%u", fta_tail_slot_source_tile(&shape));
|
|
{
|
|
/* The SAME shape without the flag — the discriminating arm. Decode leaves that tile
|
|
* unmasked, which is every position in the chunk attending to every other one. */
|
|
fta_shape decode_shape = shape;
|
|
|
|
decode_shape.flags = 0u;
|
|
ft_tt_test_check_fmt("decode: the SAME shape leaves the last live tile UNMASKED at "
|
|
"tail 0 — which is what chunk mode had to change",
|
|
fta_mask_for_tile(&decode_shape, 31u) ==
|
|
(uint32_t)kAttnCbConstPages,
|
|
"%u", fta_mask_for_tile(&decode_shape, 31u));
|
|
ft_tt_test_check_fmt("decode: the tail slot is loaded with the TAIL mask",
|
|
fta_tail_slot_source_tile(&decode_shape) ==
|
|
(uint32_t)kAttnConstMaskTile + decode_shape.tail,
|
|
"%u", fta_tail_slot_source_tile(&decode_shape));
|
|
}
|
|
|
|
/* ---- a RAGGED chunk: 7 positions at 96, the established split's last piece ---- */
|
|
fill_chunk_request(&request, 96u, 7u);
|
|
rc = ft_tt_op_attention_parameters(&request, words);
|
|
ft_tt_test_check_fmt("the ragged chunk of the 32+32+32+7 split encodes", rc == FT_TT_OK,
|
|
"rc=%d", rc);
|
|
shape = fta_decode(words);
|
|
ft_tt_test_check_fmt("ragged: live_tiles 4, tail 7, and the DIAGONAL tile is tile 3",
|
|
shape.live_tiles == 4u && shape.tail == 7u &&
|
|
fta_mask_for_tile(&shape, 3u) == (uint32_t)kAttnCbConstTailMask,
|
|
"live %u tail %u", shape.live_tiles, shape.tail);
|
|
ft_tt_test_check_fmt("ragged: the causal tile is C's, not the tail's — one tile does BOTH "
|
|
"jobs",
|
|
fta_tail_slot_source_tile(&shape) ==
|
|
(uint32_t)kAttnConstCausalTile + 6u,
|
|
"%u", fta_tail_slot_source_tile(&shape));
|
|
|
|
/* ---- the geometry refusals, on BOTH sides ---- */
|
|
fill_chunk_request(&request, 96u, 7u);
|
|
request.kv_positions = 100u; /* first = 93, not a multiple of 32 */
|
|
ft_tt_test_check("a chunk whose first position is not a multiple of 32 is refused",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
|
|
fill_chunk_request(&request, 96u, 7u);
|
|
request.query_row = 3u;
|
|
ft_tt_test_check("a chunk with a non-zero query row is refused (rows ARE positions)",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
|
|
fill_chunk_request(&request, 96u, 7u);
|
|
request.flags = 2u;
|
|
ft_tt_test_check("an unknown flag bit is refused rather than ignored",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL, "");
|
|
|
|
fill_chunk_request(&request, 992u, 32u);
|
|
(void)ft_tt_op_attention_parameters(&request, words);
|
|
words[FT_TT_OP_ATTENTION_KV_POSITIONS] = 1000u; /* first = 968, unaligned */
|
|
shape = fta_decode(words);
|
|
ft_tt_test_check("...and the DEVICE refuses the unaligned chunk too, not only the host",
|
|
!fta_shape_valid(&shape), "");
|
|
words[FT_TT_OP_ATTENTION_KV_POSITIONS] = 1024u;
|
|
words[FT_TT_OP_ATTENTION_HEADS] = 32u | (3u << 8) | (1u << 16);
|
|
shape = fta_decode(words);
|
|
ft_tt_test_check("...and a chunk carrying a query row", !fta_shape_valid(&shape), "");
|
|
words[FT_TT_OP_ATTENTION_HEADS] = 32u | (2u << 16);
|
|
shape = fta_decode(words);
|
|
ft_tt_test_check("...and an unknown flag bit", !fta_shape_valid(&shape), "");
|
|
|
|
/* ---- the causal tiles themselves ---- */
|
|
rc = ft_tt_op_attention_build_constants(MODEL_HEAD_DIM, TILE_BYTES, constants,
|
|
sizeof(constants));
|
|
ft_tt_test_check_fmt("the constants build with the causal block appended", rc == FT_TT_OK,
|
|
"rc=%d", rc);
|
|
if (rc != FT_TT_OK)
|
|
return;
|
|
ft_tt_test_check_fmt("the operand is 66 tiles — 34 decode reads, 32 appended",
|
|
ft_tt_op_attention_constants_bytes(TILE_BYTES) ==
|
|
(size_t)66u * TILE_BYTES,
|
|
"%zu B", ft_tt_op_attention_constants_bytes(TILE_BYTES));
|
|
/* THE DECODE BLOCK DID NOT MOVE. Tile 2 + tail is still the tail mask, which is what
|
|
* makes this an APPEND rather than a layout change — and what keeps the shipped decode
|
|
* arm's constants byte for byte what they were. */
|
|
ft_tt_test_check("the tail masks are still at tile 2 + tail, so decode's constants are "
|
|
"unmoved",
|
|
FT_TT_OP_ATTENTION_CONST_MASK_TILE == 2u &&
|
|
FT_TT_OP_ATTENTION_CONST_CAUSAL_TILE == 34u,
|
|
"");
|
|
for (live = 1u; live <= TILE_DIM; live++) {
|
|
char label[128];
|
|
const bool correct = causal_tile_is_correct(constants, live);
|
|
|
|
if (live == 1u || live == 7u || live == 17u || live == TILE_DIM) {
|
|
snprintf(label, sizeof(label),
|
|
"causal tile C=%u is live exactly where c <= r and c < %u, in FACE order",
|
|
live, live);
|
|
ft_tt_test_check(label, correct, "");
|
|
} else {
|
|
ft_tt_test_check("a causal tile is correct", correct, "");
|
|
}
|
|
}
|
|
{
|
|
/* A LOWER-TRIANGULAR TILE IS NOT A TAIL MASK, AND THE SUITE SHOULD SAY SO. Without
|
|
* this the whole block above would pass against a builder that emitted 32 copies of
|
|
* the all-live tile. */
|
|
const uint16_t *causal =
|
|
constants + (size_t)(FT_TT_OP_ATTENTION_CONST_CAUSAL_TILE + 31u) * TILE_DIM *
|
|
TILE_DIM;
|
|
uint32_t live_count = 0u;
|
|
uint32_t row;
|
|
|
|
for (row = 0; row < TILE_DIM; row++) {
|
|
uint32_t column;
|
|
|
|
for (column = 0; column < TILE_DIM; column++) {
|
|
if (ft_float_from_bf16(causal[ft_tt_tile_face_index(row, column)]) == 0.0f)
|
|
live_count++;
|
|
}
|
|
}
|
|
ft_tt_test_check_fmt("the C=32 causal tile is half a tile live, not a whole one",
|
|
live_count == (TILE_DIM * (TILE_DIM + 1u)) / 2u,
|
|
"%u live of %u", live_count, TILE_DIM * TILE_DIM);
|
|
}
|
|
}
|
|
|
|
/* ================================================================
|
|
* The residency and the capacity wall
|
|
* ================================================================ */
|
|
|
|
static void stage_capacity(void)
|
|
{
|
|
const size_t arena = (size_t)(FT_TT_OP_CB_ARENA_END - FT_TT_OP_CB_ARENA_BASE);
|
|
ft_tt_op_attention_residency residency;
|
|
size_t short_context;
|
|
size_t long_context;
|
|
uint32_t best;
|
|
|
|
ft_tt_test_step("the capacity wall — why the online softmax is required and not "
|
|
"preferred");
|
|
|
|
/* A STANDALONE softmax holds the row twice over, so its ceiling is a row width. A
|
|
* 32k-context attention row is Wt = 1024 and this is the assertion that says so. */
|
|
{
|
|
const size_t standalone_row = (size_t)1024u * TILE_BYTES;
|
|
|
|
ft_tt_test_check_fmt("a 32k score row does not fit L1 even ONCE, let alone twice",
|
|
2u * standalone_row > arena, "row %zu arena %zu", standalone_row,
|
|
arena);
|
|
}
|
|
|
|
residency.q_rows = MODEL_GROUP_ROWS;
|
|
residency.dim_tiles = MODEL_DIM_TILES;
|
|
residency.block_tiles = TEST_BLOCK_TILES;
|
|
residency.tile_bytes = TILE_BYTES;
|
|
short_context = ft_tt_op_attention_l1_bytes(&residency);
|
|
long_context = ft_tt_op_attention_l1_bytes(&residency);
|
|
ft_tt_test_check_fmt("attention's L1 is FLAT in the context length — the same residency "
|
|
"serves 512 and 131072",
|
|
short_context == long_context && short_context != 0u, "%zu bytes",
|
|
short_context);
|
|
ft_tt_test_check_fmt("and it fits the arena with room to spare", short_context < arena / 2u,
|
|
"%zu of %zu", short_context, arena);
|
|
|
|
best = ft_tt_op_attention_max_block_tiles(MODEL_GROUP_ROWS, MODEL_DIM_TILES, TILE_BYTES);
|
|
ft_tt_test_check_fmt("the deepest block that fits is reported and is well past 8",
|
|
best >= 16u, "block_tiles <= %u", best);
|
|
residency.block_tiles = best;
|
|
ft_tt_test_check_fmt("...and the residency at that depth still fits",
|
|
ft_tt_op_attention_l1_bytes(&residency) <= arena, "%zu of %zu",
|
|
ft_tt_op_attention_l1_bytes(&residency), arena);
|
|
residency.block_tiles = best + 1u;
|
|
ft_tt_test_check_fmt("...and one deeper does not", ft_tt_op_attention_l1_bytes(&residency) >
|
|
arena,
|
|
"%zu of %zu", ft_tt_op_attention_l1_bytes(&residency), arena);
|
|
}
|
|
|
|
/* ================================================================
|
|
* Ornith-9B's shape
|
|
* ================================================================ */
|
|
|
|
/**
|
|
* Everything this op has to be true of before Ornith's attention can move onto the card.
|
|
*
|
|
* NOT a second copy of the checks above at different numbers. Every check here is one that
|
|
* MiniCPM5-1B's shape cannot fail and Ornith's can, because the two models sit on opposite
|
|
* sides of the one product this op is bounded by:
|
|
*
|
|
* MiniCPM5-1B q_rows 8, dim_tiles 4 deep group, narrow head
|
|
* Ornith-9B q_rows 2, dim_tiles 8 shallow group, WIDE head
|
|
*
|
|
* K and V are `2 * block_tiles * dim_tiles` pages EACH, so doubling `dim_tiles` doubles the
|
|
* dominant term of the residency while halving the group count does not pay any of it back.
|
|
* That is the check that had to be run before a line of host wiring was written, and it is
|
|
* the reason it is run with no card: an L1 refusal is discovered at ARM time, which is
|
|
* after a 9.5 GB model load on a 7.1 GB host.
|
|
*/
|
|
static void stage_ornith_shape(void)
|
|
{
|
|
const size_t arena = (size_t)(FT_TT_OP_CB_ARENA_END - FT_TT_OP_CB_ARENA_BASE);
|
|
ft_tt_op_attention_residency residency;
|
|
ft_tt_kv_cache_layout layout;
|
|
size_t l1;
|
|
uint32_t best;
|
|
uint32_t head;
|
|
|
|
ft_tt_test_step("Ornith-9B's shape — 16 Q heads / 4 KV heads / head_dim 256, "
|
|
"8 attention layers");
|
|
|
|
ft_tt_test_check_fmt("the group shape is q_rows=4 over dim_tiles=8",
|
|
ORNITH_GROUP_ROWS == 4u && ORNITH_DIM_TILES == 8u, "q_rows %u "
|
|
"dim_tiles %u", ORNITH_GROUP_ROWS, ORNITH_DIM_TILES);
|
|
|
|
/* ---- the residency, at the depth pn_tt_attention.c actually ships ---- */
|
|
residency.q_rows = ORNITH_GROUP_ROWS;
|
|
residency.dim_tiles = ORNITH_DIM_TILES;
|
|
residency.block_tiles = SHIPPED_BLOCK_TILES;
|
|
residency.tile_bytes = TILE_BYTES;
|
|
l1 = ft_tt_op_attention_l1_bytes(&residency);
|
|
|
|
/* 618 pages at q_rows=4. The deeper group costs 16 pages over q_rows=2's 602 — Q, QS,
|
|
* O, PV and the output staging all carry `q_rows` rows — and it is 93.6% of the arena
|
|
* against 91.2%. Written out because a number this close to the wall must be reproduced
|
|
* by the arithmetic and not merely by the function under test. */
|
|
ft_tt_test_check_fmt("L1 at the shipped block_tiles=16 is 618 pages = 1,265,664 B",
|
|
l1 == 618u * TILE_BYTES, "%zu B (%zu pages)", l1, l1 / TILE_BYTES);
|
|
ft_tt_test_check_fmt("...and it FITS the arena — the shipped depth needs no clamp",
|
|
l1 != 0u && l1 <= arena, "%zu of %zu, %.1f%% used", l1, arena,
|
|
100.0 * (double)l1 / (double)arena);
|
|
/* MiniCPM's stage asserts `< arena/2`. Ornith cannot: at 91% it has ONE deeper block of
|
|
* headroom, and that is the finding, not a rounding detail. */
|
|
ft_tt_test_check_fmt("...with only ~6% spare, unlike MiniCPM5-1B's half-empty arena",
|
|
l1 > arena / 2u, "%zu of %zu", l1, arena);
|
|
|
|
best = ft_tt_op_attention_max_block_tiles(ORNITH_GROUP_ROWS, ORNITH_DIM_TILES, TILE_BYTES);
|
|
ft_tt_test_check_fmt("the deepest block that fits Ornith's shape is 17", best == 17u,
|
|
"block_tiles <= %u", best);
|
|
ft_tt_test_check_fmt("...so the shipped 16 is under the wall and is NOT clamped",
|
|
SHIPPED_BLOCK_TILES <= best, "shipped %u, wall %u",
|
|
SHIPPED_BLOCK_TILES, best);
|
|
residency.block_tiles = best + 1u;
|
|
ft_tt_test_check_fmt("...and one deeper does not fit",
|
|
ft_tt_op_attention_l1_bytes(&residency) > arena, "%zu of %zu",
|
|
ft_tt_op_attention_l1_bytes(&residency), arena);
|
|
ft_tt_test_check_fmt("the WIDE head is what costs it — MiniCPM's shape reaches 34",
|
|
ft_tt_op_attention_max_block_tiles(MODEL_GROUP_ROWS, MODEL_DIM_TILES,
|
|
TILE_BYTES) == 34u,
|
|
"MiniCPM wall %u",
|
|
ft_tt_op_attention_max_block_tiles(MODEL_GROUP_ROWS, MODEL_DIM_TILES,
|
|
TILE_BYTES));
|
|
|
|
/* ---- the GQA mapping: the op groups CONTIGUOUSLY, and so does Ornith ---- */
|
|
{
|
|
bool agrees = true;
|
|
|
|
for (head = 0; head < ORNITH_Q_HEADS; head++) {
|
|
/* What `pn_tt_attention()` implies: group g's query operand starts at tile
|
|
* `g * q_rows * dim_tiles`, i.e. head `g * q_rows`, so head h is served by
|
|
* group `h / q_rows`. */
|
|
const uint32_t op_group = head / ORNITH_GROUP_ROWS;
|
|
/* What Pantheon's CPU oracle does (`qwen35moe_attn.c`,
|
|
* `kv_h = qh / a->n_kv_groups` with `n_kv_groups = n_head / n_head_kv`). */
|
|
const uint32_t oracle_kv_head = head / (ORNITH_Q_HEADS / ORNITH_KV_HEADS);
|
|
|
|
if (op_group != oracle_kv_head)
|
|
agrees = false;
|
|
}
|
|
ft_tt_test_check("the op's KV grouping IS Ornith's GQA repeat — CONTIGUOUS, "
|
|
"kv_head = q_head / 4, on all 16 heads", agrees,
|
|
"checked every Q head");
|
|
}
|
|
|
|
/* ---- the joint Q+gate projection cannot be read in place ---- */
|
|
{
|
|
/* `blk.N.attn_q.weight` emits 2*q_heads*head_dim = 8,192 features: head h's Q at
|
|
* tiles [16h, 16h+8) and head h's GATE at [16h+8, 16h+16). The op reads group g as
|
|
* ONE contiguous run of `q_rows*dim_tiles = 32` tiles from `g*32`, which over that
|
|
* strip is head 4g's Q followed by head 4g's GATE — not head 4g's Q followed by head
|
|
* (4g+1)'s. This is the check that says the de-interleave is REQUIRED rather than
|
|
* tidy, and it is the only thing standing between the op and Ornith. */
|
|
const uint32_t joint_head_stride = 2u * ORNITH_DIM_TILES;
|
|
const uint32_t group = 3u;
|
|
const uint32_t run_first = group * ORNITH_GROUP_ROWS * ORNITH_DIM_TILES;
|
|
const uint32_t wanted_second_head_tile =
|
|
(group * ORNITH_GROUP_ROWS + 1u) * joint_head_stride;
|
|
|
|
ft_tt_test_check_fmt("the op's contiguous group run does NOT land on the joint "
|
|
"Q+gate strip — the de-interleave is required",
|
|
run_first + ORNITH_DIM_TILES != wanted_second_head_tile,
|
|
"group %u would read tiles %u.. but head %u's Q starts at %u",
|
|
group, run_first, group * ORNITH_GROUP_ROWS + 1u,
|
|
wanted_second_head_tile);
|
|
ft_tt_test_check_fmt("...and a DE-INTERLEAVED strip does land on it exactly",
|
|
(group * ORNITH_GROUP_ROWS + 1u) * ORNITH_DIM_TILES ==
|
|
run_first + ORNITH_DIM_TILES,
|
|
"tile %u", run_first + ORNITH_DIM_TILES);
|
|
}
|
|
|
|
/* ---- the padded read window, and what it costs at SHORT context ---- */
|
|
{
|
|
static const struct {
|
|
uint32_t kv;
|
|
uint32_t padded_at_16;
|
|
uint32_t padded_at_1;
|
|
} cases[] = {
|
|
{ 50u, 32u, 2u }, { 1024u, 32u, 32u }, { 1025u, 64u, 34u },
|
|
{ 8192u, 256u, 256u }, { 14336u, 448u, 448u }, { 32768u, 1024u, 1024u }
|
|
};
|
|
size_t index;
|
|
|
|
for (index = 0; index < sizeof(cases) / sizeof(cases[0]); index++) {
|
|
const uint32_t at16 = ft_tt_op_attention_padded_tiles(cases[index].kv,
|
|
SHIPPED_BLOCK_TILES);
|
|
const uint32_t at1 = ft_tt_op_attention_padded_tiles(cases[index].kv, 1u);
|
|
char label[96];
|
|
|
|
snprintf(label, sizeof(label), "kv=%u padded window", cases[index].kv);
|
|
ft_tt_test_check_fmt(label,
|
|
at16 == cases[index].padded_at_16 &&
|
|
at1 == cases[index].padded_at_1,
|
|
"block16 %u (want %u), block1 %u (want %u)", at16,
|
|
cases[index].padded_at_16, at1, cases[index].padded_at_1);
|
|
}
|
|
|
|
/* THE 48-TOKEN ORACLE RUNS HERE, and this is why its arm needs a different knob.
|
|
* At block_tiles=16 the window quantum is 2*16*32 = 1024 POSITIONS, so a descriptor
|
|
* at position 50 does 1024 positions of work — 16x what block_tiles=1 does. The
|
|
* shipped depth is right at context and wrong at the bench length, which is exactly
|
|
* the trade `PN_TT_ATTN_BLOCK_TILES` exists to make. */
|
|
ft_tt_test_check_fmt("at kv=50 the shipped depth does 16x the work of block_tiles=1",
|
|
ft_tt_op_attention_padded_tiles(50u, SHIPPED_BLOCK_TILES) ==
|
|
16u * ft_tt_op_attention_padded_tiles(50u, 1u),
|
|
"%u tiles vs %u",
|
|
ft_tt_op_attention_padded_tiles(50u, SHIPPED_BLOCK_TILES),
|
|
ft_tt_op_attention_padded_tiles(50u, 1u));
|
|
}
|
|
|
|
/* ---- the KV cache Ornith needs, and what it costs in card DRAM ---- */
|
|
ft_tt_test_check_fmt("a 32k KV layout for 4 heads x 256 dim is accepted",
|
|
ft_tt_kv_cache_layout_init(ORNITH_KV_HEADS, ORNITH_HEAD_DIM, 32768u,
|
|
TILE_BYTES, &layout) == FT_TT_OK,
|
|
"4 x 256 x 32768 bf16");
|
|
ft_tt_test_check_fmt("...capacity_tiles 1024, position tile 16,384 B, head stride 16 MiB",
|
|
layout.capacity_tiles == 1024u &&
|
|
layout.position_tile_bytes == 16384u &&
|
|
layout.head_stride == 16777216ull,
|
|
"tiles %u, position %u B, stride %llu B", layout.capacity_tiles,
|
|
layout.position_tile_bytes,
|
|
(unsigned long long)layout.head_stride);
|
|
ft_tt_test_check_fmt("...one region is 64 MiB, so K+V over 8 attention layers is 1 GiB",
|
|
layout.region_bytes == 67108864ull &&
|
|
layout.region_bytes * 2ull * ORNITH_ATTN_LAYERS ==
|
|
1073741824ull,
|
|
"region %llu B, family %.2f GiB",
|
|
(unsigned long long)layout.region_bytes,
|
|
(double)layout.region_bytes * 2.0 * ORNITH_ATTN_LAYERS /
|
|
(1024.0 * 1024.0 * 1024.0));
|
|
/* AT THE REAL HEAD COUNT this is a QUARTER of what the spec carries. The per-position
|
|
* cost across the 8 attention layers is 8*4*256*2 B *2 = 32,768 B on the card at bf16 and
|
|
* 65,536 B in the host fp32 mirror — not the 131,072 / 262,144 derived from 16 KV heads.
|
|
* The host-attention SLOPE (0.53 ms/position) is a measurement and stands; what changes
|
|
* is the byte attribution under it. */
|
|
ft_tt_test_check_fmt("one position of context is 32,768 B on card and 65,536 B on host",
|
|
(uint64_t)ORNITH_ATTN_LAYERS * ORNITH_KV_HEADS * ORNITH_HEAD_DIM *
|
|
2ull * 2ull == 32768ull,
|
|
"%llu B",
|
|
(unsigned long long)((uint64_t)ORNITH_ATTN_LAYERS * ORNITH_KV_HEADS *
|
|
ORNITH_HEAD_DIM * 2ull * 2ull));
|
|
|
|
/* ---- the capacity wall is enforced, not assumed ---- */
|
|
{
|
|
ft_tt_op_attention_request request;
|
|
uint32_t words[FT_TT_OP_ATTENTION_WORDS];
|
|
const ft_tt_op_operand operand = { 1u, 1u, 0x1000u, TILE_BYTES };
|
|
|
|
memset(&request, 0, sizeof(request));
|
|
request.q_rows = ORNITH_GROUP_ROWS;
|
|
request.dim_tiles = ORNITH_DIM_TILES;
|
|
request.block_tiles = SHIPPED_BLOCK_TILES;
|
|
request.capacity_tiles = layout.capacity_tiles;
|
|
request.query_row = 0u;
|
|
request.query = operand;
|
|
request.keys = operand;
|
|
request.values = operand;
|
|
request.constants = operand;
|
|
request.output = operand;
|
|
|
|
request.kv_positions = 32768u;
|
|
ft_tt_test_check("a full 32,768-position descriptor is encodable",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_OK,
|
|
"padded 1024 == capacity 1024");
|
|
request.kv_positions = 32769u;
|
|
ft_tt_test_check("...and one position past the cache is REFUSED, not clamped",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_EINVAL,
|
|
"padded 1056 > capacity 1024");
|
|
request.kv_positions = 50u;
|
|
ft_tt_test_check_fmt("the bench-length descriptor encodes with q_rows=4 in byte 0",
|
|
ft_tt_op_attention_parameters(&request, words) == FT_TT_OK &&
|
|
(words[FT_TT_OP_ATTENTION_HEADS] & 0xffu) ==
|
|
ORNITH_GROUP_ROWS &&
|
|
words[FT_TT_OP_ATTENTION_DIM_TILES] == ORNITH_DIM_TILES,
|
|
"heads word 0x%08x, dim_tiles %u", words[FT_TT_OP_ATTENTION_HEADS],
|
|
words[FT_TT_OP_ATTENTION_DIM_TILES]);
|
|
}
|
|
|
|
/* ---- the descriptor budget, which is the other thing the shape swap moves ---- */
|
|
ft_tt_test_check_fmt("Ornith issues 32 attention descriptors a token against "
|
|
"MiniCPM5-1B's 48",
|
|
ORNITH_KV_HEADS * ORNITH_ATTN_LAYERS == 32u &&
|
|
MODEL_KV_HEADS * 24u == 48u,
|
|
"%u groups x %u layers", ORNITH_KV_HEADS, ORNITH_ATTN_LAYERS);
|
|
}
|
|
|
|
/* ================================================================
|
|
* The algorithm
|
|
* ================================================================ */
|
|
|
|
#define MAX_KV 2048u
|
|
#define MAX_ELEMENTS (MODEL_GROUP_ROWS * MODEL_HEAD_DIM)
|
|
|
|
static float g_query[MAX_ELEMENTS];
|
|
static float g_keys[MAX_KV * MODEL_HEAD_DIM];
|
|
static float g_values[MAX_KV * MODEL_HEAD_DIM];
|
|
static float g_reference[MAX_ELEMENTS];
|
|
static float g_model[MAX_ELEMENTS];
|
|
|
|
/** Fill the whole PADDED window, including the region past the live length: the model reads
|
|
* it, the card reads it, and filling only the live rows would hide the mask bug. */
|
|
static void build_problem(uint32_t kv_positions, uint64_t seed, float dead_value)
|
|
{
|
|
const uint32_t padded = ft_tt_op_attention_padded_tiles(kv_positions, TEST_BLOCK_TILES) *
|
|
TILE_DIM;
|
|
uint32_t index;
|
|
|
|
ft_tt_attention_seed(seed);
|
|
for (index = 0; index < MAX_ELEMENTS; index++)
|
|
g_query[index] = ft_float_from_bf16(ft_bf16_from_float(ft_tt_attention_random()));
|
|
for (index = 0; index < padded * MODEL_HEAD_DIM; index++) {
|
|
const bool dead = (index / MODEL_HEAD_DIM) >= kv_positions;
|
|
const float value = dead ? dead_value : ft_tt_attention_random();
|
|
|
|
g_keys[index] = ft_float_from_bf16(ft_bf16_from_float(value));
|
|
g_values[index] = ft_float_from_bf16(ft_bf16_from_float(dead ? dead_value
|
|
: ft_tt_attention_random()));
|
|
}
|
|
}
|
|
|
|
static void score_model(const char *label, uint32_t kv_positions)
|
|
{
|
|
ft_tt_attention_problem problem;
|
|
double pcc;
|
|
double max_abs;
|
|
char text[160];
|
|
|
|
problem.q_rows = MODEL_GROUP_ROWS;
|
|
problem.head_dim = MODEL_HEAD_DIM;
|
|
problem.kv_positions = kv_positions;
|
|
problem.query = g_query;
|
|
problem.keys = g_keys;
|
|
problem.values = g_values;
|
|
|
|
if (!ft_tt_attention_reference(&problem, g_reference) ||
|
|
!ft_tt_attention_online_model(&problem, TEST_BLOCK_TILES, g_model)) {
|
|
ft_tt_test_check_fmt(label, false, "the reference or the model refused the shape");
|
|
return;
|
|
}
|
|
pcc = ft_tt_attention_pcc(g_reference, g_model, MAX_ELEMENTS);
|
|
max_abs = ft_tt_attention_max_abs(g_reference, g_model, MAX_ELEMENTS);
|
|
snprintf(text, sizeof(text), "%s (kv=%u)", label, kv_positions);
|
|
ft_tt_test_check_fmt(text,
|
|
pcc > 0.9999999 && max_abs < 1.0e-5 &&
|
|
ft_tt_attention_non_finite(g_model, MAX_ELEMENTS) == 0u,
|
|
"pcc %.9f max-abs %.3e non-finite %zu", pcc, max_abs,
|
|
ft_tt_attention_non_finite(g_model, MAX_ELEMENTS));
|
|
}
|
|
|
|
static void stage_algorithm(void)
|
|
{
|
|
ft_tt_test_step("the ONLINE softmax reproduces a two-pass softmax");
|
|
|
|
build_problem(1024u, 20260812ull, 0.0f);
|
|
score_model("a whole number of blocks", 1024u);
|
|
|
|
build_problem(1000u, 20260813ull, 0.0f);
|
|
score_model("a partial last tile — the additive tail mask", 1000u);
|
|
|
|
build_problem(33u, 20260814ull, 0.0f);
|
|
score_model("two live tiles inside a 16-tile padded window", 33u);
|
|
|
|
build_problem(1u, 20260815ull, 0.0f);
|
|
score_model("one live position", 1u);
|
|
|
|
ft_tt_test_step("the two invariance tests — what PCC alone cannot see");
|
|
|
|
/* 1. A WRONG RUNNING MAX. softmax(x - c) = softmax(x), so a block-local max is
|
|
* invisible to correlation on well-scaled data (docs/tt_bringup_log.md §22.2). An
|
|
* outlier in a LATE block breaks the symmetry: a max taken per block puts the early
|
|
* blocks' exponentials outside representable range. */
|
|
build_problem(1024u, 20260816ull, 0.0f);
|
|
{
|
|
uint32_t element;
|
|
|
|
for (element = 0; element < MODEL_HEAD_DIM; element++)
|
|
g_keys[(size_t)1000u * MODEL_HEAD_DIM + element] = 40.0f;
|
|
}
|
|
score_model("an outlier planted in a LATE KV block", 1024u);
|
|
|
|
/* 2. A MASK THAT DOES NOT MASK. Poison the dead region with values that would dominate
|
|
* the row max if they reached it. The reference never looks past the live length, so
|
|
* this changes the model's answer if and only if the mask is not doing its job. */
|
|
build_problem(1000u, 20260817ull, 500.0f);
|
|
score_model("garbage poisoned into the dead region past the live length", 1000u);
|
|
}
|
|
|
|
int main(void)
|
|
{
|
|
ft_tt_test_step("Phase 6 — flash-decode attention, OFF-HARDWARE");
|
|
|
|
stage_wire_contract();
|
|
stage_parameters();
|
|
stage_refusals();
|
|
stage_derived_quantities();
|
|
stage_constants();
|
|
stage_chunk_contract();
|
|
stage_capacity();
|
|
stage_ornith_shape();
|
|
stage_algorithm();
|
|
return ft_tt_test_finish();
|
|
}
|