From 8bf1c4f43df92101ae0e9d970d6ee8d7777c2dde Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Wed, 7 Oct 2026 18:36:17 +0000 Subject: [PATCH] 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 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 --- proto-cuda/nvrtc/cuda_api.h | 1 + proto-cuda/nvrtc/emu/emu_backend.cpp | 36 ++++++- proto-cuda/nvrtc/emu/packfile-test.c | 66 +++++++++++++ proto-cuda/nvrtc/emu/packfile-test.sh | 6 +- proto-cuda/nvrtc/packfile.h | 74 +++++++++++++- proto-cuda/nvrtc/worker.cpp | 52 +++++++--- proto-metal/main.swift | 58 ++++++++++- proto-opencl/host.c | 134 +++++++++++++++++++++++--- tools/class-v5/fleet-cuda-v5-bench.sh | 51 ++++++++++ tools/class-v5/kits-remote.sh | 112 +++++++++++++++++++++ tools/class-v5/pc1-amd-v5-bench.ps1 | 80 +++++++++++++++ 11 files changed, 624 insertions(+), 46 deletions(-) create mode 100755 tools/class-v5/fleet-cuda-v5-bench.sh create mode 100755 tools/class-v5/kits-remote.sh create mode 100644 tools/class-v5/pc1-amd-v5-bench.ps1 diff --git a/proto-cuda/nvrtc/cuda_api.h b/proto-cuda/nvrtc/cuda_api.h index e16e8c6e0..11d2398d2 100644 --- a/proto-cuda/nvrtc/cuda_api.h +++ b/proto-cuda/nvrtc/cuda_api.h @@ -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; diff --git a/proto-cuda/nvrtc/emu/emu_backend.cpp b/proto-cuda/nvrtc/emu/emu_backend.cpp index bb3583f77..feb9e6c02 100644 --- a/proto-cuda/nvrtc/emu/emu_backend.cpp +++ b/proto-cuda/nvrtc/emu/emu_backend.cpp @@ -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 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(params, 0), arg(params, 1)); return CUDA_SUCCESS; - case 2: emu_launch(emu_pack_a::igneum_build, gx, bx, dptr(params, 0), dptr(params, 1), arg(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(params, 0), dptr(params, 1), dptr(params, 2), arg(params, 3), arg(params, 4)); return CUDA_SUCCESS; } + if (!five && b3) { emu_launch(b3, gx, bx, dptr(params, 0), dptr(params, 1), arg(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(params, 0), dptr(params, 1), arg(params, 2), arg(params, 3), arg(params, 4)); return CUDA_SUCCESS; #ifdef IGNEUM_EMU_TWO_PACKS case 4: emu_launch(emu_pack_b::igneum_cache_fill, gx, bx, dptr(params, 0), arg(params, 1)); return CUDA_SUCCESS; - case 5: emu_launch(emu_pack_b::igneum_build, gx, bx, dptr(params, 0), dptr(params, 1), arg(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(params, 0), dptr(params, 1), dptr(params, 2), arg(params, 3), arg(params, 4)); return CUDA_SUCCESS; } + if (!five && b3) { emu_launch(b3, gx, bx, dptr(params, 0), dptr(params, 1), arg(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(params, 0), dptr(params, 1), arg(params, 2), arg(params, 3), arg(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; } diff --git a/proto-cuda/nvrtc/emu/packfile-test.c b/proto-cuda/nvrtc/emu/packfile-test.c index dca59a8a8..0bcc3d3b9 100644 --- a/proto-cuda/nvrtc/emu/packfile-test.c +++ b/proto-cuda/nvrtc/emu/packfile-test.c @@ -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; } diff --git a/proto-cuda/nvrtc/emu/packfile-test.sh b/proto-cuda/nvrtc/emu/packfile-test.sh index 5ccd2b51d..d945d0edf 100755 --- a/proto-cuda/nvrtc/emu/packfile-test.sh +++ b/proto-cuda/nvrtc/emu/packfile-test.sh @@ -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" diff --git a/proto-cuda/nvrtc/packfile.h b/proto-cuda/nvrtc/packfile.h index 48c09fab5..932be43b7 100644 --- a/proto-cuda/nvrtc/packfile.h +++ b/proto-cuda/nvrtc/packfile.h @@ -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=` and `era=` tokens (sent +/* Counter ASIC 2.0 (5 October 2026): a job or prepare line may end with `class=` and `era=` 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`. */ diff --git a/proto-cuda/nvrtc/worker.cpp b/proto-cuda/nvrtc/worker.cpp index 0439e4270..f373a81a2 100644 --- a/proto-cuda/nvrtc/worker.cpp +++ b/proto-cuda/nvrtc/worker.cpp @@ -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()); } // --------------------------------------------------------------------------------------------- diff --git a/proto-metal/main.swift b/proto-metal/main.swift index eff385acd..6c7e48823 100644 --- a/proto-metal/main.swift +++ b/proto-metal/main.swift @@ -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" { diff --git a/proto-opencl/host.c b/proto-opencl/host.c index 4abef5654..4cfb2cb02 100644 --- a/proto-opencl/host.c +++ b/proto-opencl/host.c @@ -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"); diff --git a/tools/class-v5/fleet-cuda-v5-bench.sh b/tools/class-v5/fleet-cuda-v5-bench.sh new file mode 100755 index 000000000..665bd7363 --- /dev/null +++ b/tools/class-v5/fleet-cuda-v5-bench.sh @@ -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-.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 ] diff --git a/tools/class-v5/kits-remote.sh b/tools/class-v5/kits-remote.sh new file mode 100755 index 000000000..2223ea780 --- /dev/null +++ b/tools/class-v5/kits-remote.sh @@ -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-.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 < \$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 '; echo '#include '; 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 (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" diff --git a/tools/class-v5/pc1-amd-v5-bench.ps1 b/tools/class-v5/pc1-amd-v5-bench.ps1 new file mode 100644 index 000000000..474b13406 --- /dev/null +++ b/tools/class-v5/pc1-amd-v5-bench.ps1 @@ -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 })