class v5: the fast-time gate infra/fast-time/class-v5-signal.mjs (the class v4 signal gate's shape: three signalling nodes with their executors on loopback ports, a fourth node with --evm-disable that must mine nothing after the flip, a --stale miner that keeps its first stream and must fall off at the first refresh, the v5 floor far away so the signal decides, the no-flip, flip, object-off, floor and known-failed cases, every node's state root for the flip epoch compared), the v5 floor in the 60x profile (never), proto-newpow/class-v5/bench.cu (one pack-generic 4090 bench: cache, leaves, build, the pack's self-test and vector warps, hash rate, watts), DatasetSource::refreshed_or_clone
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
6b17c79087
commit
41f055abc5
4 changed files with 619 additions and 0 deletions
|
|
@ -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<crate::state::StateLeaves>> {
|
||||
self.memhard().and_then(|m| m.leaves.as_ref())
|
||||
|
|
|
|||
362
infra/fast-time/class-v5-signal.mjs
Normal file
362
infra/fast-time/class-v5-signal.mjs
Normal file
|
|
@ -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 <daa>|never]
|
||||
// [--v4-floor 120] [--window 60] [--v3-activation 60] [--secs 780] [--epochs 11] [--stale <miner index>] [--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);
|
||||
|
|
@ -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,
|
||||
|
|
|
|||
247
proto-newpow/class-v5/bench.cu
Normal file
247
proto-newpow/class-v5/bench.cu
Normal file
|
|
@ -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 <leaves.bin>]
|
||||
//
|
||||
// Build (in the pack directory): nvcc -O3 -std=c++17 -arch=sm_89 -Xcompiler -pthread -I. -o bench ../bench.cu kernel.cu
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <chrono>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <thread>
|
||||
#include <mutex>
|
||||
|
||||
#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<double> 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<std::mutex> 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<uint32_t> 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<uint32_t> 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<uint32_t> 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<uint64_t> 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<double>(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<double>(std::chrono::steady_clock::now() - t0).count();
|
||||
s.join();
|
||||
std::lock_guard<std::mutex> 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;
|
||||
}
|
||||
Loading…
Reference in a new issue