diff --git a/igneum-pow/src/verify.rs b/igneum-pow/src/verify.rs index 1197d4096..9d1281aeb 100644 --- a/igneum-pow/src/verify.rs +++ b/igneum-pow/src/verify.rs @@ -249,6 +249,15 @@ impl DatasetSource { Self { log2_words: self.log2_words, mask: self.mask, key: self.key, key_bytes: self.key_bytes.clone(), dataset, hot: None } } + /// A copy of this source sharing its cache (and leaves), for a caller that needs an owned source from a shared one. + pub fn refreshed_or_clone(&self) -> Self { + let dataset = match &self.dataset { + Dataset::MemoryHard(m) => Dataset::MemoryHard(crate::memhard::MemhardCpu { params: m.params.clone(), cache: m.cache.clone(), leaves: m.leaves.clone() }), + Dataset::ClosedForm { d0, d1 } => Dataset::ClosedForm { d0: *d0, d1: *d1 }, + }; + Self { log2_words: self.log2_words, mask: self.mask, key: self.key, key_bytes: self.key_bytes.clone(), dataset, hot: None } + } + /// The window's state leaves, when the source carries them. pub fn leaves(&self) -> Option<&std::sync::Arc> { self.memhard().and_then(|m| m.leaves.as_ref()) diff --git a/infra/fast-time/class-v5-signal.mjs b/infra/fast-time/class-v5-signal.mjs new file mode 100644 index 000000000..0e7db5f6d --- /dev/null +++ b/infra/fast-time/class-v5-signal.mjs @@ -0,0 +1,362 @@ +#!/usr/bin/env node +// Class v5, proof of stored state and of following (docs/design/class-v5-stored-state.md section 10): the class v4 +// signal gate's shape with class v5. Three nodes on override-60x.json with CPU genesis bits, v3 from DAA 60 (epoch 1), +// the v4 FLOOR at --v4-floor (default 120, epoch 2), the class v5 floor at --floor (default 100000: far away, so v5 is +// ENABLED and the signal decides; `never` turns the object off), the signal window at --window (default 60 DAA, one +// epoch; seven windows = 420 DAA, so the first epoch whose seed block has seven full windows below it is epoch 8 at +// DAA 480), each node's object byte by IGNEUM_CLASS_SIGNAL (--signal a,b,c), each node's exec RPC on its own loopback +// port and each CPU miner's --exec-rpc pointing at its node (the state stream after each epoch's seed block). A fourth +// node runs with --evm-disable (--stateless-node, default on): it mines until the flip and then serves no class v5 +// template, which the harness counts. Ports 29760 and up, network igneum-devnet-976, data /tmp/igneum-fast-time-v5s. +// +// The cases and the known-failed cases: +// --signal 5,5,4 --expect no-flip two of three signal v5: 67 percent, the class must stay v4 (run 11 epochs) +// --signal 5,5,5 --expect flip all three: v5 from epoch 8, the first boundary with seven full windows; every +// miner's v5 id equals the CLI's --program-class v5 id and differs from the same +// seed's v4 id; the stateless node's miner accepts 0 blocks after the flip; 0 rejected +// --signal 5,5,5 --floor never --expect no-flip the object off: byte 5 counts as v4 only +// --signal 4,4,4 --floor 600 --expect floor nobody signals v5: it flips at the floor (epoch 10) and not before +// --signal 5,5,5 --expect flip --stale 2 proof of following's known-failed case: miner 2 keeps the first stream it +// fetched (--freeze-state); from the epoch after the flip every block it mines is +// rejected (accepted after the first refresh 0, rejected above 0); PASS means the +// chain threw the stale miner off +// --signal 5,5,4 --expect flip the known-failed case of the harness itself: it must report FAIL (no flip) +// +// node infra/fast-time/class-v5-signal.mjs --signal a,b,c --expect flip|no-flip|floor [--floor |never] +// [--v4-floor 120] [--window 60] [--v3-activation 60] [--secs 780] [--epochs 11] [--stale ] [--no-stateless-node] +// IGNEUMD, IGNEUM_MINER, IGNEUM_POW name the binaries (defaults: the class-v5 fork worktree's target/release and +// igneum-pow/target/release/igneum-pow, the layout on igneum-build-1 under /srv/builds/igneum-wt-class-v5). + +import { spawn, spawnSync } from 'node:child_process'; +import { mkdirSync, rmSync, writeFileSync, readFileSync, openSync, existsSync } from 'node:fs'; +import { connectRpc } from '../../tools/finality-attacks/lib/rpc.mjs'; +import { devAddress } from '../../tools/harness/lib/address.mjs'; + +const ROOT = new URL('../../', import.meta.url).pathname; +const FILE = `${ROOT}infra/fast-time/override-60x.json`; +const BIN = process.env.IGNEUM_V5_BIN || `${ROOT}vendor/igneum-node-class-v5/target/release`; +const IGNEUMD = process.env.IGNEUMD || `${BIN}/igneumd`; +const CPU_MINER = process.env.IGNEUM_MINER || `${BIN}/igneum-miner`; +const IGNEUM_POW = process.env.IGNEUM_POW || `${ROOT}igneum-pow/target/release/igneum-pow`; +const TMP = process.env.IGNEUM_V5_TMP || '/tmp/igneum-fast-time-v5s'; +const BASE = 29760, SUFFIX = 976; +const NEVER = '18446744073709551615'; +const args = process.argv.slice(2); +const flag = (name, dflt) => { const i = args.indexOf(`--${name}`); return i >= 0 ? +args[i + 1] : dflt; }; +const sflag = (name) => { const i = args.indexOf(`--${name}`); return i >= 0 ? args[i + 1] : null; }; +const actflag = (name, dflt) => { const v = sflag(name); if (v == null) return dflt; return v === 'never' ? null : +v; }; +const GENESIS_BITS = flag('genesis-bits', 0x1f010000); +const SECS = flag('secs', 780); +const EPOCHS = flag('epochs', 11); +const FLOOR = actflag('floor', 100000); +const V4_FLOOR = actflag('v4-floor', 120); +const V3_ACTIVATION = actflag('v3-activation', 60); +const WINDOW = flag('window', 60); +const STALE = sflag('stale') == null ? null : +sflag('stale'); +const STATELESS_NODE = !args.includes('--no-stateless-node'); +const WINDOWS = 7; +const SIGNAL = (sflag('signal') || '5,5,5').split(',').map(Number); +const EXPECT = sflag('expect') || 'flip'; +if (!['flip', 'no-flip', 'floor'].includes(EXPECT) || SIGNAL.length !== 3) { console.error('usage: --signal a,b,c --expect flip|no-flip|floor'); process.exit(2); } +const started = []; +const log = (...a) => console.log(new Date().toISOString().slice(11, 23), ...a); +const sleep = (ms) => new Promise(r => setTimeout(r, ms)); +for (const b of [IGNEUMD, CPU_MINER]) if (!existsSync(b)) { console.error(`missing ${b}`); process.exit(2); } + +rmSync(TMP, { recursive: true, force: true }); mkdirSync(TMP, { recursive: true }); +const baseText = readFileSync(FILE, 'utf8'); +const field = (name) => { const m = new RegExp(`"${name}":\\s*([0-9]+)`).exec(baseText); return m ? +m[1] : undefined; }; +const EPOCH = field('pow_epoch_blocks'); +const LEAD = field('pow_epoch_lead'); +const DAY_MS = field('pow_day_ms'); +const FIRST_V3_EPOCH = V3_ACTIVATION == null ? null : Math.ceil(V3_ACTIVATION / EPOCH); +const V4_FLOOR_EPOCH = V4_FLOOR == null ? null : Math.ceil(V4_FLOOR / EPOCH); +const FLOOR_EPOCH = FLOOR == null ? null : Math.ceil(FLOOR / EPOCH); +// the first epoch whose seed block (the last chain block below L*e - lead) can have DAA >= 7 x WINDOW: L*e - lead - 1 >= 7W, +// and whose previous epoch is already v4 (one step per epoch) +let FIRST_FULL_EPOCH = 0; +while (FIRST_FULL_EPOCH * EPOCH - LEAD - 1 < WINDOWS * WINDOW) FIRST_FULL_EPOCH++; +if (V4_FLOOR_EPOCH != null && FIRST_FULL_EPOCH <= V4_FLOOR_EPOCH) FIRST_FULL_EPOCH = V4_FLOOR_EPOCH + 1; +export function mergeOverrideText(text, fields) { + let out = text; + for (const k of Object.keys(fields)) out = out.replace(new RegExp(`\\s*"${k}":\\s*[^,}\\n]+,?`), ''); + const extra = Object.entries(fields).map(([k, v]) => `"${k}": ${typeof v === 'string' && !/^\d+$/.test(v) ? JSON.stringify(v) : v}`).join(', '); + return out.replace(/,?\s*}\s*$/, `,\n ${extra}\n}\n`); +} +const asText = (v) => v == null ? NEVER : String(v); +const override = `${TMP}/override.json`; +writeFileSync(override, mergeOverrideText(baseText, { genesis_bits: GENESIS_BITS, skip_proof_of_work: false, program_class_v3_activation_daa: asText(V3_ACTIVATION), program_class_v4_activation_daa: asText(V4_FLOOR), program_class_v4_signal_window_daa: String(WINDOW), program_class_v5_activation_daa: asText(FLOOR) })); +log(`signals ${SIGNAL.join('/')}, expect ${EXPECT}; v3 from ${V3_ACTIVATION ?? 'never'} (epoch ${FIRST_V3_EPOCH ?? 'none'}), v4 floor ${V4_FLOOR ?? 'never'} (epoch ${V4_FLOOR_EPOCH ?? 'none'}), v5 floor ${FLOOR ?? 'never'} (epoch ${FLOOR_EPOCH ?? 'none'}), window ${WINDOW} DAA x ${WINDOWS} (the first epoch with seven full windows after v4 is ${FIRST_FULL_EPOCH}); stale miner ${STALE ?? 'none'}; stateless node ${STATELESS_NODE ? 'n3' : 'none'}; ${EPOCH} DAA per epoch, lead ${LEAD}; run ${SECS} s or ${EPOCHS} epochs`); + +class Node { + constructor(i, connect = [], stateless = false) { + this.i = i; this.grpcPort = BASE + i * 10; this.p2pPort = BASE + i * 10 + 1; this.jsonPort = BASE + i * 10 + 2; this.execPort = BASE + i * 10 + 3; + this.connect = connect; this.dir = `${TMP}/n${i}`; this.logFile = `${this.dir}/node.log`; this.stateless = stateless; + } + get grpc() { return `grpc://127.0.0.1:${this.grpcPort}`; } + get execRpc() { return `http://127.0.0.1:${this.execPort}`; } + async start() { + mkdirSync(this.dir, { recursive: true }); + const a = ['--devnet', `--devnet-suffix=${SUFFIX}`, '--nodnsseed', '--disable-upnp', '--nologfiles', '--enable-unsynced-mining', '--utxoindex', + `--appdir=${this.dir}`, `--rpclisten=127.0.0.1:${this.grpcPort}`, `--rpclisten-json=127.0.0.1:${this.jsonPort}`, + `--listen=127.0.0.1:${this.p2pPort}`, `--override-params-file=${override}`, '--loglevel=info', '--yes']; + // class v5: each node's executor serves the state stream on its own loopback port; the stateless node has none + if (this.stateless) a.push('--evm-disable'); else a.push(`--evm-rpclisten=127.0.0.1:${this.execPort}`); + if (this.connect.length) a.push(`--connect=${this.connect.join(',')}`); else a.push('--outpeers=0'); + const out = openSync(this.logFile, 'a'); + // the node's own object byte: what its templates signal (the stateless node signals the highest byte too) + const byte = this.i < SIGNAL.length ? SIGNAL[this.i] : Math.max(...SIGNAL); + this.proc = spawn(IGNEUMD, a, { stdio: ['ignore', out, out], env: { ...process.env, IGNEUM_CLASS_SIGNAL: String(byte) } }); + started.push(this.proc); + await sleep(1200); + this.rpc = await connectRpc(`ws://127.0.0.1:${this.jsonPort}`); + log(`n${this.i} up pid ${this.proc.pid} json ${this.jsonPort} p2p ${this.p2pPort} exec ${this.stateless ? 'disabled' : this.execPort}, signals ${byte}`); + return this; + } + grepLog(re) { try { return readFileSync(this.logFile, 'utf8').split('\n').filter(l => re.test(l)); } catch { return []; } } +} +function miner(bin, argv, name, env = {}) { + const out = openSync(`${TMP}/${name}.log`, 'a'); + const p = spawn(bin, argv, { stdio: ['ignore', out, out], env: { ...process.env, ...env } }); + started.push(p); + return p; +} +async function stopAll() { + for (const p of started.reverse()) { try { p.kill('SIGINT'); } catch { } } + await sleep(1500); + for (const p of started) { try { p.kill('SIGKILL'); } catch { } } +} +process.on('SIGINT', async () => { await stopAll(); process.exit(130); }); +process.on('unhandledRejection', async (e) => { log(`FAILED: ${e?.stack || e}`); await stopAll(); process.exit(3); }); +const minerLog = (i) => { try { return readFileSync(`${TMP}/cpu${i}.log`, 'utf8').split('\n'); } catch { return []; } }; +const SIGNAL_LINE = /Program class v5 by miner signal: epoch (\d+) \(share (\d+) bps/; +const FLOOR_LINE = /Program class v5 from the override file/; +const REFUSAL_LINE = /class v5 needs the execution state/; +const WINDOW_LINE = /Program class v4 signal window from the override file/; +const OWN_LINE = /Program class signal from IGNEUM_CLASS_SIGNAL: this node signals object version (\d+)/; + +const t0 = Date.now(); +const since = () => ((Date.now() - t0) / 1000).toFixed(1); +const n0 = await new Node(0).start(); +const n1 = await new Node(1, [`127.0.0.1:${n0.p2pPort}`]).start(); +const n2 = await new Node(2, [`127.0.0.1:${n0.p2pPort}`]).start(); +const nodes = [n0, n1, n2]; +if (STATELESS_NODE) nodes.push(await new Node(3, [`127.0.0.1:${n0.p2pPort}`], true).start()); +const SIGNALLING = [n0, n1, n2]; +for (const n of nodes) log(`n${n.i}: ${n.grepLog(WINDOW_LINE).map(l => l.replace(/^.*?(Program class v4 signal window)/, '$1'))[0] || '(no window line)'} | ${n.grepLog(OWN_LINE).map(l => l.replace(/^.*?(this node signals)/, '$1'))[0] || '(no signal line)'}`); +log(`n0 digest: ${n0.grepLog(/Consensus params digest/).map(l => l.replace(/^.*?digest: /, '').slice(0, 16)).join(' ')}`); +nodes.forEach((n, i) => { + const extra = n.stateless ? [] : ['--exec-rpc', n.execRpc]; + if (STALE === i) extra.push('--freeze-state'); + miner(CPU_MINER, ['mine', n.grpc, '1', String(SECS), `cpu${i}`, '--engine', 'igneum-pow', '--payout-label', `cpu${i}`, '--status-secs', '30', '--no-vote', ...extra], `cpu${i}`, { IGNEUM_POW_DAY_MS: String(DAY_MS) }); +}); +const pay = devAddress('fast-time-v5s'); +const minerAcceptedAt = (i) => minerLog(i).filter(l => /ACCEPTED block/.test(l)).map(l => { const m = /^(\d+\.\d+) /.exec(l); return m ? +m[1] : null; }).filter(t => t != null); + +const epochs = new Map(); +let firstV4 = null, lastEpoch = -1, lastReport = 0, lastDaa = 0, endAt = null, flipWall = null; +const epochSeeds = new Map(); +const dayOfEpoch = new Map(); +const samples = []; +while (Date.now() - t0 < SECS * 1000) { + await sleep(1000); + let daa = null, epoch = null, cls = null, nextCls = null, eraSeed = null, bps = null, bps5 = null, win = null, sig = null, sigEpoch = null, seed = null, day = null; + try { + const t = await n0.rpc.call('getBlockTemplate', { payAddress: pay, extraData: [] }); + const pe = t.powEpoch || t.pow_epoch || {}; + daa = pe.virtualDaaScore ?? t.block?.header?.daaScore; epoch = pe.epochIndex; cls = pe.programClass; nextCls = pe.nextProgramClass; + eraSeed = pe.eraSeed; bps = pe.programClassV4SignalBps; bps5 = pe.programClassV5SignalBps; win = pe.programClassV4SignalWindowDaa; sig = pe.programClassSignal; sigEpoch = pe.programClassV5SignalEpoch; + seed = pe.epochSeed; day = Math.floor((+t.block?.header?.timestamp || 0) / DAY_MS); + } catch (e) { log(`template: ${e.message}`); } + if (epoch != null && epoch !== lastEpoch) { + epochs.set(epoch, { class: cls, firstSeenDaa: daa, at: +since(), eraSeed: eraSeed == null ? null : String(eraSeed), bps, bps5, signal_epoch: sigEpoch ?? null, day }); + if (seed != null) epochSeeds.set(epoch, String(seed)); + dayOfEpoch.set(epoch, day); + log(`epoch ${lastEpoch} -> ${epoch} at daa ${daa}, ${since()} s: template class ${cls}, next ${nextCls}, day ${day}, signal share at the sink v4 ${bps} v5 ${bps5} bps (window ${win}, this node signals ${sig}, v5 decided by signal at epoch ${sigEpoch ?? 'none'})`); + if (firstV4 == null && cls === 5) { firstV4 = { epoch, daa, at: +since() }; flipWall = Date.now(); log(`CLASS SWITCH: the template is class v5 from epoch ${epoch} (daa ${daa}) at ${since()} s wall`); } + lastEpoch = epoch; + } + lastDaa = daa ?? lastDaa; + if (Date.now() - lastReport > 15000) { + lastReport = Date.now(); + const counts = await Promise.all(nodes.map(async n => { try { const d = await n.rpc.call('getBlockDagInfo'); return `${d.blockCount}/${String(d.sink).slice(0, 8)}`; } catch { return '?'; } })); + log(`t=${since()} s daa ${daa} epoch ${epoch} class ${cls} signal v4 ${bps} v5 ${bps5} bps blocks/sink per node ${counts.join(' ')}`); + samples.push({ t: +since(), daa, epoch, class: cls, bps, bps5, nodes: counts }); + } + // the end: three epochs after a flip (two refreshes for the stale miner), or --epochs epochs when no flip is expected + if (firstV4 != null && daa != null && daa >= (firstV4.epoch + 3) * EPOCH) { endAt = +since(); break; } + if (firstV4 == null && daa != null && daa >= EPOCHS * EPOCH) { endAt = +since(); break; } +} +await sleep(3000); + +const dag = await Promise.all(nodes.map(async n => { try { return await n.rpc.call('getBlockDagInfo'); } catch (e) { return { error: e.message }; } })); +const genesis = dag[0].pruningPointHash; +async function allBlocks(n) { + const out = []; let low = genesis; const seen = new Set(); + for (let round = 0; round < 500; round++) { + const r = await n.rpc.call('getBlocks', { lowHash: low, includeBlocks: true, includeTransactions: false }); + const blocks = r.blocks || []; + let added = 0; + for (const b of blocks) { const h = b.verboseData?.hash || b.header?.hash; if (seen.has(h)) continue; seen.add(h); out.push({ hash: h, daa: +b.header.daaScore, version: +b.header.version, chain: !!b.verboseData?.isChainBlock }); added++; } + if (!blocks.length || added === 0) break; + low = (r.blockHashes || []).at(-1) || blocks.at(-1).verboseData?.hash; if (!low) break; + } + return out; +} +let blocks = []; +try { blocks = await allBlocks(n0); } catch (e) { log(`getBlocks: ${e.message}`); } +const BOUNDARY = firstV4 ? firstV4.epoch * EPOCH : Infinity; +const before = blocks.filter(b => b.daa < BOUNDARY), after = blocks.filter(b => b.daa >= BOUNDARY); +// the signal bytes on the chain: the share of blocks whose version high byte is 4 +const versionBytes = blocks.reduce((m, b) => { const v = b.version >> 8; m[v] = (m[v] || 0) + 1; return m; }, {}); +const signalShareOnChain = blocks.length ? Math.round(10000 * (blocks.filter(b => (b.version >> 8) >= 5).length) / blocks.length) : 0; + +const programs = new Map(); +for (const i of nodes.map(n => n.i)) for (const l of minerLog(i)) { + const m = /epoch seed ([0-9a-f]{64}) day (\d+) \(daa (\d+)\): program and 256 MiB cache ready in ([\d.]+) ms; class (v\d) program id ([0-9a-f]{16})/.exec(l); + if (!m) continue; + const k = m[1]; const e = programs.get(k) || { seed: k.slice(0, 16), epoch: Math.floor(+m[3] / EPOCH), class: m[5], id: m[6], miners: new Set() }; + if (e.id !== m[6] || e.class !== m[5]) e.disagree = true; + e.miners.add(i); programs.set(k, e); +} +const programRows = [...programs.values()].sort((a, b) => a.epoch - b.epoch).map(p => ({ epoch: p.epoch, class: p.class, program_id: p.id, seed: p.seed, miners: p.miners.size, disagree: !!p.disagree })); +// the state stream after an epoch's seed block, from node 0's exec RPC, for the CLI's --state +async function streamFile(epoch) { + const seed = epochSeeds.get(epoch); + if (!seed) return null; + try { + const r = await fetch(n0.execRpc, { method: 'POST', headers: { 'content-type': 'application/json' }, body: JSON.stringify({ jsonrpc: '2.0', id: 1, method: 'igneum_getPowStateLeaves', params: [seed] }) }); + const j = await r.json(); + if (!j.result?.streamHex) { log(`stream for epoch ${epoch}: ${JSON.stringify(j.error || j).slice(0, 200)}`); return null; } + const path = `${TMP}/state-e${epoch}.bin`; + writeFileSync(path, Buffer.from(j.result.streamHex.slice(2), 'hex')); + return { path, root: j.result.stateRoot, records: j.result.records, block: j.result.block }; + } catch (e) { log(`stream for epoch ${epoch}: ${e.message}`); return null; } +} +function cliId(seedHex, eraHex, cls, statePath) { + if (!existsSync(IGNEUM_POW)) return null; + const extra = statePath ? ['--state', statePath] : []; + const r = spawnSync(IGNEUM_POW, ['show', '--epoch-hex', seedHex, '--program-class', cls, '--era-hex', eraHex, ...extra], { encoding: 'utf8' }); + const m = /program id ([0-9a-f]{16})/.exec(r.stdout || ''); + return m ? m[1] : null; +} +const idRows = []; +const streams = {}; +for (const [k, e] of programs) { + if (e.class !== 'v5') continue; + const era = epochs.get(e.epoch)?.eraSeed; + const st = await streamFile(e.epoch); + if (st) streams[e.epoch] = { root: st.root, records: st.records, block: st.block }; + idRows.push({ epoch: e.epoch, seed: e.seed, miners_id: e.id, miners: e.miners.size, cli_v4: era ? cliId(k, era, 'v4') : null, cli_v5: era && st ? cliId(k, era, 'v5', st.path) : null, state_root: st?.root ?? null, state_records: st?.records ?? null }); +} +// every node's stream for the flip epoch must carry the same root (the serialisation agrees across nodes) +const rootsAtFlip = []; +if (firstV4) for (const n of nodes) { + if (n.stateless) continue; + try { + const r = await fetch(n.execRpc, { method: 'POST', headers: { 'content-type': 'application/json' }, body: JSON.stringify({ jsonrpc: '2.0', id: 1, method: 'igneum_getPowStateLeaves', params: [epochSeeds.get(firstV4.epoch)] }) }); + const j = await r.json(); + rootsAtFlip.push(j.result?.stateRoot ?? `error: ${JSON.stringify(j.error || j).slice(0, 120)}`); + } catch (e) { rootsAtFlip.push(`error: ${e.message}`); } +} +const minerIdx = nodes.map(n => n.i); +const accepted = minerIdx.map(i => minerLog(i).filter(l => /ACCEPTED block/.test(l)).length); +const rejectedMiner = minerIdx.map(i => minerLog(i).filter(l => /rejected nonce=|submit error/.test(l))); +const rejectedNode = SIGNALLING.map(n => n.grepLog(/PoW rejected|Rejected block|rejected block/i)); +const signalLines = SIGNALLING.map(n => n.grepLog(SIGNAL_LINE).map(l => l.replace(/^.*?(Program class v5 by miner signal)/, '$1'))[0] || null); +const signalEpochs = signalLines.map(l => { const m = l && SIGNAL_LINE.exec(l); return m ? +m[1] : null; }); +const signalShares = signalLines.map(l => { const m = l && SIGNAL_LINE.exec(l); return m ? +m[2] : null; }); +const floorLines = SIGNALLING.map(n => n.grepLog(FLOOR_LINE).map(l => l.replace(/^.*?(Program class v5 from)/, '$1'))[0] || null); +// the stateless node: its miner's blocks accepted after the flip (must be 0) and its refusal lines +const statelessNode = nodes.find(n => n.stateless) || null; +const flipAtMs = flipWall; +// accepted-after-the-flip counts read the node-side acceptance instead: each miner's ACCEPTED lines after the flip's +// wall time, by the line's own clock (the miner prints `HH:MM:SS` UTC at the start of every line) +function acceptedAfter(i, wallMs) { + if (wallMs == null) return 0; + const since0 = new Date(wallMs).toISOString().slice(11, 19); + let count = 0; + for (const l of minerLog(i)) { if (!/ACCEPTED block/.test(l)) continue; const m = /^(\d{2}:\d{2}:\d{2})/.exec(l); if (m && m[1] >= since0) count++; } + return count; +} +const statelessAcceptedAfterFlip = statelessNode ? acceptedAfter(statelessNode.i, flipAtMs) : 0; +const statelessRefusals = statelessNode ? statelessNode.grepLog(REFUSAL_LINE).length : 0; +// the stale miner: accepted after the FIRST REFRESH after the flip (one epoch later), and its rejections +const refreshWall = flipAtMs == null ? null : flipAtMs + EPOCH * 1000; +const staleAcceptedAfterRefresh = STALE == null ? null : acceptedAfter(STALE, refreshWall); +const staleRejected = STALE == null ? null : rejectedMiner[STALE].length; +const daysSeen = [...new Set([...dayOfEpoch.values()])]; +const sinks = dag.map(d => String(d.sink || '?').slice(0, 16)); +const counts = dag.map(d => d.blockCount ?? '?'); +const maxEpochSeen = Math.max(-1, ...epochs.keys()); +const classesSeen = [...epochs.values()].map(e => e.class); + +const common = { + zero_rejected_by_miners: rejectedMiner.every(r => r.length === 0), + zero_rejected_by_nodes: rejectedNode.every(r => r.length === 0), + sinks_agree: new Set(sinks).size === 1, + block_counts_agree: new Set(counts.map(String)).size === 1, + miners_agree_on_every_program: programRows.every(p => !p.disagree), + window_line_on_every_node: nodes.every(n => n.grepLog(WINDOW_LINE).length > 0), + every_node_signals_its_byte: nodes.every((n, i) => n.grepLog(OWN_LINE).some(l => +OWN_LINE.exec(l)[1] === SIGNAL[i])), + // every block's byte is one of the nodes' (genesis, made before any node, is the one byte-0 block) + chain_carries_the_bytes: blocks.length > 0 && Object.keys(versionBytes).every(v => SIGNAL.includes(+v) || +v === Math.max(...SIGNAL) || (+v === 0 && versionBytes[v] === 1)), + // class v5: every executing node served the same state root for the flip epoch's seed block + state_roots_agree_across_nodes: !firstV4 || (rootsAtFlip.length === SIGNALLING.length && rootsAtFlip.every(r => r && !r.startsWith('error')) && new Set(rootsAtFlip).size === 1), +}; +let checks; +if (EXPECT === 'flip') { + checks = { + ...common, + template_switched_to_v5: firstV4 != null, + switched_at_the_first_full_window_epoch: firstV4 != null && firstV4.epoch === FIRST_FULL_EPOCH, + switched_before_the_floor: firstV4 != null && (FLOOR_EPOCH == null || firstV4.epoch < FLOOR_EPOCH), + signal_line_on_every_node_same_epoch: signalEpochs.every(e => e != null) && new Set(signalEpochs).size === 1 && signalEpochs[0] === (firstV4 && firstV4.epoch), + signal_share_at_or_above_threshold: signalShares.every(s => s != null && s >= 9500), + blocks_on_both_sides: before.length > 0 && after.length > 0, + v5_ids_equal_the_cli_v5_id: idRows.length > 0 && idRows.every(r => r.cli_v5 != null && r.cli_v5 === r.miners_id && r.miners === 3), + v5_ids_differ_from_the_same_seed_v4_id: idRows.length > 0 && idRows.every(r => r.cli_v4 != null && r.cli_v4 !== r.miners_id), + a_v5_epoch_per_window_refresh: idRows.length >= 2 && new Set(idRows.map(r => r.state_root)).size >= 1, + stateless_node_mines_nothing_after_the_flip: !statelessNode || (statelessAcceptedAfterFlip === 0 && statelessRefusals > 0), + stale_miner_falls_off_at_the_first_refresh: STALE == null || (staleAcceptedAfterRefresh === 0 && staleRejected > 0), + }; +} else if (EXPECT === 'no-flip') { + checks = { + ...common, + template_never_v5: firstV4 == null && !classesSeen.includes(5), + no_signal_line_on_any_node: signalLines.every(l => l == null), + ran_the_epochs: maxEpochSeen >= EPOCHS - 1, + v4_programs_seen: programRows.some(p => p.class === 'v4'), + signal_share_under_threshold_on_chain_or_the_object_off: signalShareOnChain < 9500 || FLOOR == null, + }; +} else { + checks = { + ...common, + template_switched_to_v5: firstV4 != null, + switched_at_the_floor_epoch: firstV4 != null && firstV4.epoch === FLOOR_EPOCH, + no_signal_line_on_any_node: signalLines.every(l => l == null), + floor_line_names_the_floor_epoch: floorLines.every(l => l && l.includes(`the floor at epoch ${FLOOR_EPOCH} `)), + blocks_on_both_sides: before.length > 0 && after.length > 0, + v5_ids_equal_the_cli_v5_id: idRows.length > 0 && idRows.every(r => r.cli_v5 != null && r.cli_v5 === r.miners_id && r.miners === 3), + stateless_node_mines_nothing_after_the_flip: !statelessNode || (statelessAcceptedAfterFlip === 0 && statelessRefusals > 0), + }; +} +const pass = Object.values(checks).every(Boolean); +const summary = { + pass, expect: EXPECT, signals: SIGNAL, checks, window: WINDOW, windows: WINDOWS, floor: FLOOR ?? 'never', v4_floor: V4_FLOOR ?? 'never', v3_activation: V3_ACTIVATION ?? 'never', epoch_blocks: EPOCH, lead: LEAD, first_full_window_epoch: FIRST_FULL_EPOCH, floor_epoch: FLOOR_EPOCH, + stale_miner: STALE, stale_accepted_after_first_refresh: staleAcceptedAfterRefresh, stale_rejected: staleRejected, stateless_node: statelessNode ? statelessNode.i : null, stateless_accepted_after_flip: statelessAcceptedAfterFlip, stateless_refusal_lines: statelessRefusals, + state_roots_at_flip: rootsAtFlip, streams, days_seen: daysSeen, day_boundary_crossed: daysSeen.length > 1, + node: IGNEUMD, miner: CPU_MINER, template_switch: firstV4, run_ended_at_s: endAt, final_daa: lastDaa, max_epoch_seen: maxEpochSeen, + epochs: Object.fromEntries([...epochs.entries()].map(([k, v]) => [k, v])), + blocks: { total: blocks.length, before_boundary: before.length, after_boundary: after.length, version_bytes: versionBytes, signal_share_bps_on_chain: signalShareOnChain }, + programs: programRows, program_id_rows: idRows, accepted_per_miner: accepted, + rejected_by_miners: rejectedMiner.map(r => r.length), rejected_by_nodes: rejectedNode.map(r => r.length), + sinks, block_counts: counts, signal_lines: signalLines, floor_lines: floorLines, samples, +}; +writeFileSync(`${TMP}/summary.json`, JSON.stringify(summary, null, 2)); +log(`SUMMARY ${pass ? 'PASS' : 'FAIL'} (expect ${EXPECT}, signals ${SIGNAL.join('/')}, stale ${STALE ?? 'none'}): ${firstV4 ? `v5 from epoch ${firstV4.epoch} at DAA ${firstV4.daa}` : 'no v5 epoch'}; epochs seen ${[...epochs.entries()].map(([e, v]) => `e${e}:${v.class}:${v.bps5}bps:d${v.day}`).join(' ')}; chain bytes ${JSON.stringify(versionBytes)} (${signalShareOnChain} bps at byte 5+); blocks ${before.length} / ${after.length}; rejected miners ${rejectedMiner.map(r => r.length).join('/')} nodes ${rejectedNode.map(r => r.length).join('/')}; sinks ${sinks.join(' ')} (${checks.sinks_agree ? 'agree' : 'DIFFER'}); counts ${counts.join('/')}; signal lines ${signalLines.filter(Boolean).length}/3 (epochs ${signalEpochs.join('/')}, shares ${signalShares.join('/')}); floor lines ${floorLines.filter(Boolean).length}/3; state roots at the flip ${rootsAtFlip.map(r => String(r).slice(0, 18)).join(' ')}; stateless node accepted after the flip ${statelessAcceptedAfterFlip} (refusal lines ${statelessRefusals}); stale miner accepted after the first refresh ${staleAcceptedAfterRefresh ?? 'n/a'} rejected ${staleRejected ?? 'n/a'}; days ${daysSeen.join('/')}`); +for (const r of idRows) log(`PROGRAM ID epoch ${r.epoch} seed ${r.seed}: miners ${r.miners_id} (${r.miners} of 3) cli v5 ${r.cli_v5} cli v4 ${r.cli_v4}; state root ${r.state_root} (${r.state_records} records)`); +for (const [k, v] of Object.entries(checks)) if (!v) log(`FAILED CHECK ${k}`); +log(`summary: ${TMP}/summary.json`); +await stopAll(); +process.exit(pass ? 0 : 1); diff --git a/infra/fast-time/override-60x.json b/infra/fast-time/override-60x.json index 75a0e2cc5..2c4ecca85 100644 --- a/infra/fast-time/override-60x.json +++ b/infra/fast-time/override-60x.json @@ -59,6 +59,7 @@ "program_class_v3_activation_daa": 18446744073709551615, "program_class_v4_activation_daa": 18446744073709551615, "program_class_v4_signal_window_daa": 120, + "program_class_v5_activation_daa": 18446744073709551615, "latency_ladder": [{"reps": 27, "admissible": true}, {"reps": 35, "admissible": true}, {"reps": 53, "admissible": true}, {"reps": 88, "admissible": false}, {"reps": 173, "admissible": false}, {"reps": 267, "admissible": false}], "latency_ladder_activation_daa": 18446744073709551615, "latency_ladder_window_daa": 120, diff --git a/proto-newpow/class-v5/bench.cu b/proto-newpow/class-v5/bench.cu new file mode 100644 index 000000000..674d281c8 --- /dev/null +++ b/proto-newpow/class-v5/bench.cu @@ -0,0 +1,247 @@ +// Class v5 bench (docs/design/class-v5-stored-state.md section 7). TEST HARNESS ONLY: no pool, no network, no wallet. +// +// One pack directory in (program.h, vectors.h, memhard.h, kernel.cu, and leaves.bin for a class v5 pack), the shape of +// proto-newpow/state-dataset/bench.cu: cache fill on the GPU (twice, CUDA events), the dataset build (twice, the +// second pass reported; a class v5 pack's build takes the leaves of leaves.bin, IGNEUM_STATE_LEAVES x 16 words), the +// pack's dataset self-test (head, last, 64 samples) and its three vector warps, a warm-up batch of 2^24 at base +// nonce 0 (fingerprinted: FNV-1a 64 over the little-endian u64 outputs), N timed batches, then a power window where +// nvidia-smi samples at 1 Hz while the hash kernel runs back to back. The same binary shape for the class v4 control +// pack and the class v5 pack, so the two rows differ in the pack alone. +// +// bench [--batches 10] [--batch-log2 24] [--power-seconds 20] [--device 0] [--leaves ] +// +// Build (in the pack directory): nvcc -O3 -std=c++17 -arch=sm_89 -Xcompiler -pthread -I. -o bench ../bench.cu kernel.cu + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "program.h" +#include "vectors.h" +#include "memhard.h" + +#define CUDA_CHECK(call) do { cudaError_t err_ = (call); if (err_ != cudaSuccess) { \ + std::fprintf(stderr, "CUDA error: %s (%d)\n at %s:%d\n in %s\n", cudaGetErrorString(err_), (int)err_, __FILE__, __LINE__, #call); \ + std::exit(2); } } while (0) + +static uint64_t fnv1a64(const void* p, size_t n) { + const uint8_t* b = (const uint8_t*)p; + uint64_t h = 0xcbf29ce484222325ull; + for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; } + return h; +} +static float eventMs(cudaEvent_t a, cudaEvent_t b) { float ms = 0.f; CUDA_CHECK(cudaEventElapsedTime(&ms, a, b)); return ms; } + +struct Opts { int batches = 10; int batchLog2 = 24; int powerSeconds = 20; int device = 0; std::string leaves = "leaves.bin"; }; +static Opts parse(int argc, char** argv) { + Opts o; + for (int i = 1; i < argc; ++i) { + std::string a = argv[i]; + auto need = [&](int n) { if (i + n >= argc) { std::printf("%s needs %d argument(s)\n", a.c_str(), n); std::exit(2); } }; + if (a == "--batches") { need(1); o.batches = std::atoi(argv[++i]); } + else if (a == "--batch-log2") { need(1); o.batchLog2 = std::atoi(argv[++i]); } + else if (a == "--power-seconds") { need(1); o.powerSeconds = std::atoi(argv[++i]); } + else if (a == "--device") { need(1); o.device = std::atoi(argv[++i]); } + else if (a == "--leaves") { need(1); o.leaves = argv[++i]; } + else { std::printf("unknown argument %s\n", argv[i]); std::exit(2); } + } + return o; +} + +// nvidia-smi sampler: one line per second on its own thread until `timeout` ends the process. +struct Sampler { + std::vector power, sm, mem; + std::mutex m; + std::thread t; + void start(int device, int seconds) { + t = std::thread([this, device, seconds]() { + char cmd[256]; + std::snprintf(cmd, sizeof(cmd), "timeout %d nvidia-smi -i %d --query-gpu=power.draw,clocks.sm,clocks.mem --format=csv,noheader,nounits -l 1 2>/dev/null", seconds + 2, device); + FILE* f = popen(cmd, "r"); + if (!f) return; + char line[256]; + while (std::fgets(line, sizeof(line), f)) { + double p = 0, s = 0, mm = 0; + if (std::sscanf(line, "%lf, %lf, %lf", &p, &s, &mm) == 3) { std::lock_guard g(m); power.push_back(p); sm.push_back(s); mem.push_back(mm); } + } + pclose(f); + }); + } + void join() { if (t.joinable()) t.join(); } +}; + +int main(int argc, char** argv) { + Opts o = parse(argc, argv); +#ifdef IGNEUM_STATE_LEAVES + const bool v5 = true; +#else + const bool v5 = false; +#endif + std::printf("class-v5 bench pack \"%s\" class %s generator %d (test harness: no pool, no network, no wallet)\n", IGNEUM_SEED_STRING, v5 ? "v5" : "control", (int)IGNEUM_GENERATOR); + CUDA_CHECK(cudaSetDevice(o.device)); + cudaDeviceProp prop; + std::memset(&prop, 0, sizeof(prop)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, o.device)); + int drv = 0, rt = 0; + CUDA_CHECK(cudaDriverGetVersion(&drv)); + CUDA_CHECK(cudaRuntimeGetVersion(&rt)); + std::printf("GPU: %s (%d SMs, cc %d.%d, %.0f MiB), CUDA driver %d.%d runtime %d.%d\n", prop.name, prop.multiProcessorCount, prop.major, prop.minor, + (double)prop.totalGlobalMem / 1048576.0, drv / 1000, (drv % 100) / 10, rt / 1000, (rt % 100) / 10); + int regs = 0, bps = 0; + CUDA_CHECK(igneum_hash_info(®s, &bps, 1u)); + std::printf("hash kernel: %d registers/thread, %d resident blocks/SM at 1 warp/block\n", regs, bps); + size_t free0 = 0, total = 0; + CUDA_CHECK(cudaMemGetInfo(&free0, &total)); + std::printf("device memory at start: %.0f MiB used of %.0f MiB\n", (double)(total - free0) / 1048576.0, (double)total / 1048576.0); + + cudaEvent_t e0, e1; + CUDA_CHECK(cudaEventCreate(&e0)); + CUDA_CHECK(cudaEventCreate(&e1)); + + // ---- cache + const uint32_t cacheWords = 1u << IGNEUM_CACHE_LOG2_WORDS; + uint32_t* dCache = nullptr; + CUDA_CHECK(cudaMalloc(&dCache, (size_t)cacheWords * 4u)); + float cacheFill[2] = {0, 0}; + for (int pass = 0; pass < 2; ++pass) { + CUDA_CHECK(cudaEventRecord(e0)); + CUDA_CHECK(igneum_launch_cache_fill(dCache, IGNEUM_CACHE_SEGMENTS)); + CUDA_CHECK(cudaEventRecord(e1)); + CUDA_CHECK(cudaEventSynchronize(e1)); + cacheFill[pass] = eventMs(e0, e1); + } + std::printf("cache fill (GPU): %.2f ms first, %.2f ms second\n", cacheFill[0], cacheFill[1]); + { + std::vector head(16), last(16); + CUDA_CHECK(cudaMemcpy(head.data(), dCache, 64, cudaMemcpyDeviceToHost)); + CUDA_CHECK(cudaMemcpy(last.data(), dCache + (cacheWords - 16), 64, cudaMemcpyDeviceToHost)); + bool ok = std::memcmp(head.data(), IGNEUM_CACHE_HEAD, 64) == 0 && std::memcmp(last.data(), IGNEUM_CACHE_LAST, 64) == 0; + std::printf("cache head and last 16 words against the pack: %s\n", ok ? "PASS" : "FAIL"); + } + + // ---- leaves (class v5) + uint32_t* dLeaves = nullptr; + uint32_t nLeaves = 0; +#ifdef IGNEUM_STATE_LEAVES + { + FILE* f = std::fopen(o.leaves.c_str(), "rb"); + if (!f) { std::printf("FAIL: cannot open %s (the pack's leaves.bin)\n", o.leaves.c_str()); return 2; } + std::fseek(f, 0, SEEK_END); + long n = std::ftell(f); + std::fseek(f, 0, SEEK_SET); + std::vector h((size_t)n / 4); + if (std::fread(h.data(), 1, (size_t)n, f) != (size_t)n) { std::printf("FAIL: short read of %s\n", o.leaves.c_str()); return 2; } + std::fclose(f); + nLeaves = (uint32_t)(h.size() / 16); + uint64_t fnv = fnv1a64(h.data(), h.size() * 4); + std::printf("leaves: %u x 64 B from %s (%ld bytes), FNV-1a 64 %016llx against the pack's %016llx: %s; state root %s, chain block %s (%llu)\n", + nLeaves, o.leaves.c_str(), n, (unsigned long long)fnv, (unsigned long long)IGNEUM_STATE_LEAVES_FNV64, + fnv == IGNEUM_STATE_LEAVES_FNV64 && nLeaves == IGNEUM_STATE_LEAVES ? "PASS" : "FAIL", IGNEUM_STATE_ROOT_HEX, IGNEUM_STATE_BLOCK_HEX, (unsigned long long)IGNEUM_STATE_BLOCK_NUMBER); + CUDA_CHECK(cudaMalloc(&dLeaves, h.size() * 4)); + CUDA_CHECK(cudaMemcpy(dLeaves, h.data(), h.size() * 4, cudaMemcpyHostToDevice)); + } +#endif + + // ---- dataset build, twice + const uint32_t words = 1u << IGNEUM_DATASET_LOG2; + const uint32_t nItems = words >> 4; + uint32_t* dDs = nullptr; + CUDA_CHECK(cudaMalloc(&dDs, (size_t)words * 4u)); + float build[2] = {0, 0}; + for (int pass = 0; pass < 2; ++pass) { + CUDA_CHECK(cudaEventRecord(e0)); +#ifdef IGNEUM_STATE_LEAVES + CUDA_CHECK(igneum_launch_build(dDs, dCache, dLeaves, nLeaves, nItems)); +#else + CUDA_CHECK(igneum_launch_build(dDs, dCache, nItems)); +#endif + CUDA_CHECK(cudaEventRecord(e1)); + CUDA_CHECK(cudaEventSynchronize(e1)); + build[pass] = eventMs(e0, e1); + } + std::printf("dataset build (GPU): %.2f ms first, %.2f ms second (%u items, %u MiB)\n", build[0], build[1], nItems, words >> 18); + size_t free1 = 0; + CUDA_CHECK(cudaMemGetInfo(&free1, &total)); + std::printf("device memory after the build: %.0f MiB used\n", (double)(total - free1) / 1048576.0); + { + std::vector head(16); + uint32_t last = 0; + CUDA_CHECK(cudaMemcpy(head.data(), dDs, 64, cudaMemcpyDeviceToHost)); + CUDA_CHECK(cudaMemcpy(&last, dDs + IGNEUM_DS_LAST_INDEX, 4, cudaMemcpyDeviceToHost)); + int sampleOk = 0; + for (int i = 0; i < IGNEUM_DS_SAMPLES; ++i) { + uint32_t v = 0; + CUDA_CHECK(cudaMemcpy(&v, dDs + IGNEUM_DS_SAMPLE_INDEX[i], 4, cudaMemcpyDeviceToHost)); + if (v == IGNEUM_DS_SAMPLE_VALUE[i]) ++sampleOk; + } + std::printf("dataset self-test: head %s, last %s, samples %d of %d\n", std::memcmp(head.data(), IGNEUM_DS_HEAD, 64) == 0 ? "PASS" : "FAIL", last == IGNEUM_DS_LAST ? "PASS" : "FAIL", sampleOk, (int)IGNEUM_DS_SAMPLES); + } +#ifdef IGNEUM_STATE_LEAVES + if (dLeaves) { CUDA_CHECK(cudaFree(dLeaves)); dLeaves = nullptr; } +#endif + + // ---- the three vector warps + const uint32_t nonces = 1u << o.batchLog2; + uint64_t* dOut = nullptr; + CUDA_CHECK(cudaMalloc(&dOut, (size_t)nonces * 8u)); + std::vector hOut(nonces); + int vecOk = 0; + for (int w = 0; w < IGNEUM_VEC_WARPS; ++w) { + CUDA_CHECK(igneum_launch_hash(dDs, dOut, IGNEUM_VEC_BASE[w], IGNEUM_MASK, 32u, 1u)); + CUDA_CHECK(cudaDeviceSynchronize()); + CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 32 * 8, cudaMemcpyDeviceToHost)); + if (std::memcmp(hOut.data(), IGNEUM_VEC_OUT[w], 32 * 8) == 0) ++vecOk; + } + std::printf("vector warps against the pack: %d of %d PASS\n", vecOk, (int)IGNEUM_VEC_WARPS); + + // ---- warm-up and fingerprint, then timed batches + const uint32_t blockWarps = 1u; + CUDA_CHECK(igneum_launch_hash(dDs, dOut, 0u, IGNEUM_MASK, nonces, blockWarps)); + CUDA_CHECK(cudaDeviceSynchronize()); + CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, (size_t)nonces * 8u, cudaMemcpyDeviceToHost)); + std::printf("fingerprint of 2^%d outputs at base 0: %016llx (lane 0 %016llx)\n", o.batchLog2, (unsigned long long)fnv1a64(hOut.data(), (size_t)nonces * 8u), (unsigned long long)hOut[0]); + double totalMs = 0; + for (int b = 0; b < o.batches; ++b) { + CUDA_CHECK(cudaEventRecord(e0)); + CUDA_CHECK(igneum_launch_hash(dDs, dOut, (uint32_t)(b + 1) * nonces, IGNEUM_MASK, nonces, blockWarps)); + CUDA_CHECK(cudaEventRecord(e1)); + CUDA_CHECK(cudaEventSynchronize(e1)); + totalMs += eventMs(e0, e1); + } + double rate = (double)o.batches * nonces / (totalMs / 1000.0) / 1e6; + std::printf("hash rate: %.3f MH/s over %d batches of 2^%d (GPU event time %.1f ms)\n", rate, o.batches, o.batchLog2, totalMs); + + // ---- power window + if (o.powerSeconds > 0) { + Sampler s; + s.start(o.device, o.powerSeconds); + auto t0 = std::chrono::steady_clock::now(); + uint64_t done = 0; + uint32_t base = 0; + while (std::chrono::duration(std::chrono::steady_clock::now() - t0).count() < o.powerSeconds) { + CUDA_CHECK(igneum_launch_hash(dDs, dOut, base, IGNEUM_MASK, nonces, blockWarps)); + CUDA_CHECK(cudaDeviceSynchronize()); + base += nonces; + done += nonces; + } + double secs = std::chrono::duration(std::chrono::steady_clock::now() - t0).count(); + s.join(); + std::lock_guard g(s.m); + double p = 0, sm = 0, mem = 0; + int n = 0; + for (size_t i = 10; i < s.power.size(); ++i) { p += s.power[i]; sm += s.sm[i]; mem += s.mem[i]; ++n; } + if (n > 0) { p /= n; sm /= n; mem /= n; } + double rateW = done / secs / 1e6; + std::printf("power window: %.1f s, %.3f MH/s sustained, %.1f W mean after the first 10 s (%d samples of %zu), SM %.0f MHz, mem %.0f MHz, %.3f MH/s per W, %.2f microjoules per hash\n", + secs, rateW, p, n, s.power.size(), sm, mem, p > 0 ? rateW / p : 0.0, p > 0 ? p / (rateW * 1e6) * 1e6 : 0.0); + } + std::printf("RESULT pack=%s class=%s build_ms=%.2f rate_mhs=%.3f vectors=%d/%d\n", IGNEUM_SEED_STRING, v5 ? "v5" : "control", build[1], rate, vecOk, (int)IGNEUM_VEC_WARPS); + return vecOk == IGNEUM_VEC_WARPS ? 0 : 1; +}