igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-cuda/host.cu

915 lines
49 KiB
Text
Raw Blame History

// igneum-bench-cuda: host program for Igneum's random-program proof-of-work test harness on NVIDIA GPUs.
//
// TEST HARNESS ONLY. No pool, no network, no wallet, no mining protocol. It fills the dataset on the GPU,
// checks the GPU against vectors produced on the Mac (proto-metal), and times the kernel.
//
// C++17 plus the CUDA runtime API, nothing else. The kernel is compiled ahead of time by nvcc from
// packs/<seed>/kernel.cu, which was generated by proto-metal/igneum-bench --export-pack.
//
// Build (Linux, from proto-cuda/):
// nvcc -O3 -std=c++17 -arch=sm_120 -I packs/igneum-genesis -o igneum-bench-cuda-igneum-genesis host.cu packs/igneum-genesis/kernel.cu
// Windows and the -arch=native fallback are in README.md.
#include <cuda_runtime.h>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <chrono>
#include <string>
#include <vector>
#include <iostream>
#include "program.h"
#include "vectors.h"
// Packs written before 3 October 2026 have no dataset mode: they are closed-form (mode 0).
#ifndef IGNEUM_DATASET_MODE
#define IGNEUM_DATASET_MODE 0
#endif
#if IGNEUM_DATASET_MODE == 1
#include "memhard.h" // the memory-hard core, compiled here for the host reference (see proto-metal/MEMHARD.md)
#endif
#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 const uint32_t SEEDW[8] = IGNEUM_SEEDW_INIT;
#if IGNEUM_DATASET_MODE == 0
// Same closed form as ds_elem in kernel.cu and datasetElem in proto-metal/main.swift.
static uint32_t host_ds_elem(uint32_t i, uint32_t d0, uint32_t d1) {
uint32_t x = i ^ d0;
x *= 0x9E3779B1u; x ^= x >> 15;
x += d1;
x *= 0x85EBCA77u; x ^= x >> 13;
x *= 0xC2B2AE3Du; x ^= x >> 16;
return x;
}
#endif
static double wallMs();
#if IGNEUM_DATASET_MODE == 1
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;
}
// Memory-hard mode: the 256 MiB cache on the device (filled by the pack's kernel) and on the host (filled by the
// same mh_cache_segment text on one thread). Both are built once per process in setupCache(), compared word for
// word, and checked against the head, last line and FNV-1a 64 the Mac recorded in vectors.h.
static const uint32_t CACHE_WORDS_HOST = 1u << IGNEUM_CACHE_LOG2_WORDS;
static uint32_t* gCache = nullptr; // device
static std::vector<uint32_t> hCache; // host
static double gCacheFillFirstMs = 0, gCacheFillSecondMs = 0, gCacheHostMs = 0;
static bool gCachePass = false;
static bool setupCache() {
size_t bytes = (size_t)CACHE_WORDS_HOST * 4u;
CUDA_CHECK(cudaMalloc((void**)&gCache, bytes));
cudaEvent_t e0, e1;
CUDA_CHECK(cudaEventCreate(&e0));
CUDA_CHECK(cudaEventCreate(&e1));
for (int pass = 0; pass < 2; ++pass) {
CUDA_CHECK(cudaEventRecord(e0));
CUDA_CHECK(igneum_launch_cache_fill(gCache, IGNEUM_CACHE_SEGMENTS));
CUDA_CHECK(cudaEventRecord(e1));
CUDA_CHECK(cudaEventSynchronize(e1));
float msf = 0.f;
CUDA_CHECK(cudaEventElapsedTime(&msf, e0, e1));
if (pass == 0) gCacheFillFirstMs = msf; else gCacheFillSecondMs = msf;
}
CUDA_CHECK(cudaEventDestroy(e0));
CUDA_CHECK(cudaEventDestroy(e1));
std::printf("cache fill (GPU): %.2f ms first, %.2f ms second (%u chains x %u ChaCha blocks, %u MiB)\n",
gCacheFillFirstMs, gCacheFillSecondMs, (unsigned)IGNEUM_CACHE_SEGMENTS,
1u << IGNEUM_CACHE_SEGMENT_LOG2_LINES, (unsigned)(bytes >> 20));
hCache.assign(CACHE_WORDS_HOST, 0u);
double h0 = wallMs();
for (uint32_t seg = 0; seg < IGNEUM_CACHE_SEGMENTS; ++seg) mh_cache_segment(hCache.data(), seg);
gCacheHostMs = wallMs() - h0;
std::printf("cache fill (host, one thread): %.1f ms\n", gCacheHostMs);
std::vector<uint32_t> dev(CACHE_WORDS_HOST);
CUDA_CHECK(cudaMemcpy(dev.data(), gCache, bytes, cudaMemcpyDeviceToHost));
bool same = std::memcmp(dev.data(), hCache.data(), bytes) == 0;
uint64_t fnv = fnv1a64(hCache.data(), bytes);
bool fnvOk = (fnv == IGNEUM_CACHE_FNV64);
bool headOk = std::memcmp(hCache.data(), IGNEUM_CACHE_HEAD, 64) == 0;
bool lastOk = std::memcmp(hCache.data() + CACHE_WORDS_HOST - 16u, IGNEUM_CACHE_LAST, 64) == 0;
if (!same) {
for (uint32_t i = 0; i < CACHE_WORDS_HOST; ++i) if (dev[i] != hCache[i]) {
std::printf(" cache[%u]: gpu 0x%08x host 0x%08x (first difference)\n", i, dev[i], hCache[i]); break;
}
}
gCachePass = same && fnvOk && headOk && lastOk;
std::printf("cache check: %s (GPU == host all %u words %s, host FNV-1a 64 %016llx vs Mac %016llx %s, head 16 vs Mac %s, last line vs Mac %s)\n",
gCachePass ? "PASS" : "FAIL", CACHE_WORDS_HOST, same ? "PASS" : "FAIL",
(unsigned long long)fnv, (unsigned long long)IGNEUM_CACHE_FNV64, fnvOk ? "PASS" : "FAIL",
headOk ? "PASS" : "FAIL", lastOk ? "PASS" : "FAIL");
return gCachePass;
}
// dataset[w] derived on the host from the host cache through the pack's own mh_word (memhard.h), which carries the
// pack's item-to-word layout (era layout, 5 October 2026: the harness's former w >> 4 / w & 15 failed the random
// points of every interleaved pack while the Mac samples and vectors passed).
static uint32_t host_ds_word(uint32_t w) {
return mh_word(hCache.data(), w);
}
#endif
// ---------------------------------------------------------------------------------------------
// Options
struct Options {
int datasetMib = 1024;
int batchLog2 = 24;
int batches = 5;
int blockWarps = 1;
bool sweep = false;
int device = 0;
bool serve = false; // --serve: GPU worker for igneum-miner --worker (jobs on stdin), 3 October 2026
bool noPrepare = false; // --no-prepare: serve without the prepare command (ready line says "prepare 0"), to test the miner's fallback
};
static int packMib() { return (int)(((1ull << IGNEUM_DATASET_LOG2) * 4ull) >> 20); }
static void usage() {
std::printf(
"igneum-bench-cuda [--dataset-mib N] [--sweep] [--batch-log2 24] [--batches 5] [--block-warps 1] [--device 0]\n"
" --dataset-mib N dataset size in MiB, power of two (default 1024; vectors are only checked at %d MiB)\n"
" --sweep run 4, 64, 256, 512 and 1024 MiB in sequence (same sweep as the Mac)\n"
" --batch-log2 B nonces per batch = 2^B (default 24)\n"
" --batches N timed batches after one warm-up batch (default 5)\n"
" --block-warps W warps per thread block, 1..32 (default 1 = one warp per block, like the Metal run)\n"
" --device D CUDA device index (default 0)\n"
" --serve GPU worker for igneum-miner --worker: reads \"job ...\" lines on stdin, prints found/done lines\n"
" (needs the pack's kernel_bound.cu compiled in: build.bat adds it when the pack has one)\n"
" --no-prepare with --serve: no prepare support (the miner then falls back to exit 42 at a seed change)\n", packMib());
}
static bool isPow2(long long v) { return v > 0 && (v & (v - 1)) == 0; }
static Options parseArgs(int argc, char** argv) {
Options o;
for (int i = 1; i < argc; ++i) {
std::string a = argv[i];
auto next = [&](int& dst) {
if (i + 1 >= argc) { usage(); std::exit(2); }
dst = std::atoi(argv[++i]);
};
if (a == "--dataset-mib") next(o.datasetMib);
else if (a == "--batch-log2") next(o.batchLog2);
else if (a == "--batches") next(o.batches);
else if (a == "--block-warps") next(o.blockWarps);
else if (a == "--device") next(o.device);
else if (a == "--sweep") o.sweep = true;
else if (a == "--serve") o.serve = true;
else if (a == "--no-prepare") o.noPrepare = true;
else if (a == "-h" || a == "--help") { usage(); std::exit(0); }
else { std::printf("unknown argument %s\n", argv[i]); usage(); std::exit(2); }
}
if (!isPow2(o.datasetMib) || o.datasetMib < 1 || o.datasetMib > 16384) {
std::printf("--dataset-mib must be a power of two between 1 and 16384\n"); std::exit(2);
}
if (o.batchLog2 < 10 || o.batchLog2 > 28) { std::printf("--batch-log2 must be between 10 and 28\n"); std::exit(2); }
if (o.batches < 1) { std::printf("--batches must be at least 1\n"); std::exit(2); }
if (o.blockWarps < 1 || o.blockWarps > 32) { std::printf("--block-warps must be between 1 and 32\n"); std::exit(2); }
return o;
}
// ---------------------------------------------------------------------------------------------
// Helpers
static double wallMs() {
using namespace std::chrono;
return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
}
static int log2u32(uint32_t v) { int n = 0; while (v > 1u) { v >>= 1; ++n; } return n; }
static bool compareWarp(const uint64_t* got, const uint64_t* want, uint32_t base, const char* how) {
int bad = 0, first = -1;
for (int l = 0; l < 32; ++l) if (got[l] != want[l]) { if (first < 0) first = l; ++bad; }
if (bad == 0) {
std::printf("verify warp base %u (nonces %u..%u) %s: PASS\n", base, base, base + 31u, how);
} else {
std::printf("verify warp base %u (nonces %u..%u) %s: FAIL %d of 32 lanes differ, first lane %d: gpu=%016llx expected=%016llx\n",
base, base, base + 31u, how, bad, first,
(unsigned long long)got[first], (unsigned long long)want[first]);
}
return bad == 0;
}
struct SizeResult {
int mib = 0;
uint32_t words = 0;
double fillFirstMs = 0, fillSecondMs = 0;
bool dsPass = false;
bool vecChecked = false, vecPass = false;
double gpuMs = 0, wallMsTimed = 0;
double hashesPerSec = 0, gbps = 0;
};
// ---------------------------------------------------------------------------------------------
// One dataset size: fill, self-test, vectors, bench
static SizeResult runSize(const Options& o, int mib, uint64_t* dOut, uint32_t nonces) {
SizeResult r;
r.mib = mib;
uint64_t bytes = (uint64_t)mib << 20;
r.words = (uint32_t)(bytes / 4ull);
uint32_t mask = r.words - 1u;
bool atPackSize = (r.words == (1u << IGNEUM_DATASET_LOG2));
std::printf("\n=== dataset %d MiB (2^%d words, mask 0x%08x)%s ===\n", mib, log2u32(r.words), mask,
atPackSize ? "" : " [not the pack size: vectors skipped, dataset head and random points still checked]");
size_t freeB = 0, totalB = 0;
CUDA_CHECK(cudaMemGetInfo(&freeB, &totalB));
if ((uint64_t)freeB < bytes + (64ull << 20)) {
std::printf("FAIL: %llu MiB free on the device, need %d MiB for the dataset\n", (unsigned long long)(freeB >> 20), mib);
std::exit(2);
}
uint32_t* dDs = nullptr;
CUDA_CHECK(cudaMalloc((void**)&dDs, (size_t)bytes));
cudaEvent_t e0, e1;
CUDA_CHECK(cudaEventCreate(&e0));
CUDA_CHECK(cudaEventCreate(&e1));
// Fill (or build) twice: the Mac showed a first-touch cost on the first fill of a process.
for (int pass = 0; pass < 2; ++pass) {
CUDA_CHECK(cudaEventRecord(e0));
#if IGNEUM_DATASET_MODE == 1
CUDA_CHECK(igneum_launch_build(dDs, gCache, r.words / 16u));
#else
CUDA_CHECK(igneum_launch_fill(dDs, r.words, IGNEUM_DAY0, IGNEUM_DAY1));
#endif
CUDA_CHECK(cudaEventRecord(e1));
CUDA_CHECK(cudaEventSynchronize(e1));
float msf = 0.f;
CUDA_CHECK(cudaEventElapsedTime(&msf, e0, e1));
if (pass == 0) r.fillFirstMs = msf; else r.fillSecondMs = msf;
}
#if IGNEUM_DATASET_MODE == 1
std::printf("dataset build (memory-hard, from the cache): %.2f ms first, %.2f ms second -> %.1f M items/s, %.2f G cache-line reads/s (second, GPU time)\n",
r.fillFirstMs, r.fillSecondMs, (double)(r.words / 16u) / 1e6 / (r.fillSecondMs / 1000.0),
(double)(r.words / 16u) * (double)IGNEUM_ITEM_ROUNDS / 1e9 / (r.fillSecondMs / 1000.0));
#else
std::printf("dataset fill: %.2f ms first, %.2f ms second -> %.0f GB/s write (second, GPU time)\n",
r.fillFirstMs, r.fillSecondMs, (double)bytes / 1e9 / (r.fillSecondMs / 1000.0));
#endif
// Dataset self-test: head 16 (any size), element [MASK] (pack size only), 64 pseudo-random points vs the host
// formula (closed form) or the host derivation from the host cache (memory-hard), and the Mac's 64 sampled words.
{
uint32_t head[16];
CUDA_CHECK(cudaMemcpy(head, dDs, sizeof(head), cudaMemcpyDeviceToHost));
int badHead = 0;
for (int i = 0; i < 16; ++i) {
if (head[i] != IGNEUM_DS_HEAD[i]) {
if (badHead == 0) std::printf(" dataset[%d] = 0x%08x, expected 0x%08x\n", i, head[i], IGNEUM_DS_HEAD[i]);
++badHead;
}
}
bool lastOk = true;
const char* lastText = "skipped";
if (atPackSize) {
uint32_t last = 0;
CUDA_CHECK(cudaMemcpy(&last, dDs + IGNEUM_DS_LAST_INDEX, sizeof(last), cudaMemcpyDeviceToHost));
lastOk = (last == IGNEUM_DS_LAST);
lastText = lastOk ? "PASS" : "FAIL";
if (!lastOk) std::printf(" dataset[%u] = 0x%08x, expected 0x%08x\n", IGNEUM_DS_LAST_INDEX, last, IGNEUM_DS_LAST);
}
int badRnd = 0;
uint64_t s = 0x9E3779B97F4A7C15ull ^ (uint64_t)r.words;
for (int k = 0; k < 64; ++k) {
s += 0x9E3779B97F4A7C15ull;
uint64_t z = s;
z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ull;
z = (z ^ (z >> 27)) * 0x94D049BB133111EBull;
z ^= z >> 31;
uint32_t idx = (uint32_t)z & mask;
uint32_t v = 0;
CUDA_CHECK(cudaMemcpy(&v, dDs + idx, sizeof(v), cudaMemcpyDeviceToHost));
#if IGNEUM_DATASET_MODE == 1
uint32_t want = host_ds_word(idx);
#else
uint32_t want = host_ds_elem(idx, IGNEUM_DAY0, IGNEUM_DAY1);
#endif
if (v != want) {
if (badRnd == 0) std::printf(" dataset[%u] = 0x%08x, host %s 0x%08x\n", idx, v, IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", want);
++badRnd;
}
}
// The Mac's sampled words: every sample whose index lies inside this dataset size (items are the same
// at every size, the smaller dataset is a prefix of the larger one).
int badSample = 0, nSample = 0;
#ifdef IGNEUM_DS_SAMPLES
for (int k = 0; k < IGNEUM_DS_SAMPLES; ++k) {
if (IGNEUM_DS_SAMPLE_INDEX[k] > mask) continue;
++nSample;
uint32_t v = 0;
CUDA_CHECK(cudaMemcpy(&v, dDs + IGNEUM_DS_SAMPLE_INDEX[k], sizeof(v), cudaMemcpyDeviceToHost));
if (v != IGNEUM_DS_SAMPLE_VALUE[k]) {
if (badSample == 0) std::printf(" dataset[%u] = 0x%08x, Mac 0x%08x\n", IGNEUM_DS_SAMPLE_INDEX[k], v, IGNEUM_DS_SAMPLE_VALUE[k]);
++badSample;
}
}
#endif
r.dsPass = (badHead == 0 && lastOk && badRnd == 0 && badSample == 0);
std::printf("dataset self-test: %s (head 16 vs Mac %s, element [MASK] vs Mac %s, 64 random points vs host %s %s, %d Mac samples %s)\n",
r.dsPass ? "PASS" : "FAIL", badHead == 0 ? "PASS" : "FAIL", lastText,
IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", badRnd == 0 ? "PASS" : "FAIL",
nSample, nSample == 0 ? "none in pack" : (badSample == 0 ? "PASS" : "FAIL"));
}
// Vectors, standalone: one 32-thread block per base nonce, exactly like the Mac cross-check.
uint64_t got[32];
if (atPackSize) {
r.vecChecked = true;
r.vecPass = true;
for (int w = 0; w < IGNEUM_VEC_WARPS; ++w) {
CUDA_CHECK(igneum_launch_hash(dDs, dOut, IGNEUM_VEC_BASE[w], mask, 32u, 1u));
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(got, dOut, sizeof(got), cudaMemcpyDeviceToHost));
bool ok = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], "standalone, 1 warp/block");
r.vecPass = r.vecPass && ok;
}
} else {
std::printf("vectors: skipped (the pack's vectors are for %d MiB)\n", packMib());
}
// Warm-up batch at base nonce 0. With the default 2^24 nonces it contains all three vector warps,
// so the bench configuration itself (blockDim = 32 x block-warps) is also checked bit for bit.
double w0 = wallMs();
CUDA_CHECK(igneum_launch_hash(dDs, dOut, 0u, mask, nonces, (uint32_t)o.blockWarps));
CUDA_CHECK(cudaDeviceSynchronize());
double w1 = wallMs();
std::printf("warm-up batch: %u hashes in %.2f ms wall\n", nonces, w1 - w0);
if (atPackSize) {
char how[64];
std::snprintf(how, sizeof(how), "in batch, %d warp(s)/block", o.blockWarps);
for (int w = 0; w < IGNEUM_VEC_WARPS; ++w) {
if ((uint64_t)IGNEUM_VEC_BASE[w] + 32ull > (uint64_t)nonces) {
std::printf("verify warp base %u in batch: skipped (batch has %u nonces)\n", IGNEUM_VEC_BASE[w], nonces);
continue;
}
CUDA_CHECK(cudaMemcpy(got, dOut + IGNEUM_VEC_BASE[w], sizeof(got), cudaMemcpyDeviceToHost));
bool ok = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], how);
r.vecPass = r.vecPass && ok;
}
}
// Timed batches. Base nonces (b * nonces) mod 2^32, as in the Metal run.
CUDA_CHECK(cudaEventRecord(e0));
double t0 = wallMs();
for (int b = 1; b <= o.batches; ++b) {
uint32_t base = (uint32_t)((uint64_t)b * (uint64_t)nonces);
CUDA_CHECK(igneum_launch_hash(dDs, dOut, base, mask, nonces, (uint32_t)o.blockWarps));
}
CUDA_CHECK(cudaEventRecord(e1));
CUDA_CHECK(cudaEventSynchronize(e1));
double t1 = wallMs();
float gpuMsF = 0.f;
CUDA_CHECK(cudaEventElapsedTime(&gpuMsF, e0, e1));
double total = (double)nonces * (double)o.batches;
r.gpuMs = gpuMsF;
r.wallMsTimed = t1 - t0;
r.hashesPerSec = total / (r.gpuMs / 1000.0);
r.gbps = r.hashesPerSec * (double)IGNEUM_LOADS_PER_HASH * 4.0 / 1e9;
std::printf("timed: %d batches x %u hashes = %.0f hashes\n", o.batches, nonces, total);
std::printf(" GPU %.2f ms -> %.3f Mhash/s (%.0f hashes/s), %.2f GB/s useful (loads x 4 B)\n",
r.gpuMs, r.hashesPerSec / 1e6, r.hashesPerSec, r.gbps);
std::printf(" wall %.2f ms -> %.3f Mhash/s\n", r.wallMsTimed, total / (r.wallMsTimed / 1000.0) / 1e6);
CUDA_CHECK(cudaEventDestroy(e0));
CUDA_CHECK(cudaEventDestroy(e1));
CUDA_CHECK(cudaFree(dDs));
return r;
}
// ---------------------------------------------------------------------------------------------
// Serve mode: GPU worker for igneum-miner --worker (3 October 2026)
//
// Protocol, one line each. Only "found", "done" and "error" are parsed by the miner; every other line is logged.
// stdin: job <job_id> <header_prehash_hex 64> <target_hex 16> <nonce_start u64> <nonce_count u64> <epoch_seed_hex 64> <day_seed_hex>
// quit
// stdout: ready cuda <device>
// found <job_id> <nonce u64> <hash_hex 16> every nonce whose 64-bit hash is <= target
// done <job_id> <hashes> <ms> end of the job (wall ms)
// error <job_id> <text>
// prepare <epoch_seed_hex 64> <day_seed_hex> <pack_dir> build <pack_dir>/kernel.cu and kernel_bound.cu (the pack the
// miner wrote for those seeds) to cubins with nvcc in the background,
// then load them and build that pair's cache and dataset
// prepared <epoch_seed_hex> <day_seed_hex> <ms> ... the pair is resident; a job on it switches instantly
// prepare-failed <epoch_seed_hex> <day_seed_hex> <text>
// The program is compiled ahead of time from the pack (no NVRTC), so at start this worker serves exactly one epoch
// seed and one day seed: the pack's. The next pair arrives through `prepare`: the miner writes the pack for the
// prepared seeds (igneum-miner --prepare-packs <dir>) and names its directory; a background thread runs nvcc on that
// pack's kernel.cu (cache fill and build kernels, memhard.h for its day) and kernel_bound.cu (the bound hash kernel)
// to two cubins for this device's architecture, and the main loop loads them through the driver API
// (cudaGetDriverEntryPoint, so nothing new is linked), fills the cache and builds the dataset while jobs on the current
// pair keep running (at most two pairs resident; the old one is released after the first job on the new one). nvcc
// must be on PATH with a host compiler, as build.bat needs it; the ready line says "prepare 1" only when `nvcc --version`
// answers. A prepared pair's cache is not cross-checked against the host fill (memhard.h is compiled in for the
// original day); the miner's CPU re-check of every found nonce covers it. Without prepare, re-export the pack with
// `igneum-miner export-pack <node> <dir>` and rebuild. The init words of a dispatch are
// seed_words_from_bytes("igneum-block/" || prehash || nonce_hi_le32), passed by value to igneum_hash_bound
// (kernel_bound.cu); the lane nonce is baseNonce + gid as in the bench kernel.
#ifdef IGNEUM_BOUND
struct IgneumInitWords { uint32_t w[8]; };
cudaError_t igneum_launch_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask,
IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);
cudaError_t igneum_hash_bound_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps);
#endif
static void emitLine(const char* s) { std::fputs(s, stdout); std::fputc('\n', stdout); std::fflush(stdout); }
// seed_words_from_bytes of igneum-pow/src/seed.rs: FNV-1a 64 with four salts, each finalised.
static void seedWordsFromBytes(const uint8_t* b, size_t n, uint32_t out[8]) {
for (uint64_t salt = 0; salt < 4; ++salt) {
uint64_t h = 0xcbf29ce484222325ull ^ (salt * 0x9E3779B97F4A7C15ull);
for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
h ^= h >> 33; h *= 0xff51afd7ed558ccdull; h ^= h >> 33;
out[2 * salt] = (uint32_t)h;
out[2 * salt + 1] = (uint32_t)(h >> 32);
}
}
static bool unhexStr(const std::string& s, std::vector<uint8_t>& out) {
if (s.size() % 2) return false;
out.clear();
for (size_t i = 0; i < s.size(); i += 2) {
unsigned v = 0;
if (std::sscanf(s.substr(i, 2).c_str(), "%2x", &v) != 1) return false;
out.push_back((uint8_t)v);
}
return true;
}
#if defined(IGNEUM_BOUND) && IGNEUM_DATASET_MODE == 1 && defined(__has_include)
#if __has_include(<cuda.h>)
#define IGNEUM_CUDA_PREPARE 1
#include <cuda.h>
#include <thread>
#include <atomic>
#include <fstream>
#endif
#endif
#ifdef IGNEUM_CUDA_PREPARE
// The few driver API entry points the hot swap needs, fetched through the runtime so the link line is unchanged.
struct DriverApi {
CUresult (*moduleLoad)(CUmodule*, const char*) = nullptr;
CUresult (*moduleUnload)(CUmodule) = nullptr;
CUresult (*moduleGetFunction)(CUfunction*, CUmodule, const char*) = nullptr;
CUresult (*moduleGetFunctionCount)(unsigned int*, CUmodule) = nullptr; // CUDA 12.4 and newer
CUresult (*moduleEnumerateFunctions)(CUfunction*, unsigned int, CUmodule) = nullptr;
CUresult (*funcGetName)(const char**, CUfunction) = nullptr; // CUDA 12.3 and newer
CUresult (*launchKernel)(CUfunction, unsigned, unsigned, unsigned, unsigned, unsigned, unsigned, unsigned, CUstream, void**, void**) = nullptr;
CUresult (*getErrorString)(CUresult, const char**) = nullptr;
bool ok = false;
std::string why;
template <typename F> bool get(const char* name, F& fn, bool required) {
void* p = nullptr;
cudaError_t e = cudaGetDriverEntryPoint(name, &p, cudaEnableDefault);
if (e != cudaSuccess || !p) { if (required) { why = std::string("no driver entry point ") + name; } return false; }
fn = reinterpret_cast<F>(p);
return true;
}
void load() {
ok = get("cuModuleLoad", moduleLoad, true) && get("cuModuleUnload", moduleUnload, true) && get("cuModuleGetFunction", moduleGetFunction, true) &&
get("cuLaunchKernel", launchKernel, true) && get("cuGetErrorString", getErrorString, true);
get("cuModuleGetFunctionCount", moduleGetFunctionCount, false);
get("cuModuleEnumerateFunctions", moduleEnumerateFunctions, false);
get("cuFuncGetName", funcGetName, false);
}
std::string err(CUresult r) { const char* s = nullptr; if (getErrorString) getErrorString(r, &s); return s ? s : "CUDA driver error"; }
// A kernel by its plain name: the Itanium mangling nvcc gives device code first, then the enumeration (12.4+).
bool find(CUmodule m, const char* plain, const char* mangled, CUfunction* out) {
if (moduleGetFunction(out, m, mangled) == CUDA_SUCCESS) return true;
if (moduleGetFunction(out, m, plain) == CUDA_SUCCESS) return true;
if (!moduleGetFunctionCount || !moduleEnumerateFunctions || !funcGetName) return false;
unsigned int n = 0;
if (moduleGetFunctionCount(&n, m) != CUDA_SUCCESS || n == 0) return false;
std::vector<CUfunction> fns(n);
if (moduleEnumerateFunctions(fns.data(), n, m) != CUDA_SUCCESS) return false;
for (CUfunction f : fns) {
const char* name = nullptr;
if (funcGetName(&name, f) == CUDA_SUCCESS && name && std::strstr(name, plain)) { *out = f; return true; }
}
return false;
}
};
// One resident pair: the compiled-in pack (runtime launchers, gCache) or a prepared pack (two cubins, driver launches).
struct CudaPair {
std::string epochHex, dayHex;
uint32_t sw[8] = {0}, kw[8] = {0};
bool builtIn = false;
CUmodule modKernel = nullptr, modBound = nullptr;
CUfunction fCacheFill = nullptr, fBuild = nullptr, fHashBound = nullptr;
uint32_t* cache = nullptr;
uint32_t* ds = nullptr;
double nvccMs = 0, cacheMs = 0, dsMs = 0;
};
static void releasePair(DriverApi& drv, CudaPair* p) {
if (!p) return;
if (p->ds) cudaFree(p->ds);
if (p->cache) cudaFree(p->cache);
if (p->modBound) drv.moduleUnload(p->modBound);
if (p->modKernel) drv.moduleUnload(p->modKernel);
delete p;
}
// The nvcc step of a prepare, on its own thread. Only the compiler runs here; every CUDA call stays on the main thread.
struct PrepareTask {
std::string epochHex, dayHex, packDir, arch, error;
std::atomic<bool> done{false};
bool ok = false;
double t0 = 0, nvccMs = 0;
std::thread thread;
};
static bool nvccAvailable() {
#ifdef _WIN32
return std::system("nvcc --version >NUL 2>&1") == 0;
#else
return std::system("nvcc --version >/dev/null 2>&1") == 0;
#endif
}
static void prepareCompile(PrepareTask* t) {
double c0 = wallMs();
const char* files[2] = { "kernel", "kernel_bound" };
for (const char* f : files) {
std::string cmd = "nvcc -cubin -O3 -std=c++17 -arch=" + t->arch + " -allow-unsupported-compiler -I \"" + t->packDir + "\" -o \"" + t->packDir + "/" + f +
".cubin\" \"" + t->packDir + "/" + f + ".cu\" > \"" + t->packDir + "/nvcc-" + f + ".log\" 2>&1";
int rc = std::system(cmd.c_str());
if (rc != 0) {
std::string log;
std::ifstream in(t->packDir + "/nvcc-" + f + ".log");
std::string line;
while (std::getline(in, line) && log.size() < 300) { log += line; log += " | "; }
t->error = std::string("nvcc failed on ") + f + ".cu (exit " + std::to_string(rc) + "): " + log;
t->done = true;
return;
}
}
t->nvccMs = wallMs() - c0;
t->ok = true;
t->done = true;
}
// Loads the two cubins, builds the cache and dataset (main thread). Returns the pair or null with `error` set.
static CudaPair* prepareLoad(DriverApi& drv, const PrepareTask& t, uint32_t words, std::string& error) {
CudaPair* p = new CudaPair();
p->epochHex = t.epochHex; p->dayHex = t.dayHex; p->nvccMs = t.nvccMs;
{
std::vector<uint8_t> eb, db;
if (unhexStr(t.epochHex, eb) && eb.size() == 32) seedWordsFromBytes(eb.data(), 32, p->sw);
if (unhexStr(t.dayHex, db)) seedWordsFromBytes(db.data(), db.size(), p->kw);
}
CUresult r = drv.moduleLoad(&p->modKernel, (t.packDir + "/kernel.cubin").c_str());
if (r != CUDA_SUCCESS) { error = "cuModuleLoad kernel.cubin: " + drv.err(r); releasePair(drv, p); return nullptr; }
r = drv.moduleLoad(&p->modBound, (t.packDir + "/kernel_bound.cubin").c_str());
if (r != CUDA_SUCCESS) { error = "cuModuleLoad kernel_bound.cubin: " + drv.err(r); releasePair(drv, p); return nullptr; }
if (!drv.find(p->modKernel, "igneum_cache_fill", "_Z17igneum_cache_fillPjj", &p->fCacheFill)) { error = "igneum_cache_fill not found in kernel.cubin"; releasePair(drv, p); return nullptr; }
if (!drv.find(p->modKernel, "igneum_build", "_Z12igneum_buildPjPKjj", &p->fBuild)) { error = "igneum_build not found in kernel.cubin"; releasePair(drv, p); return nullptr; }
if (!drv.find(p->modBound, "igneum_hash_bound", "_Z17igneum_hash_boundPKjPyjj15IgneumInitWords", &p->fHashBound)) { error = "igneum_hash_bound not found in kernel_bound.cubin"; releasePair(drv, p); return nullptr; }
// Cache (the same segment count as the compiled-in pack: the dataset schedule is a network constant)
double c0 = wallMs();
size_t cacheBytes = (size_t)CACHE_WORDS_HOST * 4u;
if (cudaMalloc((void**)&p->cache, cacheBytes) != cudaSuccess) { error = "cudaMalloc cache"; p->cache = nullptr; releasePair(drv, p); return nullptr; }
{
uint32_t nSeg = IGNEUM_CACHE_SEGMENTS, block = 256u, grid = (nSeg + block - 1u) / block;
void* args[2] = { &p->cache, &nSeg };
r = drv.launchKernel(p->fCacheFill, grid, 1, 1, block, 1, 1, 0, nullptr, args, nullptr);
if (r != CUDA_SUCCESS || cudaDeviceSynchronize() != cudaSuccess) { error = "cache fill launch: " + drv.err(r); releasePair(drv, p); return nullptr; }
}
p->cacheMs = wallMs() - c0;
// Dataset
c0 = wallMs();
if (cudaMalloc((void**)&p->ds, (size_t)words * 4u) != cudaSuccess) { error = "cudaMalloc dataset"; p->ds = nullptr; releasePair(drv, p); return nullptr; }
{
uint32_t nItems = words / 16u, block = 256u, grid = (nItems + block - 1u) / block;
void* args[3] = { &p->ds, &p->cache, &nItems };
r = drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, nullptr, args, nullptr);
if (r != CUDA_SUCCESS || cudaDeviceSynchronize() != cudaSuccess) { error = "dataset build launch: " + drv.err(r); releasePair(drv, p); return nullptr; }
}
p->dsMs = wallMs() - c0;
return p;
}
#endif
static int runServe(const Options& o) {
#if !defined(IGNEUM_BOUND) || IGNEUM_DATASET_MODE != 1
(void)o;
emitLine("error 0 this binary was built without the pack's kernel_bound.cu (IGNEUM_BOUND) or from a closed-form pack; rebuild with build.bat from a memory-hard pack");
return 2;
#else
cudaDeviceProp prop;
std::memset(&prop, 0, sizeof(prop));
CUDA_CHECK(cudaGetDeviceProperties(&prop, o.device));
std::string devName = prop.name;
for (char& c : devName) if (c == ' ') c = '_';
// The pack's seeds: the program seed words and the day key words
static const uint32_t KEYW[8] = IGNEUM_KEY_INIT;
// Cache and dataset, once
if (!setupCache()) { emitLine("error 0 cache check failed (GPU cache differs from the host cache or the pack's FNV)"); return 1; }
const uint32_t words = 1u << IGNEUM_DATASET_LOG2;
const uint32_t mask = words - 1u;
uint32_t* dDs = nullptr;
CUDA_CHECK(cudaMalloc((void**)&dDs, (size_t)words * 4u));
CUDA_CHECK(igneum_launch_build(dDs, gCache, words / 16u));
CUDA_CHECK(cudaDeviceSynchronize());
const uint32_t batch = 1u << o.batchLog2;
uint64_t* dOut = nullptr;
CUDA_CHECK(cudaMalloc((void**)&dOut, (size_t)batch * sizeof(uint64_t)));
std::vector<uint64_t> hOut(batch);
int regs = 0, bps = 0;
igneum_hash_bound_info(&regs, &bps, (uint32_t)o.blockWarps);
// Prepare support: the driver entry points and nvcc on PATH
int prepareOk = 0;
#ifdef IGNEUM_CUDA_PREPARE
DriverApi drv;
drv.load();
std::string arch = "sm_" + std::to_string(prop.major) + std::to_string(prop.minor);
bool haveNvcc = !o.noPrepare && nvccAvailable();
prepareOk = (!o.noPrepare && drv.ok && haveNvcc) ? 1 : 0;
CudaPair* cur = new CudaPair();
cur->builtIn = true; cur->cache = gCache; cur->ds = dDs;
std::memcpy(cur->sw, SEEDW, 32); std::memcpy(cur->kw, KEYW, 32);
CudaPair* prepared = nullptr;
CudaPair* old = nullptr;
PrepareTask* task = nullptr;
#endif
std::printf("ready cuda %s pack %s dataset-log2 %d batch %u regs %d prepare %d\n", devName.c_str(), IGNEUM_SEED_STRING, IGNEUM_DATASET_LOG2, batch, regs, prepareOk);
#ifdef IGNEUM_CUDA_PREPARE
if (!o.noPrepare && !prepareOk) std::printf("info prepare unavailable: %s\n", !drv.ok ? drv.why.c_str() : "nvcc is not on PATH (open the build prompt, or install the CUDA Toolkit)");
#else
std::printf("info prepare unavailable: this binary was built without cuda.h (CPU emulation or an old toolkit)\n");
#endif
std::fflush(stdout);
std::string line;
while (std::getline(std::cin, line)) {
if (line == "quit") break;
std::vector<std::string> f;
{ size_t i = 0; while (i < line.size()) { while (i < line.size() && line[i] == ' ') ++i; size_t j = i; while (j < line.size() && line[j] != ' ') ++j; if (j > i) f.push_back(line.substr(i, j - i)); i = j; } }
if (f.empty()) continue;
#ifdef IGNEUM_CUDA_PREPARE
// A finished nvcc step is loaded here, between lines, on this thread
if (task && task->done) {
task->thread.join();
if (task->ok) {
std::string error;
CudaPair* p = prepareLoad(drv, *task, words, error);
if (p) {
if (prepared) releasePair(drv, prepared);
prepared = p;
std::printf("prepared %s %s %.1f nvcc %.1f cache %.1f dataset %.1f resident 2 programs 2 datasets\n", p->epochHex.c_str(), p->dayHex.c_str(), wallMs() - task->t0, p->nvccMs, p->cacheMs, p->dsMs);
} else {
std::printf("prepare-failed %s %s %s\n", task->epochHex.c_str(), task->dayHex.c_str(), error.c_str());
}
} else {
std::printf("prepare-failed %s %s %s\n", task->epochHex.c_str(), task->dayHex.c_str(), task->error.c_str());
}
std::fflush(stdout);
delete task; task = nullptr;
}
if (f[0] == "prepare") {
if (!prepareOk) { std::printf("info ignored (no prepare support): %s\n", line.c_str()); std::fflush(stdout); continue; }
if (f.size() < 4) { std::printf("prepare-failed %s %s this ahead-of-time worker needs a pack directory as the third field (igneum-miner --prepare-packs <dir>)\n", f.size() > 1 ? f[1].c_str() : "0", f.size() > 2 ? f[2].c_str() : "0"); std::fflush(stdout); continue; }
if (f[1].size() != 64) { std::printf("prepare-failed %s %s bad field (epoch_seed 64 hex, day_seed hex)\n", f[1].c_str(), f[2].c_str()); std::fflush(stdout); continue; }
if (task) { std::printf("prepare-failed %s %s a prepare is still running\n", f[1].c_str(), f[2].c_str()); std::fflush(stdout); continue; }
if (prepared && prepared->epochHex == f[1] && prepared->dayHex == f[2]) { std::printf("prepared %s %s 0 (already resident)\n", f[1].c_str(), f[2].c_str()); std::fflush(stdout); continue; }
task = new PrepareTask();
task->epochHex = f[1]; task->dayHex = f[2]; task->packDir = f[3]; task->arch = arch; task->t0 = wallMs();
task->thread = std::thread(prepareCompile, task);
std::printf("info prepare started for epoch %.16s day %s from %s (nvcc -arch=%s in the background)\n", f[1].c_str(), f[2].c_str(), f[3].c_str(), arch.c_str()); std::fflush(stdout);
continue;
}
#else
if (f[0] == "prepare") { std::printf("info ignored (no prepare support): %s\n", line.c_str()); std::fflush(stdout); continue; }
#endif
if (f[0] != "job") { std::printf("info ignored: %s\n", line.c_str()); std::fflush(stdout); continue; }
std::string jobId = f.size() > 1 ? f[1] : "0";
if (f.size() < 8) { std::printf("error %s malformed job line (need 7 fields after job)\n", jobId.c_str()); std::fflush(stdout); continue; }
std::vector<uint8_t> prehash, epochSeed, daySeed;
uint64_t target = 0, nonceStart = 0, nonceCount = 0;
if (!unhexStr(f[2], prehash) || prehash.size() != 32 || std::sscanf(f[3].c_str(), "%llx", (unsigned long long*)&target) != 1 ||
std::sscanf(f[4].c_str(), "%llu", (unsigned long long*)&nonceStart) != 1 || std::sscanf(f[5].c_str(), "%llu", (unsigned long long*)&nonceCount) != 1 ||
!unhexStr(f[6], epochSeed) || epochSeed.size() != 32 || !unhexStr(f[7], daySeed)) {
std::printf("error %s bad field (prehash 64 hex, target 16 hex, nonce_start u64, nonce_count u64, epoch_seed 64 hex, day_seed hex)\n", jobId.c_str()); std::fflush(stdout); continue;
}
if (nonceCount == 0 || nonceCount % 32 != 0 || (nonceStart & 31) != 0) { std::printf("error %s nonce_start must be 32-aligned and nonce_count a non-zero multiple of 32\n", jobId.c_str()); std::fflush(stdout); continue; }
uint32_t sw[8], kw[8];
seedWordsFromBytes(epochSeed.data(), epochSeed.size(), sw);
seedWordsFromBytes(daySeed.data(), daySeed.size(), kw);
double t0 = wallMs();
#ifdef IGNEUM_CUDA_PREPARE
bool switched = false;
if (std::memcmp(sw, cur->sw, 32) != 0 || std::memcmp(kw, cur->kw, 32) != 0) {
if (prepared && std::memcmp(sw, prepared->sw, 32) == 0 && std::memcmp(kw, prepared->kw, 32) == 0) {
if (old) releasePair(drv, old);
old = cur; cur = prepared; prepared = nullptr; switched = true;
std::printf("info switched to the prepared pair epoch %.16s day %s in %.2f ms\n", cur->epochHex.c_str(), cur->dayHex.c_str(), wallMs() - t0); std::fflush(stdout);
} else if (std::memcmp(sw, cur->sw, 32) != 0) {
std::printf("error %s epoch seed mismatch: this worker holds %s%s (seed words %08x %08x ...)%s, the job's epoch seed %s gives %08x %08x ...; send prepare with a pack directory, or run igneum-miner export-pack and rebuild\n",
jobId.c_str(), cur->builtIn ? "pack \"" IGNEUM_SEED_STRING "\"" : "prepared epoch ", cur->builtIn ? "" : cur->epochHex.c_str(), cur->sw[0], cur->sw[1], prepared ? " plus one prepared pair" : "", f[6].substr(0, 16).c_str(), sw[0], sw[1]);
std::fflush(stdout); continue;
} else {
std::printf("error %s day seed mismatch: this worker's cache is for key %08x %08x ..., the job's day seed %s gives %08x %08x ...; send prepare with a pack directory, or run igneum-miner export-pack and rebuild\n",
jobId.c_str(), cur->kw[0], cur->kw[1], f[7].c_str(), kw[0], kw[1]);
std::fflush(stdout); continue;
}
}
const uint32_t* jobDs = cur->ds;
#else
if (std::memcmp(sw, SEEDW, 32) != 0) {
std::printf("error %s epoch seed mismatch: this worker was built for pack \"%s\" (seed words %08x %08x ...), the job's epoch seed %s gives %08x %08x ...; run igneum-miner export-pack and rebuild\n",
jobId.c_str(), IGNEUM_SEED_STRING, SEEDW[0], SEEDW[1], f[6].substr(0, 16).c_str(), sw[0], sw[1]);
std::fflush(stdout); continue;
}
if (std::memcmp(kw, KEYW, 32) != 0) {
std::printf("error %s day seed mismatch: this worker's cache is for key %08x %08x ..., the job's day seed %s gives %08x %08x ...; run igneum-miner export-pack and rebuild\n",
jobId.c_str(), KEYW[0], KEYW[1], f[7].c_str(), kw[0], kw[1]);
std::fflush(stdout); continue;
}
const uint32_t* jobDs = dDs;
#endif
uint64_t remaining = nonceCount, hashes = 0;
uint32_t hi = (uint32_t)(nonceStart >> 32), lo = (uint32_t)nonceStart;
bool failed = false;
while (remaining > 0) {
uint64_t room = (uint64_t)(0xffffffffu - lo) + 1ull;
uint64_t chunk64 = remaining < batch ? remaining : batch;
if (chunk64 > room) chunk64 = room;
uint32_t chunk = (uint32_t)chunk64;
IgneumInitWords iw;
{
uint8_t b[49];
std::memcpy(b, "igneum-block/", 13);
std::memcpy(b + 13, prehash.data(), 32);
b[45] = (uint8_t)hi; b[46] = (uint8_t)(hi >> 8); b[47] = (uint8_t)(hi >> 16); b[48] = (uint8_t)(hi >> 24);
seedWordsFromBytes(b, 49, iw.w);
}
cudaError_t e = cudaSuccess;
#ifdef IGNEUM_CUDA_PREPARE
if (!cur->builtIn) {
// A prepared pair: the same launch shape as igneum_launch_hash_bound, through the driver API
uint32_t block = 32u * (uint32_t)o.blockWarps;
uint32_t baseNonce = lo, maskArg = mask;
void* args[5] = { (void*)&jobDs, (void*)&dOut, &baseNonce, &maskArg, &iw };
CUresult r = drv.launchKernel(cur->fHashBound, chunk / block, 1, 1, block, 1, 1, 0, nullptr, args, nullptr);
if (r != CUDA_SUCCESS) { std::printf("error %s dispatch failed: %s\n", jobId.c_str(), drv.err(r).c_str()); std::fflush(stdout); failed = true; break; }
} else
#endif
e = igneum_launch_hash_bound(jobDs, dOut, lo, mask, iw, chunk, (uint32_t)o.blockWarps);
if (e == cudaSuccess) e = cudaDeviceSynchronize();
if (e == cudaSuccess) e = cudaMemcpy(hOut.data(), dOut, (size_t)chunk * sizeof(uint64_t), cudaMemcpyDeviceToHost);
if (e != cudaSuccess) { std::printf("error %s dispatch failed: %s\n", jobId.c_str(), cudaGetErrorString(e)); std::fflush(stdout); failed = true; break; }
for (uint32_t i = 0; i < chunk; ++i) if (hOut[i] <= target) {
uint64_t nonce = ((uint64_t)hi << 32) | (uint64_t)(uint32_t)(lo + i);
std::printf("found %s %llu %016llx\n", jobId.c_str(), (unsigned long long)nonce, (unsigned long long)hOut[i]);
}
std::fflush(stdout);
hashes += chunk;
remaining -= chunk;
if (chunk64 == room) { hi += 1u; lo = 0u; } else lo += chunk;
}
if (failed) continue;
std::printf("done %s %llu %.2f\n", jobId.c_str(), (unsigned long long)hashes, wallMs() - t0);
#ifdef IGNEUM_CUDA_PREPARE
if (switched && old) { releasePair(drv, old); old = nullptr; std::printf("info dropped the previous pair (its program, cache and dataset)\n"); }
#endif
std::fflush(stdout);
}
cudaFree(dOut);
#ifdef IGNEUM_CUDA_PREPARE
if (task) { task->thread.join(); delete task; }
if (old) releasePair(drv, old);
if (prepared) releasePair(drv, prepared);
releasePair(drv, cur); // frees gCache and dDs when the built-in pair is still current
#else
cudaFree(dDs);
cudaFree(gCache);
#endif
return 0;
#endif
}
// ---------------------------------------------------------------------------------------------
// Main
int main(int argc, char** argv) {
Options o = parseArgs(argc, argv);
std::printf("igneum-bench-cuda pack \"%s\" (test harness: no pool, no network, no wallet)\n", IGNEUM_SEED_STRING);
int count = 0;
CUDA_CHECK(cudaGetDeviceCount(&count));
if (count == 0) { std::printf("FAIL: no CUDA device\n"); return 2; }
if (o.device < 0 || o.device >= count) { std::printf("FAIL: device %d out of range (%d devices)\n", o.device, count); return 2; }
CUDA_CHECK(cudaSetDevice(o.device));
// Blocking sync, set before the context exists: the host thread sleeps in cudaDeviceSynchronize instead of
// spinning (one full core per worker process at the default spin schedule, measured on the RTX 5090 with
// eight workers, 3 Oct 2026). The microseconds of wake-up latency are nothing against a 100 ms dispatch.
#ifdef cudaDeviceScheduleBlockingSync
CUDA_CHECK(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync));
#endif
if (o.serve) {
if (o.batchLog2 == 24) { Options s2 = o; s2.batchLog2 = 22; return runServe(s2); } // 2^22 nonces per dispatch by default
return runServe(o);
}
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));
int warp = 0, clk = 0, memclk = 0, bus = 0, l2 = 0, thrSM = 0;
cudaDeviceGetAttribute(&warp, cudaDevAttrWarpSize, o.device);
cudaDeviceGetAttribute(&clk, cudaDevAttrClockRate, o.device);
cudaDeviceGetAttribute(&memclk, cudaDevAttrMemoryClockRate, o.device);
cudaDeviceGetAttribute(&bus, cudaDevAttrGlobalMemoryBusWidth, o.device);
cudaDeviceGetAttribute(&l2, cudaDevAttrL2CacheSize, o.device);
cudaDeviceGetAttribute(&thrSM, cudaDevAttrMaxThreadsPerMultiProcessor, o.device);
cudaGetLastError(); // attribute queries are informational; clear any error they left
std::printf("GPU: %s (%d SMs, compute capability %d.%d, %.0f MiB global memory)\n",
prop.name, prop.multiProcessorCount, prop.major, prop.minor, (double)prop.totalGlobalMem / 1048576.0);
std::printf(" SM clock %d MHz, memory clock %d MHz, bus %d bits, L2 %d MiB, max %d threads/SM, warp size %d\n",
clk / 1000, memclk / 1000, bus, l2 / 1048576, thrSM, warp);
std::printf("CUDA: driver %d.%d, runtime %d.%d\n", drv / 1000, (drv % 100) / 10, rt / 1000, (rt % 100) / 10);
if (warp != 32) {
std::printf("WARNING: warp size is %d, not 32. The 32-lane shuffle model does not hold on this device.\n", warp);
}
int regs = 0, blocksPerSM = 0;
CUDA_CHECK(igneum_hash_info(&regs, &blocksPerSM, (uint32_t)o.blockWarps));
std::printf("kernel: %d registers/thread, %d resident blocks/SM at %d warp(s)/block = %d resident warps/SM\n",
regs, blocksPerSM, o.blockWarps, blocksPerSM * o.blockWarps);
std::printf("program: %d instructions x %d iterations, loads/hash %d, op mix %s\n",
IGNEUM_INSTR_COUNT, IGNEUM_ITERATIONS, IGNEUM_LOADS_PER_HASH, IGNEUM_OP_MIX);
std::printf("seed words: %08x %08x %08x %08x %08x %08x %08x %08x\n",
SEEDW[0], SEEDW[1], SEEDW[2], SEEDW[3], SEEDW[4], SEEDW[5], SEEDW[6], SEEDW[7]);
std::printf("day \"%s\" (d0 0x%08x, d1 0x%08x), pack dataset 2^%d words = %d MiB\n",
IGNEUM_DAY_STRING, IGNEUM_DAY0, IGNEUM_DAY1, IGNEUM_DATASET_LOG2, packMib());
#if IGNEUM_DATASET_MODE == 1
std::printf("dataset construction: memory-hard (256 MiB ChaCha cache, %d dependent cache reads per 64-byte item; proto-metal/MEMHARD.md)\n", IGNEUM_ITEM_ROUNDS);
bool cachePass = setupCache();
#else
std::printf("dataset construction: closed-form ds_elem (the original prototype dataset, not memory-hard)\n");
bool cachePass = true;
#endif
uint32_t nonces = 1u << o.batchLog2;
if (nonces % (32u * (uint32_t)o.blockWarps) != 0u) {
std::printf("FAIL: 2^%d nonces is not a multiple of %d threads per block\n", o.batchLog2, 32 * o.blockWarps);
return 2;
}
uint64_t* dOut = nullptr;
CUDA_CHECK(cudaMalloc((void**)&dOut, (size_t)nonces * sizeof(uint64_t)));
std::vector<int> sizes;
if (o.sweep) { sizes.push_back(4); sizes.push_back(64); sizes.push_back(256); sizes.push_back(512); sizes.push_back(1024); }
else sizes.push_back(o.datasetMib);
std::vector<SizeResult> results;
for (size_t i = 0; i < sizes.size(); ++i) results.push_back(runSize(o, sizes[i], dOut, nonces));
CUDA_CHECK(cudaFree(dOut));
std::printf("\n=== summary (%s, pack %s, batch 2^%d x %d, %d warp(s)/block, GPU-event time) ===\n",
prop.name, IGNEUM_SEED_STRING, o.batchLog2, o.batches, o.blockWarps);
std::printf("| dataset MiB | %s ms (second) | Mhash/s | GB/s useful | random loads/s (G) | loads/hash | dataset self-test | vectors |\n",
IGNEUM_DATASET_MODE == 1 ? "build" : "fill");
std::printf("|---|---|---|---|---|---|---|---|\n");
bool overall = cachePass, anyVec = false;
for (size_t i = 0; i < results.size(); ++i) {
const SizeResult& r = results[i];
overall = overall && r.dsPass && (!r.vecChecked || r.vecPass);
anyVec = anyVec || r.vecChecked;
std::printf("| %d | %.2f | %.3f | %.2f | %.2f | %d | %s | %s |\n",
r.mib, r.fillSecondMs, r.hashesPerSec / 1e6, r.gbps,
r.hashesPerSec * (double)IGNEUM_LOADS_PER_HASH / 1e9, IGNEUM_LOADS_PER_HASH,
r.dsPass ? "PASS" : "FAIL",
r.vecChecked ? (r.vecPass ? "PASS (3 warps, standalone and in batch)" : "FAIL") : "skipped (not pack size)");
}
#if IGNEUM_DATASET_MODE == 1
std::printf("cache: GPU fill %.2f ms (second), host fill %.1f ms one thread, cache check %s\n", gCacheFillSecondMs, gCacheHostMs, cachePass ? "PASS" : "FAIL");
CUDA_CHECK(cudaFree(gCache));
#endif
if (!anyVec) std::printf("NOTE: no vectors were checked. Run at %d MiB (the default) to verify against the Mac.\n", packMib());
std::printf("OVERALL: %s\n", overall ? "PASS" : "FAIL");
return overall ? 0 : 1;
}
Reference in a new issue View git blame Copy permalink