igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity

class v5 kits: every worker host uploads the pack's state leaves for igneum_build (the NVRTC worker, the OpenCL host for AMD, Intel and Apple, the Metal worker), checked against the pack's count and FNV first, freed after the build; a v5 pack without its leaves builds nothing

Browse source 
packfile.h reads IGNEUM_STATE_LEAVES, IGNEUM_STATE_LEAVES_FNV64, the file name, root and block (generator 5 = class v5; a generator 5 pack without the count, or another class with one, is refused) and pf_load_leaves reads and checks leaves.bin; the loader test carries the class v5 cases known-failed first (no count, no file, a short file, one flipped bit, a v4 pack with a count) and the checked-in v5-dn3-epoch0 pack. The CUDA driver gains cuMemcpyHtoD; the emulation takes both igneum_build shapes (weak overloads) and refuses a shape mismatch. The OpenCL host sets the build arguments in one place for both shapes, uploads the leaves in --bench-pack, --serve and the compiled-in bench (the host reference derives words from the same leaves). The Metal worker's servePackDataset binds buffers 2 and 3 as packbench.swift does and self-tests the dataset head and last word against vectors.json. tools/class-v5: kits-remote.sh (the kit build and tests on igneum-build-1, the zip to /srv/artefacts/packs), fleet-cuda-v5-bench.sh (the fleet card), pc1-amd-v5-bench.ps1 (PC 1's 9070 XT through the coordinator).

Measured on the Mac under the measure lock, 7 October 2026: v5-dn3-epoch0 fingerprint 82b19cbde8557ea5 on Metal (packbench, M5 Max, 18:34:47Z, 14.163 MH/s GPU time) and on Apple OpenCL (--bench-pack, 18:34:59Z, 15.265 MH/s wall; the compiled-in bench's dataset self-test PASS at 18:34:54Z), cache FNV 7334fa46e5d972eb PASS, vectors 3 of 3 on both.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs 2026-10-07 18:36:17 +00:00
parent 6fdd357434
commit 8bf1c4f43d
11 changed files with 624 additions and 46 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
proto-cuda/nvrtc/cuda_api.h
View file

@ -23,6 +23,7 @@ struct Drv {
decltype(&cuMemAlloc_v2) memAlloc = nullptr;
decltype(&cuMemFree_v2) memFree = nullptr;
decltype(&cuMemcpyDtoH_v2) memcpyDtoH = nullptr;
decltype(&cuMemcpyHtoD_v2) memcpyHtoD = nullptr; // class v5: the state leaves go up for igneum_build (7 October 2026)
decltype(&cuModuleLoadData) moduleLoadData = nullptr;
decltype(&cuModuleUnload) moduleUnload = nullptr;
decltype(&cuModuleGetFunction) moduleGetFunction = nullptr;

36
proto-cuda/nvrtc/emu/emu_backend.cpp
View file

@ -24,20 +24,28 @@
#include "../cuda_api.h" // the real cuda.h / nvrtc.h types
#include "cuda_runtime.h" // the shim (proto-cuda/emu): emu_launch and the kernel symbols' world
// igneum_build has two shapes: (ds, cache, nItems) for classes v2 to v4 and (ds, cache, leaves, nLeaves, nItems) for class v5
// (7 October 2026). A pack defines one of them; both are declared weak so the one the pack lacks is a null symbol, and
// d_launch picks the shape by the argument count the worker passed (5 for a class v5 pack).
#define EMU_WEAK __attribute__((weak))
namespace emu_pack_a {
struct IgneumInitWords { uint32_t w[8]; }; // the same definition as the pack's kernel_bound.cu, inside its namespace
void igneum_cache_fill(uint32_t* cache, uint32_t nSegments);
void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems);
EMU_WEAK void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems);
EMU_WEAK void igneum_build(uint32_t* ds, const uint32_t* cache, const uint32_t* leaves, uint32_t nLeaves, uint32_t nItems);
void igneum_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw);
}
#ifdef IGNEUM_EMU_TWO_PACKS
namespace emu_pack_b {
struct IgneumInitWords { uint32_t w[8]; };
void igneum_cache_fill(uint32_t* cache, uint32_t nSegments);
void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems);
EMU_WEAK void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems);
EMU_WEAK void igneum_build(uint32_t* ds, const uint32_t* cache, const uint32_t* leaves, uint32_t nLeaves, uint32_t nItems);
void igneum_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw);
}
#endif
typedef void (*EmuBuild3)(uint32_t*, const uint32_t*, uint32_t);
typedef void (*EmuBuild5)(uint32_t*, const uint32_t*, const uint32_t*, uint32_t, uint32_t);
static std::string readAll(const std::string& path) {
FILE* f = std::fopen(path.c_str(), "rb");
@ -198,6 +206,7 @@ static CUresult d_memInfo(size_t* f, size_t* t) { *f = 8ull << 30; *t = 16ull <<
static CUresult d_alloc(CUdeviceptr* p, size_t n) { void* m = std::malloc(n); if (!m) return CUDA_ERROR_OUT_OF_MEMORY; *p = (CUdeviceptr)(uintptr_t)m; return CUDA_SUCCESS; }
static CUresult d_free(CUdeviceptr p) { std::free((void*)(uintptr_t)p); return CUDA_SUCCESS; }
static CUresult d_dtoh(void* dst, CUdeviceptr src, size_t n) { std::memcpy(dst, (const void*)(uintptr_t)src, n); return CUDA_SUCCESS; }
static CUresult d_htod(CUdeviceptr dst, const void* src, size_t n) { std::memcpy((void*)(uintptr_t)dst, src, n); return CUDA_SUCCESS; }
static CUresult d_modLoad(CUmodule* m, const void* img) {
const char* s = (const char*)img;
if (std::strncmp(s, "EMU-IMAGE:", 10) != 0) return CUDA_ERROR_INVALID_IMAGE;
@ -222,11 +231,28 @@ template <class T> static T* dptr(void** params, int i) { return (T*)(uintptr_t)
static CUresult d_launch(CUfunction f, unsigned gx, unsigned, unsigned, unsigned bx, unsigned, unsigned, unsigned, CUstream, void** params, void**) {
switch ((int)(uintptr_t)f) {
case 1: emu_launch(emu_pack_a::igneum_cache_fill, gx, bx, dptr<uint32_t>(params, 0), arg<uint32_t>(params, 1)); return CUDA_SUCCESS;
case 2: emu_launch(emu_pack_a::igneum_build, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), arg<uint32_t>(params, 2)); return CUDA_SUCCESS;
case 2: {
// the worker passes 5 arguments for a class v5 pack (params[3] is the leaf count, params[4] the item count) and 3 otherwise;
// the pack's kernel.cu defines exactly one shape: a mismatch (a v5 pack built without its leaves, or leaves handed to a
// v4 kernel) is the known-failed case and fails the launch instead of running the wrong kernel
EmuBuild5 b5 = (EmuBuild5)emu_pack_a::igneum_build; EmuBuild3 b3 = (EmuBuild3)emu_pack_a::igneum_build;
bool five = params[3] != nullptr && params[4] != nullptr;
if (five && b5) { emu_launch(b5, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), dptr<const uint32_t>(params, 2), arg<uint32_t>(params, 3), arg<uint32_t>(params, 4)); return CUDA_SUCCESS; }
if (!five && b3) { emu_launch(b3, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), arg<uint32_t>(params, 2)); return CUDA_SUCCESS; }
std::fprintf(stderr, "emu: igneum_build called with %d arguments but pack A's kernel has the %s shape\n", five ? 5 : 3, b5 ? "class v5 (leaves)" : "class v2 to v4");
return CUDA_ERROR_INVALID_VALUE;
}
case 3: emu_launch(emu_pack_a::igneum_hash_bound, gx, bx, dptr<const uint32_t>(params, 0), dptr<uint64_t>(params, 1), arg<uint32_t>(params, 2), arg<uint32_t>(params, 3), arg<emu_pack_a::IgneumInitWords>(params, 4)); return CUDA_SUCCESS;
#ifdef IGNEUM_EMU_TWO_PACKS
case 4: emu_launch(emu_pack_b::igneum_cache_fill, gx, bx, dptr<uint32_t>(params, 0), arg<uint32_t>(params, 1)); return CUDA_SUCCESS;
case 5: emu_launch(emu_pack_b::igneum_build, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), arg<uint32_t>(params, 2)); return CUDA_SUCCESS;
case 5: {
EmuBuild5 b5 = (EmuBuild5)emu_pack_b::igneum_build; EmuBuild3 b3 = (EmuBuild3)emu_pack_b::igneum_build;
bool five = params[3] != nullptr && params[4] != nullptr;
if (five && b5) { emu_launch(b5, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), dptr<const uint32_t>(params, 2), arg<uint32_t>(params, 3), arg<uint32_t>(params, 4)); return CUDA_SUCCESS; }
if (!five && b3) { emu_launch(b3, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), arg<uint32_t>(params, 2)); return CUDA_SUCCESS; }
std::fprintf(stderr, "emu: igneum_build called with %d arguments but pack B's kernel has the %s shape\n", five ? 5 : 3, b5 ? "class v5 (leaves)" : "class v2 to v4");
return CUDA_ERROR_INVALID_VALUE;
}
case 6: emu_launch(emu_pack_b::igneum_hash_bound, gx, bx, dptr<const uint32_t>(params, 0), dptr<uint64_t>(params, 1), arg<uint32_t>(params, 2), arg<uint32_t>(params, 3), arg<emu_pack_b::IgneumInitWords>(params, 4)); return CUDA_SUCCESS;
#endif
default: return CUDA_ERROR_INVALID_HANDLE;
@ -244,7 +270,7 @@ void emu_fill_driver(Drv& d) {
d.init = d_init; d.driverGetVersion = d_driverVersion; d.deviceGetCount = d_count; d.deviceGet = d_get; d.deviceGetName = d_name;
d.deviceGetAttribute = d_attr; d.deviceTotalMem = d_totalMem; d.primaryCtxSetFlags = d_ctxFlags; d.primaryCtxRetain = d_ctxRetain;
d.primaryCtxRelease = d_ctxRelease; d.ctxSetCurrent = d_ctxSet; d.ctxSynchronize = d_ctxSync; d.memGetInfo = d_memInfo;
d.memAlloc = d_alloc; d.memFree = d_free; d.memcpyDtoH = d_dtoh; d.moduleLoadData = d_modLoad; d.moduleUnload = d_modUnload;
d.memAlloc = d_alloc; d.memFree = d_free; d.memcpyDtoH = d_dtoh; d.memcpyHtoD = d_htod; d.moduleLoadData = d_modLoad; d.moduleUnload = d_modUnload;
d.moduleGetFunction = d_getFn; d.launchKernel = d_launch; d.streamCreate = d_streamCreate; d.streamSynchronize = d_streamSync;
d.streamDestroy = d_streamDestroy; d.funcGetAttribute = d_funcAttr; d.occupancy = d_occ; d.getErrorString = d_errStr; d.getErrorName = d_errName;
}

66
proto-cuda/nvrtc/emu/packfile-test.c
View file

@ -172,6 +172,72 @@ int main(int argc, char** argv) {
CHECK(pf_load(dir, &pk, err, sizeof(err)) == 0 && strstr(err, "generator 1 is not") != NULL, "a pack with no generator line (generator 1) is refused");
}
// Class v5 (docs/design/class-v5-stored-state.md, 7 October 2026): the state leaves. Known-failed first: a generator 5
// pack without the leaf count, one whose leaves file is missing, one whose file has another FNV, one of another class
// carrying a count; then the checked-in v5 pack (argv[2]) loads and its leaves load with the pack's FNV.
{
char sw[200], kw[200], text[2600], why[256], path[1024];
uint32_t* leaves = NULL;
size_t bytes = 0;
FILE* f;
int i;
words_hex(bare, sw); words_hex(keyw, kw);
snprintf(text, sizeof(text), "#define IGNEUM_SEED_BYTES_HEX \"%s\"\n#define IGNEUM_DAY_BYTES_HEX \"%s\"\n#define IGNEUM_GENERATOR 5\n#define IGNEUM_PROGRAM_CLASS \"v5\"\n#define IGNEUM_ERA_SEED_HEX \"%s\"\n#define IGNEUM_SHADOW_INSTRS 256\n#define IGNEUM_SHADOW_REPS 27\n#define IGNEUM_DATASET_LOG2 28\n#define IGNEUM_DATASET_MODE 1\n#define IGNEUM_SEEDW_INIT { %s }\n#define IGNEUM_KEY_INIT { %s }\n#define IGNEUM_CACHE_LOG2_WORDS 26\n#define IGNEUM_CACHE_SEGMENTS 4096u\n", EPOCH_34, DAY_20731, EPOCH_33, sw, kw);
write_file(dir, "program.h", text);
err[0] = 0;
CHECK(pf_load(dir, &pk, err, sizeof(err)) == 0 && strstr(err, "without IGNEUM_STATE_LEAVES") != NULL, "known-failed: a generator 5 pack without the leaf count is refused in plain words");
// the count and the FNV of two leaves of 64 bytes (0x00..0x3f, 0x40..0x7f), the file not yet written
{
uint8_t raw[128];
uint64_t fnv;
for (i = 0; i < 128; ++i) raw[i] = (uint8_t)i;
fnv = pf_fnv1a64(raw, sizeof(raw));
snprintf(text + strlen(text), sizeof(text) - strlen(text), "#define IGNEUM_STATE_LEAVES 2\n#define IGNEUM_STATE_LEAVES_FNV64 0x%016llxull\n#define IGNEUM_STATE_ROOT_HEX \"%s\"\n", (unsigned long long)fnv, EPOCH_33);
write_file(dir, "program.h", text);
snprintf(path, sizeof(path), "%s/leaves.bin", dir); remove(path);
err[0] = 0;
CHECK(pf_load(dir, &pk, err, sizeof(err)) == 1 && pk.stateLeaves == 2 && pk.stateLeavesFnv == fnv && strcmp(pk.programClass, "v5") == 0 && strcmp(pk.stateLeavesFile, "leaves.bin") == 0, "a generator 5 pack with the leaf count loads as class v5 (count, FNV and file name read back)");
CHECK(pf_pack_class_ok(pk.programClass, pk.eraHex, "v5", EPOCH_33, why, sizeof(why)) == 1, "the v5 pack matches a job naming class v5 and its era");
CHECK(pf_pack_class_ok(pk.programClass, pk.eraHex, "v4", EPOCH_33, why, sizeof(why)) == 0 && strstr(why, "program class mismatch") == why, "a job naming class v4 refuses the v5 pack");
err[0] = 0;
CHECK(pf_load_leaves(dir, &pk, &leaves, &bytes, err, sizeof(err)) == 0 && leaves == NULL && strstr(err, "without its leaves file") != NULL, "known-failed: the leaves file missing refuses the build in plain words");
f = fopen(path, "wb"); fwrite(raw, 1, 64, f); fclose(f);
err[0] = 0;
CHECK(pf_load_leaves(dir, &pk, &leaves, &bytes, err, sizeof(err)) == 0 && leaves == NULL && strstr(err, "is 64 bytes, the pack says 2 leaves of 64") != NULL, "known-failed: a short leaves file is refused with both sizes named");
raw[5] ^= 0x80;
f = fopen(path, "wb"); fwrite(raw, 1, 128, f); fclose(f);
raw[5] ^= 0x80;
err[0] = 0;
CHECK(pf_load_leaves(dir, &pk, &leaves, &bytes, err, sizeof(err)) == 0 && leaves == NULL && strstr(err, "is not the pack's IGNEUM_STATE_LEAVES_FNV64") != NULL, "known-failed: leaves of another root (one bit) are refused by the FNV");
f = fopen(path, "wb"); fwrite(raw, 1, 128, f); fclose(f);
err[0] = 0;
CHECK(pf_load_leaves(dir, &pk, &leaves, &bytes, err, sizeof(err)) == 1 && leaves != NULL && bytes == 128 && memcmp(leaves, raw, 128) == 0, "known-good: the right leaves load as 2 x 16 words");
free(leaves); leaves = NULL;
remove(path);
}
// a class v4 pack carrying a leaf count is a mis-stamped export
snprintf(text, sizeof(text), "#define IGNEUM_SEED_BYTES_HEX \"%s\"\n#define IGNEUM_DAY_BYTES_HEX \"%s\"\n#define IGNEUM_GENERATOR 4\n#define IGNEUM_PROGRAM_CLASS \"v4\"\n#define IGNEUM_ERA_SEED_HEX \"%s\"\n#define IGNEUM_SHADOW_INSTRS 256\n#define IGNEUM_STATE_LEAVES 2\n#define IGNEUM_DATASET_LOG2 28\n#define IGNEUM_DATASET_MODE 1\n#define IGNEUM_SEEDW_INIT { %s }\n#define IGNEUM_KEY_INIT { %s }\n#define IGNEUM_CACHE_LOG2_WORDS 26\n#define IGNEUM_CACHE_SEGMENTS 4096u\n", EPOCH_34, DAY_20731, EPOCH_33, sw, kw);
write_file(dir, "program.h", text);
err[0] = 0;
CHECK(pf_load(dir, &pk, err, sizeof(err)) == 0 && strstr(err, "state leaves belong to class v5") != NULL, "known-failed: a generator 4 pack carrying IGNEUM_STATE_LEAVES is refused");
// a class v4 pack without leaves loads with stateLeaves 0 and pf_load_leaves hands back nothing
snprintf(text, sizeof(text), "#define IGNEUM_SEED_BYTES_HEX \"%s\"\n#define IGNEUM_DAY_BYTES_HEX \"%s\"\n#define IGNEUM_GENERATOR 4\n#define IGNEUM_PROGRAM_CLASS \"v4\"\n#define IGNEUM_ERA_SEED_HEX \"%s\"\n#define IGNEUM_SHADOW_INSTRS 256\n#define IGNEUM_DATASET_LOG2 28\n#define IGNEUM_DATASET_MODE 1\n#define IGNEUM_SEEDW_INIT { %s }\n#define IGNEUM_KEY_INIT { %s }\n#define IGNEUM_CACHE_LOG2_WORDS 26\n#define IGNEUM_CACHE_SEGMENTS 4096u\n", EPOCH_34, DAY_20731, EPOCH_33, sw, kw);
write_file(dir, "program.h", text);
err[0] = 0;
CHECK(pf_load(dir, &pk, err, sizeof(err)) == 1 && pk.stateLeaves == 0 && pf_load_leaves(dir, &pk, &leaves, &bytes, err, sizeof(err)) == 1 && leaves == NULL && bytes == 0, "a class v4 pack has no leaves and the leaf loader hands back none");
if (argc > 2) {
// the checked-in class v5 pack (proto-cuda/packs-ca3-v5/v5-dn3-epoch0): loads, its leaves load under the pack's FNV
err[0] = 0;
CHECK(pf_load(argv[2], &pk, err, sizeof(err)) == 1 && strcmp(pk.programClass, "v5") == 0 && pk.generator == 5 && pk.stateLeaves > 0 && pk.haveVectors, "known-good: the checked-in class v5 pack loads with its leaf count and vectors");
if (err[0]) printf(" %s\n", err);
err[0] = 0;
CHECK(pf_load_leaves(argv[2], &pk, &leaves, &bytes, err, sizeof(err)) == 1 && leaves != NULL && bytes == (size_t)pk.stateLeaves * 64u, "known-good: the checked-in pack's leaves.bin loads under the pack's FNV-1a 64");
if (err[0]) printf(" %s\n", err);
printf(" v5 pack: %u leaves, FNV-1a 64 %016llx, state root %s\n", (unsigned)pk.stateLeaves, (unsigned long long)pk.stateLeavesFnv, pk.stateRootHex);
free(leaves); leaves = NULL;
}
}
printf("%s: %d failure(s)\n", argv[0], failures);
return failures ? 1 : 0;
}

6
proto-cuda/nvrtc/emu/packfile-test.sh
View file

@ -1,6 +1,6 @@
#!/usr/bin/env bash
# The pack loader's seed rule (packfile.h) on a known-good and a known-mismatched pack: emu/packfile-test.c, C99,
# no GPU. Runs on the Mac in a second and in CI. Usage: emu/packfile-test.sh
# The pack loader's seed rule (packfile.h) on a known-good and a known-mismatched pack, and the class v5 leaf rule on the
# checked-in v5 pack (7 October 2026): emu/packfile-test.c, C99, no GPU. Runs on the Mac in a second and in CI. Usage: emu/packfile-test.sh
set -euo pipefail
HERE="$(cd "$(dirname "$0")" && pwd)"
ROOT="$(cd "$HERE/../../.." && pwd)"
@ -8,4 +8,4 @@ OUT="${TMPDIR:-/tmp}/igneum-packfile-test"
mkdir -p "$OUT"
CC="${CC:-cc}"
"$CC" -std=c99 -Wall -Wextra -Wno-unused-function -O1 -o "$OUT/packfile-test" "$HERE/packfile-test.c"
"$OUT/packfile-test" "$ROOT/proto-cuda/packs/igneum-devnet-v4-epoch0"
"$OUT/packfile-test" "$ROOT/proto-cuda/packs/igneum-devnet-v4-epoch0" "$ROOT/proto-cuda/packs-ca3-v5/v5-dn3-epoch0"

74
proto-cuda/nvrtc/packfile.h
View file

@ -31,6 +31,15 @@ typedef struct {
char loadClass[64];
char programClass[8]; /* IGNEUM_PROGRAM_CLASS: "v2", "v3" (Counter ASIC 2.0) or "v4" (Counter ASIC 3.0); absent = the generator's class */
char eraHex[65]; /* IGNEUM_ERA_SEED_HEX of a class v3 or v4 chain pack; empty otherwise */
/* Class v5 (docs/design/class-v5-stored-state.md, 7 October 2026): the window's state leaves. leaves.bin beside program.h
* holds IGNEUM_STATE_LEAVES leaves of 16 little-endian words (64 B each), leaf(t) = leaves[t mod stateLeaves]; every item
* of the dataset is keyed by one, so a host that has not uploaded them builds nothing (pf_load_leaves). stateLeaves is 0
* for every other class; a generator 5 pack without them, or a pack of another class with them, is refused by pf_load. */
uint32_t stateLeaves;
uint64_t stateLeavesFnv; /* IGNEUM_STATE_LEAVES_FNV64: FNV-1a 64 over leaves.bin's bytes */
char stateLeavesFile[64]; /* IGNEUM_STATE_LEAVES_FILE, "leaves.bin" when absent */
char stateRootHex[65]; /* IGNEUM_STATE_ROOT_HEX: the state root the leaves hash under */
char stateBlockHex[65]; /* IGNEUM_STATE_BLOCK_HEX: the window's reference chain block */
// Counter ASIC 2.0 (5 October 2026): the mixer multiplier of the item derivation (IGNEUM_MIXER_MULT, 1 when absent:
// version 2; 4 under class v3). The emitted memhard.h / kernel.cl carry it in their text; this is for the log lines.
uint32_t mixerMult;
@ -117,6 +126,15 @@ static int pf_define_u32(const char* text, const char* name, uint32_t* out) {
return 1;
}
/* A 64-bit define (`0x...ull`), the pack's FNV lines in program.h. */
static int pf_define_u64(const char* text, const char* name, uint64_t* out) {
const char* p = pf_find_define(text, name);
const char* e;
if (!p) return 0;
while (*p == ' ' || *p == '\t') ++p;
return pf_number(p, out, &e);
}
static int pf_define_str(const char* text, const char* name, char* out, size_t cap) {
const char* p = pf_find_define(text, name);
const char* q;
@ -286,12 +304,14 @@ static int pf_load(const char* dir, PfPack* pk, char* err, size_t cap) {
/* Spec 01 section 1.4.5: a pack whose generator version is not one this worker runs is refused. Generator 2 is
* program class v2 (the lottery hash of 4 October 2026), generator 3 is class v3 (Counter ASIC 2.0), generator 4
* is class v4 (Counter ASIC 3.0, 6 October 2026: class v3 plus the latency-shadow block, emitted in the pack's own
* kernel text, so this loader needs nothing new beyond the number and the class token). */
if (pk->generator != 2 && pk->generator != 3 && pk->generator != 4) {
char m[200]; snprintf(m, sizeof(m), "program pack generator %u is not a generator version this worker runs (2, 3 or 4)", (unsigned)pk->generator);
* kernel text, so this loader needs nothing new beyond the number and the class token), generator 5 is class v5
* (proof of stored state, 7 October 2026: class v4 over a dataset keyed by the window's state leaves, leaves.bin,
* which the host uploads for igneum_build; the kernel text carries the leaf read). */
if (pk->generator != 2 && pk->generator != 3 && pk->generator != 4 && pk->generator != 5) {
char m[200]; snprintf(m, sizeof(m), "program pack generator %u is not a generator version this worker runs (2, 3, 4 or 5)", (unsigned)pk->generator);
free(prog); return pf_fail(err, cap, m);
}
strcpy(pk->programClass, pk->generator == 4 ? "v4" : pk->generator == 3 ? "v3" : "v2");
strcpy(pk->programClass, pk->generator == 5 ? "v5" : pk->generator == 4 ? "v4" : pk->generator == 3 ? "v3" : "v2");
{
/* Counter ASIC 3.0 (6 October 2026): the shadow block marks class v4. A generator 3 pack with IGNEUM_SHADOW_INSTRS
* is a v4 program stamped as v3 (the old export path; it carried the v3 control's program id) and is refused;
@ -299,6 +319,7 @@ static int pf_load(const char* dir, PfPack* pk, char* err, size_t cap) {
uint32_t shadow = 0;
if (!pf_define_u32(prog, "IGNEUM_SHADOW_INSTRS", &shadow)) shadow = 0;
if (pk->generator == 4 && shadow == 0) { free(prog); return pf_fail(err, cap, "program pack generator 4 (class v4) without IGNEUM_SHADOW_INSTRS: not a class v4 pack"); }
if (pk->generator == 5 && shadow == 0) { free(prog); return pf_fail(err, cap, "program pack generator 5 (class v5) without IGNEUM_SHADOW_INSTRS: not a class v5 pack (class v5 is class v4 over the state leaves)"); }
if (pk->generator == 3 && shadow != 0) { /* a generator 2 pack with a shadow is the measurement ladder (a class-bearing id) and loads */
char m[220]; snprintf(m, sizeof(m), "program pack generator %u with a shadow block (IGNEUM_SHADOW_INSTRS %u): a class v4 program is generator 4 (export the pack as class v4)", (unsigned)pk->generator, (unsigned)shadow);
free(prog); return pf_fail(err, cap, m);
@ -312,6 +333,24 @@ static int pf_load(const char* dir, PfPack* pk, char* err, size_t cap) {
}
}
pk->eraHex[0] = 0; pf_define_str(prog, "IGNEUM_ERA_SEED_HEX", pk->eraHex, sizeof(pk->eraHex));
/* Class v5: the state leaves. The count, the FNV and the file name come from program.h; the bytes are read by
* pf_load_leaves when the host builds. A v5 pack without the count (or with 0) is refused: nothing could key its items;
* a pack of another class carrying a count is a mis-stamped export and is refused too. */
pk->stateLeaves = 0; pk->stateLeavesFnv = 0; pk->stateRootHex[0] = 0; pk->stateBlockHex[0] = 0;
strcpy(pk->stateLeavesFile, "leaves.bin");
if (!pf_define_u32(prog, "IGNEUM_STATE_LEAVES", &pk->stateLeaves)) pk->stateLeaves = 0;
if (pk->generator == 5 && pk->stateLeaves == 0) { free(prog); return pf_fail(err, cap, "program pack generator 5 (class v5) without IGNEUM_STATE_LEAVES: no state leaves to key the dataset (export the pack with --state)"); }
if (pk->generator != 5 && pk->stateLeaves != 0) {
char m[200]; snprintf(m, sizeof(m), "program pack generator %u carries IGNEUM_STATE_LEAVES %u: state leaves belong to class v5 (generator 5)", (unsigned)pk->generator, (unsigned)pk->stateLeaves);
free(prog); return pf_fail(err, cap, m);
}
if (pk->stateLeaves) {
if (pk->stateLeaves > 0x02000000u) { free(prog); return pf_fail(err, cap, "program.h IGNEUM_STATE_LEAVES is above 2^25 (the designed dataset's item cap)"); }
if (!pf_define_u64(prog, "IGNEUM_STATE_LEAVES_FNV64", &pk->stateLeavesFnv)) { free(prog); return pf_fail(err, cap, "program pack class v5 without IGNEUM_STATE_LEAVES_FNV64: the leaves cannot be checked"); }
pf_define_str(prog, "IGNEUM_STATE_LEAVES_FILE", pk->stateLeavesFile, sizeof(pk->stateLeavesFile));
pf_define_str(prog, "IGNEUM_STATE_ROOT_HEX", pk->stateRootHex, sizeof(pk->stateRootHex));
pf_define_str(prog, "IGNEUM_STATE_BLOCK_HEX", pk->stateBlockHex, sizeof(pk->stateBlockHex));
}
if (!pf_define_u32(prog, "IGNEUM_PROGRAM_ATTEMPT", &pk->attempt)) pk->attempt = 0;
if (pf_define_words(prog, "IGNEUM_SEEDW_INIT", pk->seedw, 8) != 8) { free(prog); return pf_fail(err, cap, "program.h has no IGNEUM_SEEDW_INIT with 8 words"); }
pk->loadsPerHash = 128; pf_define_u32(prog, "IGNEUM_LOADS_PER_HASH", &pk->loadsPerHash);
@ -405,6 +444,31 @@ static int pf_load(const char* dir, PfPack* pk, char* err, size_t cap) {
return 1;
}
/* Class v5: reads the pack's leaves file into a malloc'd buffer of stateLeaves x 64 bytes (16 little-endian words per leaf,
* the layout igneum_build reads), checked against the pack's count (the byte length) and its FNV-1a 64
* (IGNEUM_STATE_LEAVES_FNV64) before anything is uploaded. Returns 1 with *out and *bytes set (NULL and 0 for a pack of
* another class, which has no leaves); 0 with err and nothing allocated. The caller frees *out after the build. */
static int pf_load_leaves(const char* dir, const PfPack* pk, uint32_t** out, size_t* bytes, char* err, size_t cap) {
char* raw;
size_t n = 0;
uint64_t fnv;
*out = NULL; *bytes = 0;
if (pk->stateLeaves == 0) return 1;
raw = pf_read_pack_file(dir, pk->stateLeavesFile, &n);
if (!raw) { char m[600]; snprintf(m, sizeof(m), "class v5 pack without its leaves file %.400s/%.60s (the state leaves key every item: nothing can be built without them)", dir, pk->stateLeavesFile); return pf_fail(err, cap, m); }
if (n != (size_t)pk->stateLeaves * 64u) {
char m[300]; snprintf(m, sizeof(m), "%.60s is %llu bytes, the pack says %u leaves of 64 (%llu bytes)", pk->stateLeavesFile, (unsigned long long)n, (unsigned)pk->stateLeaves, (unsigned long long)pk->stateLeaves * 64ull);
free(raw); return pf_fail(err, cap, m);
}
fnv = pf_fnv1a64(raw, n);
if (fnv != pk->stateLeavesFnv) {
char m[300]; snprintf(m, sizeof(m), "%.60s FNV-1a 64 %016llx is not the pack's IGNEUM_STATE_LEAVES_FNV64 %016llx (the leaves are of another state root; export the pack again)", pk->stateLeavesFile, (unsigned long long)fnv, (unsigned long long)pk->stateLeavesFnv);
free(raw); return pf_fail(err, cap, m);
}
*out = (uint32_t*)raw; *bytes = n;
return 1;
}
// The self-test verdict from values the host read back from the device. `vec` holds vecWarps x 32 outputs of the
// bound kernel run with the pack's own seed words as init words (that is igneum_hash of kernel.cu). Writes one line.
// A hot-table pack (pk->hotMb) also hands the hot table's head, last line and FNV-1a 64 (NULL and 0 otherwise); a
@ -443,7 +507,7 @@ static int pf_selftest(const PfPack* pk, const uint32_t* cacheHead, const uint32
}
/* Counter ASIC 2.0 (5 October 2026): a job or prepare line may end with `class=<v2|v3|v4>` and `era=<hex>` tokens (sent
/* Counter ASIC 2.0 (5 October 2026): a job or prepare line may end with `class=<v2|v3|v4|v5>` and `era=<hex>` tokens (sent
* only when the chain is on class v3 or v4, so every v2 line is the line of before). A pack matches the line when its
* class is the named class and, when an era is named, its era seed is that era (every class after v2 carries one).
* Empty wanted strings accept any pack. Returns 1 on a match, else 0 with the reason in `why`. */

52
proto-cuda/nvrtc/worker.cpp
View file

@ -166,6 +166,7 @@ static bool loadDriver(Drv& d, std::string& err, std::string& libName) {
LOAD_SYM(d, memAlloc, "cuMemAlloc_v2");
LOAD_SYM(d, memFree, "cuMemFree_v2");
LOAD_SYM(d, memcpyDtoH, "cuMemcpyDtoH_v2");
LOAD_SYM(d, memcpyHtoD, "cuMemcpyHtoD_v2");
LOAD_SYM(d, moduleLoadData, "cuModuleLoadData");
LOAD_SYM(d, moduleUnload, "cuModuleUnload");
LOAD_SYM(d, moduleGetFunction, "cuModuleGetFunction");
@ -438,6 +439,8 @@ struct Pair {
// pack's igneum_hot_fill, the argument after the init words
uint32_t hotMb = 0, hotWords = 0, hotSegments = 0, hotSlots = 0;
CUfunction fHotFill = nullptr;
// class v5 (7 October 2026): the pack's state leaves, uploaded for igneum_build and freed after it (0 for other classes)
uint32_t stateLeaves = 0;
CUdeviceptr hot = 0;
double hotMs = 0;
};
@ -841,7 +844,14 @@ static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string&
}
p->loadClass = pk.loadClass; p->loadsPerHash = pk.loadsPerHash; p->bytesPerHash = pk.bytesPerHash; p->scratchOps = pk.scratchOps;
p->programClass = pk.programClass; p->eraHex = pk.eraHex;
p->stateLeaves = pk.stateLeaves;
p->persistent = pk.persistent != 0;
// Class v5 (docs/design/class-v5-stored-state.md): the window's leaves (leaves.bin), read and checked against the pack's
// count and FNV-1a 64 before anything is allocated; uploaded for igneum_build below and freed right after it, so device
// memory while hashing is the class v4 worker's. A v5 pack without its leaves builds nothing (the known-failed case).
uint32_t* hLeaves = nullptr;
size_t leavesBytes = 0;
{ char lerr[700]; if (!pf_load_leaves(dir.c_str(), &pk, &hLeaves, &leavesBytes, lerr, sizeof(lerr))) { err = "pack " + dir + ": " + lerr; releasePair(c, p); return nullptr; } }
p->hotMb = pk.hotMb; p->hotWords = pk.hotWords; p->hotSegments = pk.hotSegments; p->hotSlots = pk.hotSlots;
p->residentWarps = p->blocksPerSM * c.blockWarps * c.sms;
size_t scratchBytes = 0, hotBytes = (size_t)pk.hotWords * 4u;
@ -860,14 +870,14 @@ static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string&
size_t cacheBytes = (size_t)p->cacheWords * 4u, dsBytes = (size_t)p->words * 4u;
{
size_t freeB = 0, totalB = 0;
if (c.drv.memGetInfo(&freeB, &totalB) == CUDA_SUCCESS && freeB < cacheBytes + dsBytes + scratchBytes + hotBytes + (64u << 20)) {
err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu + scratch %llu + hot %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes + scratchBytes + hotBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20), (unsigned long long)(scratchBytes >> 20), (unsigned long long)(hotBytes >> 20));
releasePair(c, p); return nullptr;
if (c.drv.memGetInfo(&freeB, &totalB) == CUDA_SUCCESS && freeB < cacheBytes + dsBytes + scratchBytes + hotBytes + leavesBytes + (64u << 20)) {
err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu + scratch %llu + hot %llu + leaves %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes + scratchBytes + hotBytes + leavesBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20), (unsigned long long)(scratchBytes >> 20), (unsigned long long)(hotBytes >> 20), (unsigned long long)(leavesBytes >> 20));
std::free(hLeaves); releasePair(c, p); return nullptr;
}
}
if (p->persistent) {
CUresult r = c.drv.memAlloc(&p->scratch, scratchBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc scratch: " + c.err(r); p->scratch = 0; releasePair(c, p); return nullptr; }
if (r != CUDA_SUCCESS) { err = "cuMemAlloc scratch: " + c.err(r); p->scratch = 0; std::free(hLeaves); releasePair(c, p); return nullptr; }
p->scratchBytes = scratchBytes;
int after = 0;
c.drv.occupancy(&after, p->fHashBound, 32 * c.blockWarps, 0);
@ -876,24 +886,41 @@ static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string&
}
{
CUresult r = c.drv.memAlloc(&p->cache, cacheBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc cache: " + c.err(r); p->cache = 0; releasePair(c, p); return nullptr; }
if (r != CUDA_SUCCESS) { err = "cuMemAlloc cache: " + c.err(r); p->cache = 0; std::free(hLeaves); releasePair(c, p); return nullptr; }
uint32_t nSeg = p->cacheSegments, block = 256u, grid = (nSeg + block - 1u) / block;
void* args[2] = { &p->cache, &nSeg };
r = c.drv.launchKernel(p->fCacheFill, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
if (r != CUDA_SUCCESS) { err = "cache fill: " + c.err(r); releasePair(c, p); return nullptr; }
if (r != CUDA_SUCCESS) { err = "cache fill: " + c.err(r); std::free(hLeaves); releasePair(c, p); return nullptr; }
}
p->cacheMs = wallMs() - t0;
// Dataset
t0 = wallMs();
{
CUresult r = c.drv.memAlloc(&p->ds, dsBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc dataset: " + c.err(r); p->ds = 0; releasePair(c, p); return nullptr; }
if (r != CUDA_SUCCESS) { err = "cuMemAlloc dataset: " + c.err(r); p->ds = 0; std::free(hLeaves); releasePair(c, p); return nullptr; }
uint32_t nItems = p->words / 16u, block = 256u, grid = (nItems + block - 1u) / block;
void* args[3] = { &p->ds, &p->cache, &nItems };
r = c.drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
if (r != CUDA_SUCCESS) { err = "dataset build: " + c.err(r); releasePair(c, p); return nullptr; }
if (hLeaves) {
// class v5: igneum_build(ds, cache, leaves, nLeaves, nItems), the leaf buffer freed once the build has run
CUdeviceptr dLeaves = 0;
uint32_t nLeaves = pk.stateLeaves;
r = c.drv.memAlloc(&dLeaves, leavesBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc state leaves: " + c.err(r); std::free(hLeaves); releasePair(c, p); return nullptr; }
r = c.drv.memcpyHtoD(dLeaves, hLeaves, leavesBytes);
std::free(hLeaves); hLeaves = nullptr;
if (r == CUDA_SUCCESS) {
void* args[5] = { &p->ds, &p->cache, &dLeaves, &nLeaves, &nItems };
r = c.drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
}
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
c.drv.memFree(dLeaves);
if (r != CUDA_SUCCESS) { err = "dataset build (class v5, " + std::to_string(nLeaves) + " state leaves): " + c.err(r); releasePair(c, p); return nullptr; }
} else {
void* args[5] = { &p->ds, &p->cache, &nItems, nullptr, nullptr }; // five slots: the driver reads the kernel's three, the emulation reads the shape from the two null tails
r = c.drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
if (r != CUDA_SUCCESS) { err = "dataset build: " + c.err(r); releasePair(c, p); return nullptr; }
}
}
p->dsMs = wallMs() - t0;
// Hot table (hot-table experiment): filled from the epoch seed by the pack's own kernel, never shipped
@ -958,7 +985,8 @@ static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string&
}
static std::string pairSummary(const Pair* p) {
return fmt("nvrtc %.0f cache %.0f dataset %.0f hot %.0f check %.0f race %.0f ms variant %s; %s", p->compileMs, p->cacheMs, p->dsMs, p->hotMs, p->checkMs, p->raceMs, p->variant.c_str(), p->check.c_str());
return fmt("nvrtc %.0f cache %.0f dataset %.0f hot %.0f check %.0f race %.0f ms variant %s class %s%s; %s", p->compileMs, p->cacheMs, p->dsMs, p->hotMs, p->checkMs, p->raceMs, p->variant.c_str(), p->programClass.c_str(),
p->stateLeaves ? fmt(" (state leaves %u, uploaded for the build and freed)", p->stateLeaves).c_str() : "", p->check.c_str());
}
// ---------------------------------------------------------------------------------------------

58
proto-metal/main.swift
View file

@ -2992,9 +2992,9 @@ final class ServeDataset {
/// Counter ASIC 2.0 and 3.0: the classes this worker runs from a prepared pack only (never from the Swift version 2
/// generator): class v3 (generator 3) and class v4 (generator 4, class v3 plus the latency-shadow block, which is in
/// the pack's own program_bound.metal). Each carries an era seed.
func isPackClass(_ cls: String) -> Bool { cls == "v3" || cls == "v4" }
func isPackClass(_ cls: String) -> Bool { cls == "v3" || cls == "v4" || cls == "v5" } // class v5 (7 October 2026): class v4 over the state leaves, from a pack
/// The class name of a pack's IGNEUM_GENERATOR (packfile.h's rule).
func packClassOf(generator: UInt32) -> String { generator == 4 ? "v4" : generator == 3 ? "v3" : "v2" }
func packClassOf(generator: UInt32) -> String { generator == 5 ? "v5" : generator == 4 ? "v4" : generator == 3 ? "v3" : "v2" }
// The resident programs and datasets, shared by the job loop (main thread) and the prepare queue (background).
final class ServeStore {
@ -3085,13 +3085,17 @@ func servePackProgram(_ gpu: GPU, _ store: ServeStore, seedHex: String, seed: [U
}
func refuse(_ why: String) -> NSError { NSError(domain: "pack", code: 2, userInfo: [NSLocalizedDescriptionKey: "pack \(dir): \(why)"]) }
let generator = defineU32("IGNEUM_GENERATOR") ?? 1
guard generator == 2 || generator == 3 || generator == 4 else { throw refuse("program pack generator \(generator) is not a generator version this worker runs (2, 3 or 4)") }
guard generator == 2 || generator == 3 || generator == 4 || generator == 5 else { throw refuse("program pack generator \(generator) is not a generator version this worker runs (2, 3, 4 or 5)") }
let packClass = packClassOf(generator: generator)
if let named = defineStr("IGNEUM_PROGRAM_CLASS"), named != packClass { throw refuse("program pack IGNEUM_PROGRAM_CLASS \"\(named)\" does not match IGNEUM_GENERATOR \(generator)") }
// Counter ASIC 3.0 (6 October 2026): the shadow block marks class v4 (packfile.h's rule): a generator 3 pack with
// IGNEUM_SHADOW_INSTRS is a v4 program stamped v3 (the v3 control's program id) and a generator 4 pack without it is no v4 pack
let shadow = defineU32("IGNEUM_SHADOW_INSTRS") ?? 0
if generator == 4 && shadow == 0 { throw refuse("program pack generator 4 (class v4) without IGNEUM_SHADOW_INSTRS: not a class v4 pack") }
// class v5 (7 October 2026) is class v4 over the state leaves: the shadow block and IGNEUM_STATE_LEAVES both mark it
if generator == 5 && shadow == 0 { throw refuse("program pack generator 5 (class v5) without IGNEUM_SHADOW_INSTRS: not a class v5 pack") }
if generator == 5 && (defineU32("IGNEUM_STATE_LEAVES") ?? 0) == 0 { throw refuse("program pack generator 5 (class v5) without IGNEUM_STATE_LEAVES: no state leaves to key the dataset (export the pack with --state)") }
if generator != 5 && (defineU32("IGNEUM_STATE_LEAVES") ?? 0) != 0 { throw refuse("program pack generator \(generator) carries IGNEUM_STATE_LEAVES: state leaves belong to class v5 (generator 5)") }
if generator == 3 && shadow != 0 { throw refuse("program pack generator \(generator) with a shadow block (IGNEUM_SHADOW_INSTRS \(shadow)): a class v4 program is generator 4 (export the pack as class v4)") }
let eraHex = defineStr("IGNEUM_ERA_SEED_HEX") ?? ""
if wantClass != "" && wantClass != packClass { throw refuse("program class mismatch: this pack is class \(packClass), the line names class \(wantClass) (export the pack again)") }
@ -3159,6 +3163,29 @@ func servePackDataset(_ gpu: GPU, _ store: ServeStore, dayHex: String, dir: Stri
let fillPipe = try gpu.device.makeComputePipelineState(function: fillFn)
let buildPipe = try gpu.device.makeComputePipelineState(function: buildFn)
let words = 1 << datasetLog2
// Class v5 (docs/design/class-v5-stored-state.md, 7 October 2026): the window's state leaves, leaves.bin beside program.h
// (IGNEUM_STATE_LEAVES leaves of 16 little-endian words), checked against the pack's count and FNV-1a 64
// (IGNEUM_STATE_LEAVES_FNV64) and bound as buffer 2 of igneum_build with the count in buffer 3 (packbench.swift's shape,
// the kernel text the Rust emitter wrote). A v5 pack without its leaves builds nothing (the known-failed case); the buffer
// is dropped after the build, so device memory while hashing is the class v4 worker's.
func defineU64(_ name: String) -> UInt64? {
guard let re = try? NSRegularExpression(pattern: "#define \(name) 0x([0-9a-fA-F]+)ull"), let m = re.firstMatch(in: programH, range: NSRange(programH.startIndex..., in: programH)) else { return nil }
return UInt64(programH[Range(m.range(at: 1), in: programH)!], radix: 16)
}
let stateLeaves = defineU32("IGNEUM_STATE_LEAVES") ?? 0
if (packClass == "v5") != (stateLeaves > 0) { throw refuse(packClass == "v5" ? "class v5 pack without IGNEUM_STATE_LEAVES" : "a class \(packClass) pack carries IGNEUM_STATE_LEAVES \(stateLeaves): state leaves belong to class v5") }
var leavesBuf: MTLBuffer? = nil
var nLeaves: UInt32 = 0
if stateLeaves > 0 {
let file = defineStr("IGNEUM_STATE_LEAVES_FILE") ?? "leaves.bin"
guard let data = FileManager.default.contents(atPath: dir + "/" + file) else { throw refuse("class v5 pack without its leaves file \(file) (the state leaves key every item: nothing can be built without them)") }
if data.count != Int(stateLeaves) * 64 { throw refuse("\(file) is \(data.count) bytes, the pack says \(stateLeaves) leaves of 64") }
guard let want = defineU64("IGNEUM_STATE_LEAVES_FNV64") else { throw refuse("class v5 pack without IGNEUM_STATE_LEAVES_FNV64: the leaves cannot be checked") }
let got = fnv1a64Bytes([UInt8](data))
if got != want { throw refuse("\(file) FNV-1a 64 \(h64(got)) is not the pack's IGNEUM_STATE_LEAVES_FNV64 \(h64(want)) (the leaves are of another state root; export the pack again)") }
guard let b = gpu.device.makeBuffer(bytes: (data as NSData).bytes, length: data.count, options: .storageModeShared) else { throw refuse("cannot allocate the \(data.count) byte leaf buffer") }
leavesBuf = b; nLeaves = stateLeaves
}
guard let cache = gpu.device.makeBuffer(length: (1 << cacheLog2) * 4, options: .storageModePrivate) else { throw refuse("cannot allocate the 2^\(cacheLog2) word cache") }
guard let dataset = gpu.device.makeBuffer(length: words * 4, options: .storageModePrivate) else { throw refuse("cannot allocate the 2^\(datasetLog2) word dataset") }
func run(_ body: (MTLComputeCommandEncoder) -> Void) throws -> Double {
@ -3174,8 +3201,29 @@ func servePackDataset(_ gpu: GPU, _ store: ServeStore, dayHex: String, dir: Stri
}
let buildMs = try run { enc in
enc.setComputePipelineState(buildPipe); enc.setBuffer(cache, offset: 0, index: 0); enc.setBuffer(dataset, offset: 0, index: 1)
if let lb = leavesBuf { enc.setBuffer(lb, offset: 0, index: 2); var n = nLeaves; enc.setBytes(&n, length: 4, index: 3) }
enc.dispatchThreadgroups(MTLSize(width: words / 16 / 256, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: 256, height: 1, depth: 1))
}
leavesBuf = nil
// The build's self-test when the pack carries vectors.json (igneum-pow export writes it): the dataset's head 16 words and
// its last word against the CPU reference, as the one-click workers' pf_selftest does; a mismatch refuses the day (the
// miner's CPU re-check would catch every found nonce later, but a wrong dataset is wrong on every hash, so it stops here)
if let vdata = FileManager.default.contents(atPath: dir + "/vectors.json"), let vj = (try? JSONSerialization.jsonObject(with: vdata)) as? [String: Any],
let headHex = vj["dataset_head"] as? [String], headHex.count == 16, let lastIdx = (vj["dataset_last_index"] as? NSNumber)?.uint64Value, let lastHex = vj["dataset_last"] as? String {
func hex32(_ x: String) -> UInt32? { UInt32(x.hasPrefix("0x") ? String(x.dropFirst(2)) : x, radix: 16) }
let wantHead = headHex.compactMap(hex32)
if wantHead.count == 16, let wantLast = hex32(lastHex), lastIdx < UInt64(words) {
let copy = gpu.device.makeBuffer(length: 64 + 4, options: .storageModeShared)!
let cb = gpu.queue.makeCommandBuffer()!; let bl = cb.makeBlitCommandEncoder()!
bl.copy(from: dataset, sourceOffset: 0, to: copy, destinationOffset: 0, size: 64)
bl.copy(from: dataset, sourceOffset: Int(lastIdx) * 4, to: copy, destinationOffset: 64, size: 4)
bl.endEncoding(); cb.commit(); cb.waitUntilCompleted()
let p = copy.contents().bindMemory(to: UInt32.self, capacity: 17)
let headOk = (0..<16).allSatisfy { p[$0] == wantHead[$0] }
let lastOk = p[16] == wantLast
if !headOk || !lastOk { throw refuse("dataset self-test FAIL against vectors.json (head \(headOk ? "ok" : "BAD"), word [\(lastIdx)] \(lastOk ? "ok" : "BAD")\(stateLeaves > 0 ? "; class v5 with \(stateLeaves) state leaves" : ""))") }
}
}
let sd = ServeDataset(dayHex: dayHex, packClass: packClass, eraHex: eraHex, buffer: dataset, cacheFillGPUms: fillMs, buildGPUms: buildMs)
store.add(sd)
return (sd, ms(t0, nowNs()))
@ -3206,8 +3254,8 @@ func runServe(_ opts: Options) -> Never {
}
if wantClass != "" && wantClass != "v2" && !isPackClass(wantClass) {
let id = f.count > 1 ? f[1] : "0"
if f[0] == "prepare" { emit("prepare-failed \(id) \(f.count > 2 ? f[2] : "0") program class \(wantClass) is not one this worker runs (v2, v3 or v4)") }
else { emit("error \(id) program class \(wantClass) is not one this worker runs (v2, v3 or v4)") }
if f[0] == "prepare" { emit("prepare-failed \(id) \(f.count > 2 ? f[2] : "0") program class \(wantClass) is not one this worker runs (v2, v3, v4 or v5)") }
else { emit("error \(id) program class \(wantClass) is not one this worker runs (v2, v3, v4 or v5)") }
continue
}
if f[0] == "prepare" {

134
proto-opencl/host.c
View file

@ -184,10 +184,19 @@ static uint64_t fnv1a64(const void* p, size_t n) {
}
static const uint32_t CACHE_WORDS_HOST = 1u << IGNEUM_CACHE_LOG2_WORDS;
static uint32_t* hCache = NULL;
#ifdef IGNEUM_STATE_LEAVES
/* class v5 compiled-in pack: the host copy of the pack's leaves (leaves.bin beside IGNEUM_KERNEL_PATH, loaded and checked by
* compiledInLeavesBuffer before the build), so the host reference derives the same words the device does */
static uint32_t* hLeaves = NULL;
#endif
/* dataset[w] through the pack's own mh_word (memhard.h), which carries the pack's item-to-word layout (era layout,
* 5 October 2026; the former w >> 4 / w & 15 here failed the random points of every interleaved pack). */
static uint32_t host_ds_word(uint32_t w) {
#ifdef IGNEUM_STATE_LEAVES
return mh_word(hCache, hLeaves, IGNEUM_STATE_LEAVES, w); /* class v5: every item keyed by leaf(t) = leaves[t mod n] */
#else
return mh_word(hCache, w);
#endif
}
#endif
@ -505,6 +514,13 @@ typedef struct {
int groupSize; // work-group size the program was built for (IGNEUM_GROUP = 32 x group-warps)
} Device;
/* class v5 build arguments and leaf buffers (defined after the serve pair below) */
static cl_int setBuildArgs(cl_kernel k, cl_mem ds, cl_mem cache, cl_mem leaves, cl_uint nLeaves, cl_uint nItems);
static int packLeavesBuffer(Device* dv, const char* packDir, cl_mem* leaves, cl_uint* nLeaves, char* err, size_t errCap);
#ifdef IGNEUM_STATE_LEAVES
static int compiledInLeavesBuffer(Device* dv, cl_mem* leaves, cl_uint* nLeaves, char* err, size_t errCap);
#endif
static const char* exchangeName(int m) { return m == 1 ? "sub_group_shuffle_xor (cl_khr_subgroup_shuffle)" : m == 2 ? "intel_sub_group_shuffle_xor (cl_intel_subgroups)" : "local-memory exchange with barrier"; }
// Returns 0 on success, 1 on build failure (log printed).
@ -819,10 +835,19 @@ static SizeResult runSize(Device* dv, const DeviceInfo* di, const Options* o, in
#if IGNEUM_DATASET_MODE == 1
cl_uint nItems = r.words / 16u;
size_t local = kernelMaxLocal(dv, dv->kBuild, di, 256);
CL_CHECK(clSetKernelArg(dv->kBuild, 0, sizeof(cl_mem), &dDs));
CL_CHECK(clSetKernelArg(dv->kBuild, 1, sizeof(cl_mem), &gCache));
CL_CHECK(clSetKernelArg(dv->kBuild, 2, sizeof(cl_uint), &nItems));
cl_mem leaves = NULL;
cl_uint nLeaves = 0;
#ifdef IGNEUM_STATE_LEAVES
{ /* class v5 compiled-in pack: its leaves from the pack directory beside the kernel */
char lerr[800];
if (!compiledInLeavesBuffer(dv, &leaves, &nLeaves, lerr, sizeof(lerr))) { printf("FAIL: %s\n", lerr); exit(2); }
if (pass == 0) printf("class v5: %u state leaves uploaded for the build (FNV-1a 64 %016llx checked against the pack)\n", (unsigned)nLeaves, (unsigned long long)IGNEUM_STATE_LEAVES_FNV64);
}
#endif
CL_CHECK(setBuildArgs(dv->kBuild, dDs, gCache, leaves, nLeaves, nItems));
ev = launch1D(dv, dv->kBuild, nItems, local);
CL_CHECK(clFinish(dv->q));
if (leaves) { clReleaseMemObject(leaves); ++gMemReleased; }
#else
cl_uint n = r.words, d0 = IGNEUM_DAY0, d1 = IGNEUM_DAY1;
size_t local = kernelMaxLocal(dv, dv->kFill, di, 256);
@ -1088,6 +1113,7 @@ typedef struct {
int checked;
char programClass[8]; /* Counter ASIC 2.0: the pack's class ("v2" or "v3") and era seed, from packfile.h */
char eraHex[65];
cl_uint stateLeaves; /* class v5: the state leaves the dataset was built from (uploaded for the build, released after); 0 otherwise */
} ServePair;
static int hexEq(const char* a, const char* b) {
@ -1132,6 +1158,65 @@ static void releasePair(ServePair* p) {
free(p);
}
/* Class v5 (docs/design/class-v5-stored-state.md, 7 October 2026): igneum_build takes the window's state leaves. The
* emitted kernel of a class v5 pack is igneum_build(ds, cache, leaves, nLeaves, nItems); every other class's is
* igneum_build(ds, cache, nItems). One place sets the arguments for both shapes. */
static cl_int setBuildArgs(cl_kernel k, cl_mem ds, cl_mem cache, cl_mem leaves, cl_uint nLeaves, cl_uint nItems) {
cl_int e = clSetKernelArg(k, 0, sizeof(cl_mem), &ds);
if (e == CL_SUCCESS) e = clSetKernelArg(k, 1, sizeof(cl_mem), &cache);
if (leaves) {
if (e == CL_SUCCESS) e = clSetKernelArg(k, 2, sizeof(cl_mem), &leaves);
if (e == CL_SUCCESS) e = clSetKernelArg(k, 3, sizeof(cl_uint), &nLeaves);
if (e == CL_SUCCESS) e = clSetKernelArg(k, 4, sizeof(cl_uint), &nItems);
} else if (e == CL_SUCCESS) e = clSetKernelArg(k, 2, sizeof(cl_uint), &nItems);
return e;
}
/* The leaves of a pack directory as a device buffer (NULL with *nLeaves 0 for a pack of another class): read and checked
* against the pack's count and FNV-1a 64 by pf_load_leaves before the upload. The caller releases the buffer after the
* build (the leaves are not needed while hashing). Returns 1, or 0 with err set. */
static int packLeavesBuffer(Device* dv, const char* packDir, cl_mem* leaves, cl_uint* nLeaves, char* err, size_t errCap) {
PfPack pk;
char perr[700];
uint32_t* h = NULL;
size_t bytes = 0;
cl_int e = 0;
*leaves = NULL; *nLeaves = 0;
if (!packDir || !packDir[0]) return 1;
if (!pf_load(packDir, &pk, perr, sizeof(perr))) { snprintf(err, errCap, "pack %s: %s", packDir, perr); return 0; }
if (!pf_load_leaves(packDir, &pk, &h, &bytes, perr, sizeof(perr))) { snprintf(err, errCap, "pack %s: %s", packDir, perr); return 0; }
if (!h) return 1;
*leaves = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, bytes, h, &e);
free(h);
if (e != CL_SUCCESS) { *leaves = NULL; snprintf(err, errCap, "clCreateBuffer state leaves (%s)", clErrName(e)); return 0; }
++gMemCreated;
*nLeaves = pk.stateLeaves;
return 1;
}
#ifdef IGNEUM_STATE_LEAVES
/* The compiled-in pack is a class v5 pack: its leaves live beside IGNEUM_KERNEL_PATH (the directory build.sh compiled
* against). The plain bench and the placeholder serve path build from them; without the file they build nothing. */
static int compiledInLeavesBuffer(Device* dv, cl_mem* leaves, cl_uint* nLeaves, char* err, size_t errCap) {
static char dir[1200];
const char* k = IGNEUM_KERNEL_PATH;
const char* slash = strrchr(k, '/');
#ifdef _WIN32
const char* bslash = strrchr(k, '\\');
if (bslash && (!slash || bslash > slash)) slash = bslash;
#endif
if (slash) { size_t n = (size_t)(slash - k); if (n >= sizeof(dir)) n = sizeof(dir) - 1; memcpy(dir, k, n); dir[n] = 0; } else strcpy(dir, ".");
if (!hLeaves) {
/* the host copy for host_ds_word (the 64 random points of the dataset self-test), checked by pf_load_leaves */
PfPack pk; char perr[700]; size_t bytes = 0;
if (!pf_load(dir, &pk, perr, sizeof(perr))) { snprintf(err, errCap, "pack %s: %s", dir, perr); return 0; }
if (!pf_load_leaves(dir, &pk, &hLeaves, &bytes, perr, sizeof(perr))) { snprintf(err, errCap, "pack %s: %s", dir, perr); return 0; }
if (!hLeaves || pk.stateLeaves != IGNEUM_STATE_LEAVES) { snprintf(err, errCap, "pack %s: program.h on disk says %u state leaves, this exe was built for %u", dir, (unsigned)pk.stateLeaves, (unsigned)IGNEUM_STATE_LEAVES); return 0; }
}
return packLeavesBuffer(dv, dir, leaves, nLeaves, err, errCap);
}
#endif
/* The prepare request and its result, handed between the main loop and the prepare thread. */
typedef struct {
Device* dv;
@ -1254,12 +1339,19 @@ static int pairBuffers(Device* dv, const DeviceInfo* di, cl_command_queue q, Ser
++gMemCreated;
local = kernelMaxLocal(dv, p->kBuild, di, 256);
g = ((nItems + local - 1) / local) * local;
e = clSetKernelArg(p->kBuild, 0, sizeof(cl_mem), &p->ds);
if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 1, sizeof(cl_mem), &p->cache);
if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 2, sizeof(cl_uint), &nItems);
if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kBuild, 1, NULL, &g, &local, 0, NULL, NULL);
if (e == CL_SUCCESS) e = clFinish(q);
if (e != CL_SUCCESS) { snprintf(err, errCap, "dataset build (%s)", clErrName(e)); return 0; }
{
/* class v5: the pack's state leaves, uploaded for the build (checked by pf_load_leaves) and released after it;
* a v5 pack without them returns 0 here and the pair is never served (the known-failed case) */
cl_mem leaves = NULL;
cl_uint nLeaves = 0;
if (!packLeavesBuffer(dv, packDir, &leaves, &nLeaves, err, errCap)) return 0;
e = setBuildArgs(p->kBuild, p->ds, p->cache, leaves, nLeaves, nItems);
if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kBuild, 1, NULL, &g, &local, 0, NULL, NULL);
if (e == CL_SUCCESS) e = clFinish(q);
if (leaves) { clReleaseMemObject(leaves); ++gMemReleased; }
if (e != CL_SUCCESS) { snprintf(err, errCap, "dataset build (%s%s)", clErrName(e), nLeaves ? ", class v5 with the state leaves" : ""); return 0; }
p->stateLeaves = nLeaves;
}
p->datasetMs = wallMs() - tb;
if (p->kHotFill && p->hotWords) {
/* hot-table experiment: the epoch's table from the pack's own fill kernel (one work-item per segment) */
@ -1395,7 +1487,8 @@ static int runBenchPack(Device* dv, const DeviceInfo* di, const Options* o) {
memcpy(cur->sw, gPack.seedw, 32); memcpy(cur->kw, gPack.keyw, 32);
if (!pairHotKernel(cur, o->packDir, perr, sizeof(perr))) { printf("FAIL: pack %s: %s\n", o->packDir, perr); return 1; }
if (!pairBuffers(dv, di, dv->q, cur, words, gServeCacheWords, gServeSegments, o->packDir, perr, sizeof(perr))) { printf("FAIL: pack %s: %s\n", o->packDir, perr); return 1; }
printf("pack %s: cache %.0f dataset %.0f hot %.0f check %.0f ms (%.0f ms in all); %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->hotMs, cur->checkMs, wallMs() - t0, cur->check);
printf("pack %s: cache %.0f dataset %.0f hot %.0f check %.0f ms (%.0f ms in all); class %s%s; %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->hotMs, cur->checkMs, wallMs() - t0,
gPack.programClass, cur->stateLeaves ? " (class v5: the state leaves uploaded for the build and released)" : "", cur->check);
printKernelInfo(di, cur->kHashBound, "igneum_hash_bound", (int)groupSize, "");
dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)nonces * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out");
dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, cur->sw, &err); CL_CHECK_ERR(err, "clCreateBuffer init words");
@ -1478,17 +1571,26 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
strncpy(cur->programClass, gPack.programClass, sizeof(cur->programClass) - 1); strncpy(cur->eraHex, gPack.eraHex, sizeof(cur->eraHex) - 1);
if (!pairHotKernel(cur, o->packDir, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o->packDir, perr); fflush(stdout); return 1; }
if (!pairBuffers(dv, di, dv->q, cur, words, gServeCacheWords, gServeSegments, o->packDir, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o->packDir, perr); fflush(stdout); return 1; }
printf("info first pack %s: cache %.0f dataset %.0f check %.0f ms (%.0f ms in all); %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->checkMs, wallMs() - t0, cur->check);
printf("info first pack %s: cache %.0f dataset %.0f check %.0f ms (%.0f ms in all); class %s%s; %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->checkMs, wallMs() - t0,
cur->programClass[0] ? cur->programClass : "v2", cur->stateLeaves ? " (class v5: the state leaves uploaded for the build and released)" : "", cur->check);
} else {
if (!setupCache(dv, di)) { printf("error 0 cache check failed (device cache differs from the host cache or the pack's FNV)\n"); fflush(stdout); return 1; }
memcpy(cur->sw, SEEDW, 32); memcpy(cur->kw, KEYW, 32);
cur->cache = gCache; gCache = NULL;
cur->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer dataset");
CL_CHECK(clSetKernelArg(cur->kBuild, 0, sizeof(cl_mem), &cur->ds));
CL_CHECK(clSetKernelArg(cur->kBuild, 1, sizeof(cl_mem), &cur->cache));
CL_CHECK(clSetKernelArg(cur->kBuild, 2, sizeof(cl_uint), &nItems));
countRelease(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256)));
CL_CHECK(clFinish(dv->q));
{
cl_mem leaves = NULL;
cl_uint nLeaves = 0;
#ifdef IGNEUM_STATE_LEAVES
char lerr[800];
if (!compiledInLeavesBuffer(dv, &leaves, &nLeaves, lerr, sizeof(lerr))) { printf("error 0 %s\n", lerr); fflush(stdout); return 1; }
#endif
CL_CHECK(setBuildArgs(cur->kBuild, cur->ds, cur->cache, leaves, nLeaves, nItems));
countRelease(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256)));
CL_CHECK(clFinish(dv->q));
if (leaves) { clReleaseMemObject(leaves); ++gMemReleased; }
cur->stateLeaves = nLeaves;
}
}
dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)batch * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out");
dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY, 32, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer init words");

51
tools/class-v5/fleet-cuda-v5-bench.sh Executable file
View file

@ -0,0 +1,51 @@
#!/usr/bin/env bash
# Class v5 kit bench on one fleet NVIDIA card (7 October 2026, docs/design/class-v5-stored-state.md section 13, "the kit:
# CUDA"). Run ON the pod by the fleet lane, in a directory holding the unzipped kit (packs-ca3-v5-<stamp>.zip from
# igneum-build-1:/srv/artefacts/packs/): the kit's Linux NVRTC worker (bin/linux/igneum-worker-cuda, THIS tree's
# proto-cuda/nvrtc/worker.cpp with the class v5 leaf upload) runs --check and --bench on the first class v5 pack,
# packs/v5-dn3-epoch0 (program id e5a4ac5978462156, 11 leaves under state root 7e37a9fb..., cache FNV 7334fa46e5d972eb);
# then src/bench.cu (the nvcc harness of the 4090 rows) on the same pack. Expected on every line: the self-test PASS and
# the fingerprint of the 2^24 outputs at base nonce 0 equal to the Metal reading 82b19cbde8557ea5 (M5 Max, 18:07:09Z).
# One card, about a minute. Prints RESULT lines; exit 0 only when both fingerprints match.
# bash tools/fleet-cuda-v5-bench.sh [device index, default 0] [sm arch for nvcc, default from nvidia-smi]
set -u
DEV="${1:-0}"
EXPECTED=82b19cbde8557ea5
KIT="$(cd "$(dirname "$0")/.." && pwd)"
export PATH=/usr/local/cuda/bin:$PATH
echo "RESULT start $(date -u +%Y-%m-%dT%H:%M:%SZ) host=$(hostname) kit=$KIT expected_fingerprint=$EXPECTED"
nvidia-smi -i "$DEV" --query-gpu=name,driver_version,power.limit,memory.total --format=csv,noheader | sed 's/^/RESULT card /'
ARCH="${2:-}"
if [ -z "$ARCH" ]; then cc=$(nvidia-smi -i "$DEV" --query-gpu=compute_cap --format=csv,noheader | tr -d ' .'); ARCH="sm_$cc"; fi
PACK="$KIT/packs/v5-dn3-epoch0"
[ -f "$PACK/leaves.bin" ] || { echo "RESULT error leaves.bin missing in $PACK (a class v5 pack without its leaves builds nothing)"; exit 2; }
echo "RESULT leaves $(wc -c < "$PACK/leaves.bin") bytes sha256 $(sha256sum "$PACK/leaves.bin" | cut -c1-64)"
W="$KIT/bin/linux/igneum-worker-cuda"
chmod +x "$W" 2>/dev/null
echo "RESULT worker $W sha256 $(sha256sum "$W" | cut -c1-64)"
ok=0
echo "RESULT check start $(date -u +%H:%M:%SZ) cmd=igneum-worker-cuda --check --pack packs/v5-dn3-epoch0 --device $DEV"
"$W" --check --pack "$PACK" --device "$DEV" 2>&1 | sed 's/^/RESULT check out /'
echo "RESULT bench start $(date -u +%H:%M:%SZ) cmd=igneum-worker-cuda --bench --pack packs/v5-dn3-epoch0 --device $DEV --batch-log2 24 --batches 5"
out=$("$W" --bench --pack "$PACK" --device "$DEV" --batch-log2 24 --batches 5 2>&1); rc=$?
echo "$out" | grep -E '^(pack |RESULT|FAIL|warm-up)' | sed 's/^/RESULT bench out /'
fp=$(echo "$out" | grep -o 'fingerprint=[0-9a-f]*' | tail -1 | cut -d= -f2)
mhs=$(echo "$out" | grep -o 'mhs=[0-9.]*' | tail -1 | cut -d= -f2)
chk=$(echo "$out" | grep -o 'check=[A-Za-z]*' | tail -1 | cut -d= -f2)
match=no; [ "$fp" = "$EXPECTED" ] && [ "$chk" = PASS ] && { match=yes; ok=$((ok + 1)); }
echo "RESULT worker-v5 fingerprint=$fp expected=$EXPECTED match=$match check=$chk mhs=$mhs exit=$rc $(date -u +%H:%M:%SZ)"
# the nvcc harness (the 4090 rows of the design page, section 7): compiled in the pack directory against its kernel.cu
if command -v nvcc > /dev/null; then
( cd "$PACK" && nvcc -O3 -std=c++17 -arch="$ARCH" -Xcompiler -pthread -I. -o bench "$KIT/src/bench.cu" kernel.cu 2>&1 | grep -v '^$' | sed 's/^/RESULT nvcc /'; echo "RESULT nvcc rc ${PIPESTATUS[0]} arch=$ARCH" )
if [ -x "$PACK/bench" ]; then
out2=$(cd "$PACK" && CUDA_VISIBLE_DEVICES="$DEV" ./bench --batches 10 --power-seconds 0 2>&1); rc2=$?
echo "$out2" | sed 's/^/RESULT bench.cu out /'
fp2=$(echo "$out2" | grep -o 'fingerprint of 2^24 outputs at base 0: [0-9a-f]*' | grep -o '[0-9a-f]\{16\}$')
m2=no; [ "$fp2" = "$EXPECTED" ] && echo "$out2" | grep -q 'vector warps against the pack: 3 of 3' && { m2=yes; ok=$((ok + 1)); }
echo "RESULT bench.cu-v5 fingerprint=$fp2 expected=$EXPECTED match=$m2 exit=$rc2 $(date -u +%H:%M:%SZ)"
fi
else
echo "RESULT nvcc absent: the bench.cu row is skipped (the worker row stands)"; ok=$((ok + 1))
fi
echo "RESULT end $(date -u +%Y-%m-%dT%H:%M:%SZ) rows_matched=$ok of 2"
[ "$ok" = 2 ]

112
tools/class-v5/kits-remote.sh Executable file
View file

@ -0,0 +1,112 @@
#!/usr/bin/env bash
# The class v5 kits for every platform, built and tested on igneum-build-1 (docs/design/class-v5-stored-state.md section 13,
# the kit rows; 7 October 2026). From THIS worktree's commit (HEAD through the bare mirror, as tools/build-remote.sh moves
# sources; nothing is built on the Mac): the pack loader test with the class v5 cases (proto-cuda/nvrtc/emu/packfile-test.c),
# the CPU emulation of the NVRTC worker running --check on the class v5 pack (the real kernel text on host threads, the
# dataset built from the pack's leaves and checked against the pack's cache FNV, dataset head, last word, 64 samples and
# three vector warps: the "exactly as igneum-pow's verifier" line without a GPU), the Linux and Windows builds of the two
# one-click workers (proto-cuda/nvrtc/worker.cpp through NVRTC, proto-opencl/host.c for AMD, Intel and Apple), then one kit
# zip at /srv/artefacts/packs/packs-ca3-v5-<stamp>.zip holding the three packs (v4-genesis, v5-genesis, v5-dn3-epoch0), the
# binaries, the sources, the bench scripts for the fleet and PC 1, and SHA256SUMS. Prints the zip's sha256 last.
#
# IGNEUM_AGENT=v5-kits tools/class-v5/kits-remote.sh [--box 1] [--skip-emu]
#
# Takes one of the box's build slots through infra/build-server/remote-run.sh (bs_remote_run), box 1 by default (the
# build-server lane's word of 7 October 2026, 19:4x BST: the kit builds on build-1 explicitly).
set -euo pipefail
HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
BS_TOOL=class-v5-kits
# shellcheck source=../../infra/build-server/lib.sh
. "$HERE/../../infra/build-server/lib.sh"
BOX=1; SKIP_EMU=0
while [ $# -gt 0 ]; do
case "$1" in --box) BOX="$2"; shift 2 ;; --skip-emu) SKIP_EMU=1; shift ;; *) echo "unknown argument $1" >&2; exit 2 ;; esac
done
export IGNEUM_AGENT="${IGNEUM_AGENT:-v5-kits}"
bs_host "$BOX"
pushd "$HERE/../../igneum-pow" > /dev/null; bs_context; popd > /dev/null
WT="$BS_REMOTE_WT"
STAMP=$(date -u +%Y%m%dT%H%M%SZ)
ZIP="/srv/artefacts/packs/packs-ca3-v5-$STAMP.zip"
bs_toolchain_check
bs_sync_sources
bs_log "sources at $WT (commit $BS_SHA on $BS_BRANCH); building the class v5 kits on box $BOX"
read -r -d '' CMD <<EOF2 || true
set -u
cd '$WT'
T=\$(mktemp -d /tmp/v5-kits.XXXXXX); S=\$T/stage; mkdir -p \$S/bin/linux \$S/bin/windows \$S/src \$S/packs \$S/tools
rc=0
step() { echo "STEP \$1 \$(date -u +%H:%M:%SZ)"; }
fail() { echo "FAIL \$1"; rc=1; }
V5=proto-cuda/packs-ca3-v5/v5-dn3-epoch0
step packfile-test
cc -std=c99 -Wall -Wextra -Wno-unused-function -O1 -o \$T/packfile-test proto-cuda/nvrtc/emu/packfile-test.c && \$T/packfile-test proto-cuda/packs/igneum-devnet-v4-epoch0 \$V5 > \$T/packfile-test.log 2>&1 || fail "packfile-test (\$(tail -1 \$T/packfile-test.log))"
grep -E '^(FAIL| v5 pack)' \$T/packfile-test.log; grep -c '^ok' \$T/packfile-test.log | sed 's/^/packfile-test ok lines /'
step linux-opencl-worker
gcc -std=c99 -O2 -Wall -Wextra -Wno-stringop-truncation -Wno-format-truncation -DIGNEUM_CL_DYNAMIC -DCL_TARGET_OPENCL_VERSION=120 -I proto-cuda/packs/igneum-devnet-v4-epoch0 -DIGNEUM_KERNEL_PATH='"kernel_bound.cl"' -o \$S/bin/linux/igneum-worker-opencl proto-opencl/host.c -ldl -lpthread 2> \$T/cl-linux.log || { fail "linux opencl worker"; head -20 \$T/cl-linux.log; }
step linux-cuda-worker
g++ -std=c++17 -O2 -Wall -Wextra -I proto-cuda/nvrtc -I /usr/local/cuda/include -o \$S/bin/linux/igneum-worker-cuda proto-cuda/nvrtc/worker.cpp -ldl -lpthread 2> \$T/cuda-linux.log || { fail "linux cuda worker"; head -20 \$T/cuda-linux.log; }
step windows-cuda-worker
x86_64-w64-mingw32-g++ -std=c++17 -O2 -Wall -Wextra -static -I proto-cuda/nvrtc -I /usr/local/cuda/include -o \$S/bin/windows/igneum-worker-cuda.exe proto-cuda/nvrtc/worker.cpp 2> \$T/cuda-win.log && x86_64-w64-mingw32-strip \$S/bin/windows/igneum-worker-cuda.exe || { fail "windows cuda worker"; head -20 \$T/cuda-win.log; }
step windows-opencl-worker
x86_64-w64-mingw32-gcc -std=c99 -O2 -Wall -Wextra -Wno-stringop-truncation -Wno-format-truncation -static -DIGNEUM_CL_DYNAMIC -DCL_TARGET_OPENCL_VERSION=120 -I /usr/include -I proto-cuda/packs/igneum-devnet-v4-epoch0 -DIGNEUM_KERNEL_PATH='"kernel_bound.cl"' -o \$S/bin/windows/igneum-worker-opencl.exe proto-opencl/host.c 2> \$T/cl-win.log && x86_64-w64-mingw32-strip \$S/bin/windows/igneum-worker-opencl.exe || { fail "windows opencl worker"; head -20 \$T/cl-win.log; }
if [ $SKIP_EMU = 0 ]; then
step emu-check-v5
# the CPU emulation (proto-cuda/nvrtc/emu/test.sh's build) with pack A = the class v5 pack and pack B = the v4 control
E=\$T/emu; mkdir -p \$E
emu_kernel() { { echo '#include <cuda_runtime.h>'; echo '#include <cstdint>'; echo "namespace \$2 {"; sed -E 's/([A-Za-z_0-9]+)<<<([^,]+), ([^>]+)>>>\\(/emu_launch(\\1, \\2, \\3, /' "\$1/\$3.cu"; echo "}"; } > "\$E/\$3_\$2.cpp"; g++ -std=c++17 -O2 -w -I proto-cuda/emu -I "\$1" -c "\$E/\$3_\$2.cpp" -o "\$E/\$3_\$2.o"; }
emu_kernel \$V5 emu_pack_a kernel && emu_kernel \$V5 emu_pack_a kernel_bound && emu_kernel proto-cuda/packs-ca3-v5/v4-genesis emu_pack_b kernel && emu_kernel proto-cuda/packs-ca3-v5/v4-genesis emu_pack_b kernel_bound \\
&& g++ -std=c++17 -O2 -Wall -Wextra -DIGNEUM_EMU -I proto-cuda/nvrtc -I /usr/local/cuda/include -c proto-cuda/nvrtc/worker.cpp -o \$E/worker.o \\
&& g++ -std=c++17 -O2 -Wall -Wextra -DIGNEUM_EMU -DIGNEUM_EMU_TWO_PACKS -I proto-cuda/nvrtc -I proto-cuda/emu -I /usr/local/cuda/include -c proto-cuda/nvrtc/emu/emu_backend.cpp -o \$E/emu_backend.o \\
&& g++ -std=c++17 -O2 -w -I proto-cuda/emu -c proto-cuda/emu/shim.cpp -o \$E/shim.o \\
&& g++ -o \$E/igneum-worker-cuda-emu \$E/*.o -pthread 2> \$T/emu-build.log || { fail "emu build"; head -30 \$T/emu-build.log; }
if [ -x \$E/igneum-worker-cuda-emu ]; then
( IGNEUM_EMU_PACK=\$PWD/\$V5 IGNEUM_EMU_PACK2=\$PWD/proto-cuda/packs-ca3-v5/v4-genesis timeout 1500 nice -n 10 \$E/igneum-worker-cuda-emu --check --pack \$V5 > \$T/emu-check.log 2>&1 ) || fail "emu --check on the class v5 pack (rc \$?)"
grep -E '^check PASS|self-test|FAIL|error' \$T/emu-check.log | cut -c1-400
fi
fi
step stage
for p in v4-genesis v5-genesis v5-dn3-epoch0; do mkdir -p \$S/packs/\$p; cp proto-cuda/packs-ca3-v5/\$p/* \$S/packs/\$p/; done
cp proto-cuda/nvrtc/worker.cpp proto-cuda/nvrtc/packfile.h proto-cuda/nvrtc/cuda_api.h proto-opencl/host.c proto-opencl/cl_dynamic.h proto-newpow/class-v5/bench.cu proto-newpow/class-v5/run.sh proto-metal/packbench.swift \$S/src/
cp tools/class-v5/fleet-cuda-v5-bench.sh tools/class-v5/pc1-amd-v5-bench.ps1 \$S/tools/
cat > \$S/README.txt <<'R'
packs-ca3-v5: the class v5 kit (Igneum, 7 October 2026; docs/design/class-v5-stored-state.md)
packs/v5-dn3-epoch0 the first class v5 pack: Devnet 3 epoch 0, program id e5a4ac5978462156, 11 state leaves (leaves.bin, 704 B) under
state root 7e37a9fb19b154d32daf5bf30a50d339a75029fbc9eec9ea20e95439dba5a311, cache FNV-1a 64 7334fa46e5d972eb
expected fingerprint of the 2^24 outputs at base nonce 0: 82b19cbde8557ea5 (Metal, M5 Max, 18:07:09Z) on every platform
packs/v5-genesis the string-seed class v5 pack over the devnet's 93 leaves; packs/v4-genesis the class v4 control (sub-version 3)
bin/linux igneum-worker-cuda (NVRTC, libcuda + libnvrtc.so.12 at run time), igneum-worker-opencl (libOpenCL.so.1 at run time)
bin/windows igneum-worker-cuda.exe (nvcuda.dll + nvrtc64 at run time), igneum-worker-opencl.exe (OpenCL.dll at run time)
Every worker reads a class v5 pack's leaves.bin, checks it against the pack's IGNEUM_STATE_LEAVES_FNV64, uploads it for igneum_build and
frees it after the build; a pack without its leaves builds nothing. Self-test: cache head, last line and FNV; dataset head, last word and
64 samples; 96 vector lanes. Bench: igneum-worker-cuda --bench --pack packs/v5-dn3-epoch0 --batch-log2 24 (fleet-cuda-v5-bench.sh);
igneum-worker-opencl --bench-pack --pack packs/v5-dn3-epoch0 --batch-log2 24 --device <n> (pc1-amd-v5-bench.ps1); src/bench.cu with nvcc.
R
( cd \$S && find . -type f ! -name SHA256SUMS | sort | xargs sha256sum > SHA256SUMS )
mkdir -p /srv/artefacts/packs
python3 - "\$S" '$ZIP' <<'PY'
import os, sys, zipfile
stage, out = sys.argv[1], sys.argv[2]
with zipfile.ZipFile(out, 'w', zipfile.ZIP_DEFLATED) as z:
for dp, dn, fn in os.walk(stage):
dn.sort()
for f in sorted(fn):
p = os.path.join(dp, f); rel = os.path.relpath(p, stage)
zi = zipfile.ZipInfo(rel, date_time=(2026, 10, 7, 0, 0, 0)); zi.compress_type = zipfile.ZIP_DEFLATED
zi.external_attr = (0o755 if rel.startswith('bin/') or rel.endswith('.sh') else 0o644) << 16
with open(p, 'rb') as fh: z.writestr(zi, fh.read())
PY
sha256sum '$ZIP' | cut -c1-64 > '$ZIP.sha256'
echo "KIT $ZIP bytes \$(stat -c %s '$ZIP') files \$(python3 -c "import zipfile,sys; print(len(zipfile.ZipFile(sys.argv[1]).namelist()))" '$ZIP')"
echo "KIT sha256 \$(cat '$ZIP.sha256')"
for b in \$S/bin/linux/* \$S/bin/windows/*; do echo "BIN \$(basename \$(dirname \$b))/\$(basename \$b) \$(stat -c %s \$b) bytes sha256 \$(sha256sum \$b | cut -c1-64)"; done
cp \$T/*.log /srv/builds/_log/v5-class/kits-$STAMP/ 2>/dev/null || { mkdir -p /srv/builds/_log/v5-class/kits-$STAMP && cp \$T/*.log /srv/builds/_log/v5-class/kits-$STAMP/; }
echo "LOGS /srv/builds/_log/v5-class/kits-$STAMP"
rm -rf \$T
( exit \$rc )
EOF2
BR_KIND=build BR_COMMAND="class v5 kits (packfile-test, emu --check v5, Linux and Windows workers, kit zip)" BR_TARGET="x86_64-linux+windows" BR_ARTEFACTS="" \
bs_remote_run "$WT" "$BS_WT class-v5 kits" "$CMD" 2>&1 | tee "/tmp/v5-kits-$STAMP.log" | grep -E '^(STEP|FAIL|KIT|BIN|LOGS|check PASS|self-test|packfile-test|emu:| v5|build-remote: RESULT)' || true
rc=${PIPESTATUS[0]}
bs_wt_unlock
[ "$rc" = 0 ] && bs_log "class v5 kits built: $ZIP" || bs_die "class v5 kits FAILED (rc $rc); full log /tmp/v5-kits-$STAMP.log"

80
tools/class-v5/pc1-amd-v5-bench.ps1 Normal file
View file

@ -0,0 +1,80 @@
# Class v5 kit bench on PC 1's RX 9070 XT (7 October 2026, docs/design/class-v5-stored-state.md section 13, "the kit:
# OpenCL (AMD)"): the kit's igneum-worker-opencl.exe (THIS tree's proto-opencl/host.c with the class v5 leaf upload:
# igneum_build(ds, cache, leaves, nLeaves, nItems), the leaves of leaves.bin checked against the pack's FNV-1a 64 first)
# runs --bench-pack on the first class v5 pack, proto-cuda/packs-ca3-v5/v5-dn3-epoch0 (program id e5a4ac5978462156, 11
# leaves under state root 7e37a9fb..., cache FNV 7334fa46e5d972eb), on the gfx1201 device. Expected: the self-test
# PASS line (cache head, last line and FNV; dataset head, word [268435455] and 64 samples; 96 of 96 vector lanes) and the
# fingerprint of the 2^24 outputs at base nonce 0 equal to the Metal reading, 82b19cbde8557ea5 (the M5 Max, 7 October
# 2026, 18:07:09Z). The v4-genesis control pack runs after it so the run has a known class v4 row beside the v5 row.
# Beside the miners (a correctness and fingerprint gate; the rate row is labelled loaded when the card mines): no card is
# switched, nothing is posted to the installed app, nothing is built on the PC. Published by the Counter ASIC coordinator
# only (the PC 1 queue is its); the kit is fetch job $kitId. Every result line starts with RESULT; SUMMARY {json} ends it.
$ErrorActionPreference = 'Continue'
$jobName = 'pc1-amd-v5-bench'
$kitId = $env:IGNEUM_V5_KIT_ID; if (-not $kitId) { $kitId = 'fetch-ca3-v5-kit-20261007' }
$expected = '82b19cbde8557ea5'
$started = Get-Date
function Stamp { (Get-Date).ToUniversalTime().ToString('yyyy-MM-ddTHH:mm:ssZ') }
function Summary([string] $status, [hashtable] $extra) {
$o = [ordered]@{ job = $jobName; status = $status; duration_s = [int]((Get-Date) - $started).TotalSeconds; finished_at = (Stamp) }
foreach ($k in $extra.Keys) { $o[$k] = $extra[$k] }
'SUMMARY ' + ($o | ConvertTo-Json -Compress -Depth 4)
}
"RESULT start $(Stamp) job=$jobName machine=$env:COMPUTERNAME app_version=$env:IGNEUM_APP_VERSION expected_fingerprint=$expected"
$jobs = Split-Path $env:IGNEUM_JOB_DIR
$kit = Join-Path $jobs $kitId
if (-not (Test-Path $kit)) { "RESULT error kit missing at $kit (the fetch job $kitId runs first; republish it after any app update)"; Summary 'failed' @{ error = 'kit missing' }; exit 2 }
$exe = Join-Path $kit 'bin\windows\igneum-worker-opencl.exe'
$packs = Join-Path $kit 'packs'
if (-not (Test-Path $exe)) { "RESULT error worker missing at $exe"; Summary 'failed' @{ error = 'worker missing' }; exit 2 }
"RESULT worker kit $exe sha256 $((Get-FileHash -Algorithm SHA256 $exe).Hash.ToLower()) bytes $((Get-Item $exe).Length)"
foreach ($pk in @('v5-dn3-epoch0', 'v4-genesis')) {
$d = Join-Path $packs $pk
if (-not (Test-Path (Join-Path $d 'kernel_bound.cl'))) { "RESULT error pack $pk missing at $d"; Summary 'failed' @{ error = "pack $pk missing" }; exit 2 }
"RESULT pack $pk kernel_bound.cl sha256 $((Get-FileHash -Algorithm SHA256 (Join-Path $d 'kernel_bound.cl')).Hash.ToLower()) program.h sha256 $((Get-FileHash -Algorithm SHA256 (Join-Path $d 'program.h')).Hash.ToLower())"
}
$leaves = Join-Path $packs 'v5-dn3-epoch0\leaves.bin'
if (-not (Test-Path $leaves)) { "RESULT error leaves.bin missing at $leaves (a class v5 pack without its leaves builds nothing)"; Summary 'failed' @{ error = 'leaves missing' }; exit 2 }
"RESULT leaves $leaves bytes $((Get-Item $leaves).Length) sha256 $((Get-FileHash -Algorithm SHA256 $leaves).Hash.ToLower())"
# the OpenCL device index of the 9070 XT (the installed worker's list when present: the app's own indices)
$inst = @("$env:LOCALAPPDATA\Programs\Igneum Miner", "$env:ProgramFiles\Igneum Miner") | Where-Object { Test-Path (Join-Path $_ 'igneum-app.exe') } | Select-Object -First 1
$listExe = $exe
if ($inst -and (Test-Path (Join-Path $inst 'igneum-worker-opencl.exe'))) { $listExe = Join-Path $inst 'igneum-worker-opencl.exe' }
$list = @(& $listExe --list 2>&1 | ForEach-Object { "$_" })
$list | ForEach-Object { "RESULT list $_" }
$dev = $null
foreach ($l in $list) { if ($l -match '^\s*\[(\d+)\].*gfx1201' -and $l -notmatch 'dup') { $dev = [int]$Matches[1]; break } }
if ($null -eq $dev) { "RESULT error no gfx1201 device in --list (the eGPU is off the bus: the AMD v5 row stays OWED)"; Summary 'failed' @{ error = 'no gfx1201' }; exit 2 }
"RESULT device $dev gfx1201 (list from $listExe)"
function Workers { @(Get-CimInstance Win32_Process -Filter "Name = 'igneum-worker-opencl.exe' OR Name = 'igneum-worker-cuda.exe'" -ErrorAction SilentlyContinue | ForEach-Object { "$($_.Name):$($_.ProcessId):[$($_.CommandLine -replace '\s+', ' ')]" }) }
$w = @(Workers)
$loaded = ($w | Where-Object { $_ -match "igneum-worker-opencl.*--device\s+$dev(\s|$)" }).Count -gt 0
"RESULT workers_before $(Stamp) $($w -join ' ')"
$state = if ($loaded) { 'loaded' } else { 'quiet' }
"RESULT context card_state=$state (a fingerprint gate: the load changes the rate row, never the bytes)"
$rows = @{}
$fpOk = $false
foreach ($pk in @('v5-dn3-epoch0', 'v4-genesis')) {
$d = Join-Path $packs $pk
$t0 = Get-Date
"RESULT bench $pk start $(Stamp) cmd=igneum-worker-opencl.exe --bench-pack --pack $d --device $dev --batch-log2 24 --batches 5"
$out = @(& $exe --bench-pack --pack $d --device $dev --batch-log2 24 --batches 5 2>&1 | ForEach-Object { "$_" })
$code = $LASTEXITCODE
$secs = [int]((Get-Date) - $t0).TotalSeconds
foreach ($l in $out) { if ($l -match '^(pack |class v5|RESULT |FAIL|error|warm-up)') { "RESULT bench $pk out $l" } }
$res = $out | Where-Object { $_ -match '^RESULT ' } | Select-Object -Last 1
$fp = ''; $mhs = ''; $check = ''
if ($res -match 'fingerprint=([0-9a-f]{16})') { $fp = $Matches[1] }
if ($res -match 'mhs=([0-9.]+)') { $mhs = $Matches[1] }
if ($res -match 'check=(\w+)') { $check = $Matches[1] }
$rows[$pk] = [ordered]@{ exit = $code; seconds = $secs; fingerprint = $fp; mhs = $mhs; check = $check; card_state = $state }
if ($pk -eq 'v5-dn3-epoch0') {
$fpOk = ($fp -eq $expected -and $check -eq 'PASS')
"RESULT v5 fingerprint=$fp expected=$expected match=$fpOk check=$check mhs=$mhs card_state=$state exit=$code seconds=$secs $(Stamp)"
} else {
"RESULT v4-control fingerprint=$fp check=$check mhs=$mhs card_state=$state exit=$code seconds=$secs $(Stamp)"
}
}
"RESULT workers_after $(Stamp) $((@(Workers)) -join ' ')"
Summary $(if ($fpOk) { 'done' } else { 'failed' }) @{ expected = $expected; v5 = $rows['v5-dn3-epoch0']; v4_control = $rows['v4-genesis']; device = $dev; kit = $kitId }
exit $(if ($fpOk) { 0 } else { 1 })