From b4207ab2d0ec1d937bf6c5a3bd442f0b8ee29c54 Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Mon, 5 Oct 2026 20:05:57 +0000 Subject: [PATCH] read-width: CUDA worker --bench and --memprobe, scratch arena from the occupancy capacity; PC playbooks (card under test off in the app, restored after); harness arena label Co-Authored-By: Claude Fable 5.1 --- proto-cuda/nvrtc/worker.cpp | 268 ++++++++++++++++++++++++++++- proto-metal/packbench.swift | 2 +- relay/playbooks/readwidth-5090.ps1 | 67 ++++++++ relay/playbooks/readwidth-9070.ps1 | 69 ++++++++ 4 files changed, 400 insertions(+), 6 deletions(-) create mode 100644 relay/playbooks/readwidth-5090.ps1 create mode 100644 relay/playbooks/readwidth-9070.ps1 diff --git a/proto-cuda/nvrtc/worker.cpp b/proto-cuda/nvrtc/worker.cpp index 11597fd2a..94705562d 100644 --- a/proto-cuda/nvrtc/worker.cpp +++ b/proto-cuda/nvrtc/worker.cpp @@ -241,6 +241,10 @@ static bool loadNvrtc(Rtc& r, std::string& err, std::string& libName) { struct Ctx { Drv drv; Rtc rtc; + // read-width experiment, variant 5 (5 October 2026): persistent warps and their scratch; --warps caps the launch + int warps = 0; // 0 = the resident capacity from the occupancy query, rounded down to a power of two + uint32_t salt = 1; // the running per-unit tag salt (+= units per launch) + int batches = 5; // --bench: timed dispatches CUdevice dev = 0; CUcontext ctx = nullptr; std::string name; @@ -409,6 +413,14 @@ struct Pair { std::string variant = "base"; // the bound kernel in service: a variant name (see allVariants) std::string raceLine; // the race's one-line report, emitted by the main thread with "prepared" double raceMs = 0; + // read-width experiment (5 October 2026): the pack's load class and, for variant 5, the persistent-warp scratch + std::string loadClass = "v2"; + uint32_t loadsPerHash = 128, bytesPerHash = 512, scratchOps = 0; + bool persistent = false; + CUdeviceptr scratch = 0; + int warps = 0; // persistent warps launched (the arena holds this many) + int residentWarps = 0; // the occupancy query's capacity: blocks/SM x warps/block x SMs + size_t scratchBytes = 0; }; // The job loop and a race take turns on the card: a variant is timed with no job running (exclusive numbers), and @@ -419,6 +431,7 @@ static void releasePair(Ctx& c, Pair* p) { if (!p) return; if (p->ds) c.drv.memFree(p->ds); if (p->cache) c.drv.memFree(p->cache); + if (p->scratch) c.drv.memFree(p->scratch); if (p->modBound) c.drv.moduleUnload(p->modBound); if (p->modKernel) c.drv.moduleUnload(p->modKernel); delete p; @@ -430,6 +443,17 @@ struct IgneumInitWordsArg { uint32_t w[8]; }; static bool launchHash(Ctx& c, Pair* p, CUdeviceptr out, uint32_t baseNonce, const uint32_t iw[8], uint32_t nonces, uint32_t block, CUstream s, std::string& err) { uint32_t mask = p->words - 1u; IgneumInitWordsArg a; std::memcpy(a.w, iw, 32); + if (p->persistent) { + // Variant 5: N persistent warps over nonces / 32 units; the arena was sized for p->warps warps in buildPair. + uint32_t units = nonces / 32u, warps = (uint32_t)p->warps; + if (warps > units) warps = units; + while (warps > 1u && units % warps != 0u) warps >>= 1; + if (block != 32u) { err = "a variant-5 pack runs one warp per block (--block-warps 1)"; return false; } + uint32_t salt = c.salt; c.salt += units; + void* args[8] = { &p->ds, &out, &baseNonce, &mask, &a, &p->scratch, &units, &salt }; + DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, warps, 1, 1, 32, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound (persistent)"); + return true; + } void* args[5] = { &p->ds, &out, &baseNonce, &mask, &a }; DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, nonces / block, 1, 1, block, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound"); return true; @@ -768,16 +792,39 @@ static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string& c.drv.funcGetAttribute(&p->regs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->fHashBound); c.drv.occupancy(&p->blocksPerSM, p->fHashBound, 32 * c.blockWarps, 0); } + p->loadClass = pk.loadClass; p->loadsPerHash = pk.loadsPerHash; p->bytesPerHash = pk.bytesPerHash; p->scratchOps = pk.scratchOps; + p->persistent = pk.persistent != 0; + p->residentWarps = p->blocksPerSM * c.blockWarps * c.sms; + size_t scratchBytes = 0; + if (p->persistent) { + // Variant 5: one arena per launched warp. The launch is the resident capacity (the occupancy query), rounded + // down to a power of two so it divides every batch, or --warps; the allocation cannot change the occupancy + // (registers and shared memory decide it), and the number is re-queried after the allocation below to show it. + if (c.blockWarps != 1) { err = "a variant-5 pack runs one warp per block: use --block-warps 1"; releasePair(c, p); return nullptr; } + int w = c.warps > 0 ? c.warps : p->residentWarps; + int pw = 1; while (pw * 2 <= w) pw *= 2; + p->warps = pw; + scratchBytes = (size_t)p->warps * 32u * (size_t)pk.scratchWordsPerLane * 4u; + } // Cache double t0 = wallMs(); size_t cacheBytes = (size_t)p->cacheWords * 4u, dsBytes = (size_t)p->words * 4u; { size_t freeB = 0, totalB = 0; - if (c.drv.memGetInfo(&freeB, &totalB) == CUDA_SUCCESS && freeB < cacheBytes + dsBytes + (64u << 20)) { - err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20)); + if (c.drv.memGetInfo(&freeB, &totalB) == CUDA_SUCCESS && freeB < cacheBytes + dsBytes + scratchBytes + (64u << 20)) { + err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu + scratch %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes + scratchBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20), (unsigned long long)(scratchBytes >> 20)); releasePair(c, p); return nullptr; } } + if (p->persistent) { + CUresult r = c.drv.memAlloc(&p->scratch, scratchBytes); + if (r != CUDA_SUCCESS) { err = "cuMemAlloc scratch: " + c.err(r); p->scratch = 0; releasePair(c, p); return nullptr; } + p->scratchBytes = scratchBytes; + int after = 0; + c.drv.occupancy(&after, p->fHashBound, 32 * c.blockWarps, 0); + info(fmt("variant 5: %d persistent warps (resident capacity %d = %d blocks/SM x %d warps/block x %d SMs; occupancy query after the allocation %d blocks/SM), scratch %llu MiB (%u KiB per warp)", + p->warps, p->residentWarps, p->blocksPerSM, c.blockWarps, c.sms, after, (unsigned long long)(scratchBytes >> 20), pk.scratchWordsPerLane * 4u * 32u / 1024u)); + } { CUresult r = c.drv.memAlloc(&p->cache, cacheBytes); if (r != CUDA_SUCCESS) { err = "cuMemAlloc cache: " + c.err(r); p->cache = 0; releasePair(c, p); return nullptr; } @@ -876,6 +923,8 @@ static void prepareRun(Ctx* c, PrepareTask* t) { struct Options { bool serve = false, check = false, raceOnly = false; + bool bench = false, memprobe = false; // read-width experiment (5 October 2026) + int batches = 5, warps = 0, probeMib = 0; int device = 0, batchLog2 = 22, blockWarps = 1; std::string pack, arch = "auto"; std::string race = "on", pinned, tuningPath; @@ -890,6 +939,12 @@ static void usage() { " --batch-log2 B nonces per dispatch = 2^B (default 22)\n" " --block-warps W warps per thread block (default 1)\n" " --arch sm_XY|compute_XY|auto NVRTC target (default auto: the device's architecture)\n" + " --bench --pack read-width experiment: build and self-test the pack, time --batches dispatches of 2^B nonces,\n" + " print the 2^B fingerprint at base nonce 0 (one RESULT line); a variant-5 pack runs --warps persistent warps\n" + " --memprobe [--probe-mib N] no pack: dependent random 4, 16 and 64-byte reads, independent reads, a coalesced stream and an\n" + " integer chain at 4, 64 and 1024 MiB (the same table as igneum-worker-opencl --memprobe)\n" + " --batches N --bench: timed dispatches (default 5)\n" + " --warps N --bench on a variant-5 pack: persistent warps (default: the occupancy capacity, rounded down to a power of two)\n" " --race --pack the variant race alone (3 rounds): one line per variant, the race line, exit 0 or 1\n" " --race on|off|a,b,c in --serve: race every variant (default), none, or these names\n" " --race-bench-ms N timed window per variant (default 2000)\n" @@ -906,6 +961,11 @@ static Options parseArgs(int argc, char** argv) { auto next = [&]() -> std::string { if (i + 1 >= argc) { usage(); std::exit(2); } return argv[++i]; }; if (a == "--serve") o.serve = true; else if (a == "--check") o.check = true; + else if (a == "--bench") o.bench = true; + else if (a == "--memprobe") o.memprobe = true; + else if (a == "--batches") o.batches = std::atoi(next().c_str()); + else if (a == "--warps") o.warps = std::atoi(next().c_str()); + else if (a == "--probe-mib") o.probeMib = std::atoi(next().c_str()); else if (a == "--race" && (i + 1 >= argc || std::string(argv[i + 1]).rfind("--", 0) == 0)) o.raceOnly = true; else if (a == "--race") o.race = next(); else if (a == "--race-bench-ms") o.raceBenchMs = std::atoi(next().c_str()); @@ -924,12 +984,12 @@ static Options parseArgs(int argc, char** argv) { } if (o.batchLog2 < 10 || o.batchLog2 > 28) { std::printf("--batch-log2 must be between 10 and 28\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); } - if (!o.serve && !o.check && !o.raceOnly) { usage(); std::exit(2); } + if (!o.serve && !o.check && !o.raceOnly && !o.bench && !o.memprobe) { usage(); std::exit(2); } if (o.raceBenchMs < 200 || o.raceBenchMs > 20000) { std::printf("--race-bench-ms must be between 200 and 20000\n"); std::exit(2); } if (o.raceBudgetS < 5 || o.raceBudgetS > 540) { std::printf("--race-budget-s must be between 5 and 540 (the prepare lead is 600 DAA)\n"); std::exit(2); } if (o.raceRounds == 0) o.raceRounds = o.raceOnly ? 3 : 1; if (o.tuningPath.empty()) if (const char* t = std::getenv("IGNEUM_TUNING_FILE")) o.tuningPath = t; - if (o.pack.empty()) { std::printf("--pack is required (igneum-miner export-pack writes one)\n"); std::exit(2); } + if (o.pack.empty() && !o.memprobe) { std::printf("--pack is required (igneum-miner export-pack writes one)\n"); std::exit(2); } while (o.pack.size() > 1 && (o.pack.back() == '/' || o.pack.back() == '\\')) o.pack.pop_back(); return o; } @@ -1127,11 +1187,201 @@ static int runServe(Ctx& c, const Options& o, Pair* cur) { // --------------------------------------------------------------------------------------------- // Main +// --------------------------------------------------------------------------------------------- +// Read-width experiment (5 October 2026, docs/plans/read-width.md): --bench and --memprobe + +static uint64_t fnv1a64Bytes(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; +} + +// --bench: the pair is built and self-tested (vectors through the bound kernel with the seed words); then a warm-up +// dispatch at base nonce 0 (fingerprinted) and --batches timed dispatches of 2^B nonces, wall time around +// cuStreamSynchronize (the driver API path loads no event symbols; a 2^24 dispatch is 60 to 900 ms on the cards here, +// so the launch overhead is under 1 percent). +static int runBench(Ctx& c, const Options& o, Pair* p) { + uint32_t nonces = 1u << o.batchLog2, block = 32u * (uint32_t)c.blockWarps; + if (p->persistent) { uint32_t unit = 32u * (uint32_t)p->warps; nonces = (nonces / unit) * unit; if (nonces == 0) nonces = unit; } + CUdeviceptr dOut = 0; + std::string err; + if (c.drv.memAlloc(&dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: cuMemAlloc out\n"); return 2; } + std::vector hOut(nonces); + double sum = 0, warm = 0; + uint64_t fp = 0; + for (int b = -1; b < o.batches; ++b) { + double t0 = wallMs(); + if (!launchHash(c, p, dOut, (uint32_t)(b + 1) * nonces, p->sw, nonces, block, nullptr, err)) { std::printf("FAIL: %s\n", err.c_str()); return 2; } + CUresult r = c.drv.streamSynchronize(nullptr); + if (r != CUDA_SUCCESS) { std::printf("FAIL: dispatch %d: %s\n", b, c.err(r).c_str()); return 2; } + double ms = wallMs() - t0; + if (b < 0) { + warm = ms; + if (c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: read-back\n"); return 2; } + fp = fnv1a64Bytes(hOut.data(), (size_t)nonces * 8u); + } else sum += ms; + } + c.drv.memFree(dOut); + std::string dev = c.name; for (char& ch : dev) if (ch == ' ') ch = '_'; + std::printf("warm-up dispatch (base 0): %.2f ms; %d timed dispatches of %u nonces: mean %.2f ms\n", warm, o.batches, nonces, sum / o.batches); + std::printf("RESULT pack=%s class=%s device=%s arch=%s regs=%d blocks_per_sm=%d warps=%d resident=%d arena_mib=%llu nonces=%u batches=%d check=%s fingerprint=%016llx mhs=%.3f loads=%u bytes=%u scratch_ops=%u time=wall\n", + p->dir.c_str(), p->loadClass.c_str(), dev.c_str(), c.archOpt.c_str(), p->regs, p->blocksPerSM, p->warps, p->residentWarps, (unsigned long long)(p->scratchBytes >> 20), nonces, o.batches, + p->checked ? (p->checkPass ? "PASS" : "FAIL") : "skipped", (unsigned long long)fp, (double)nonces * (double)o.batches / (sum / 1000.0) / 1e6, + p->loadsPerHash, p->bytesPerHash, p->scratchOps * 8u); + return 0; +} + +// --memprobe: the OpenCL worker's table (proto-opencl/host.c, 5 October 2026) in CUDA C through NVRTC, so the two +// vendors are probed with the same access patterns: a dependent chain of random 4-byte reads (the hash's pattern), +// eight independent chains per lane, dependent random 16-byte and 64-byte reads, a coalesced stream and an integer +// chain, at 4, 64 and 1024 MiB. Wall time around cuStreamSynchronize, best of 3, a fresh seed per repetition. +static const char* PROBE_CUDA = + "#include \n" + "__device__ __forceinline__ uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n" + "extern \"C\" __global__ void probe_fill(uint32_t* ds, uint32_t n) { uint32_t i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) ds[i] = pm_mix(i ^ 0x9E3779B9u); }\n" + "extern \"C\" __global__ void probe_chase(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n" + " uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n" + " for (uint32_t s = 0u; s < steps; ++s) x = ds[x & mask] ^ (x * 0x9E3779B1u + s);\n" + " out[g] = x;\n" + "}\n" + "extern \"C\" __global__ void probe_indep(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n" + " uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;\n" + " uint32_t x0 = pm_mix(g * 8u ^ seed), x1 = pm_mix((g * 8u + 1u) ^ seed), x2 = pm_mix((g * 8u + 2u) ^ seed), x3 = pm_mix((g * 8u + 3u) ^ seed);\n" + " uint32_t x4 = pm_mix((g * 8u + 4u) ^ seed), x5 = pm_mix((g * 8u + 5u) ^ seed), x6 = pm_mix((g * 8u + 6u) ^ seed), x7 = pm_mix((g * 8u + 7u) ^ seed);\n" + " for (uint32_t s = 0u; s < steps; ++s) {\n" + " x0 = ds[x0 & mask] ^ (x0 * 0x9E3779B1u + s); x1 = ds[x1 & mask] ^ (x1 * 0x9E3779B1u + s);\n" + " x2 = ds[x2 & mask] ^ (x2 * 0x9E3779B1u + s); x3 = ds[x3 & mask] ^ (x3 * 0x9E3779B1u + s);\n" + " x4 = ds[x4 & mask] ^ (x4 * 0x9E3779B1u + s); x5 = ds[x5 & mask] ^ (x5 * 0x9E3779B1u + s);\n" + " x6 = ds[x6 & mask] ^ (x6 * 0x9E3779B1u + s); x7 = ds[x7 & mask] ^ (x7 * 0x9E3779B1u + s);\n" + " }\n" + " out[g] = x0 ^ x1 ^ x2 ^ x3 ^ x4 ^ x5 ^ x6 ^ x7;\n" + "}\n" + "extern \"C\" __global__ void probe_line16(const uint4* ds, uint32_t vecMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n" + " uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n" + " for (uint32_t s = 0u; s < steps; ++s) { uint4 a = ds[x & vecMask]; x = (a.x ^ a.y ^ a.z ^ a.w) ^ (x * 0x9E3779B1u + s); }\n" + " out[g] = x;\n" + "}\n" + "extern \"C\" __global__ void probe_line(const uint4* ds, uint32_t lineMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n" + " uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n" + " for (uint32_t s = 0u; s < steps; ++s) { uint32_t l = (x & lineMask) * 4u; uint4 a = ds[l], b = ds[l + 1u], c = ds[l + 2u], d = ds[l + 3u]; x = (a.x ^ b.y ^ c.z ^ d.w) ^ (x * 0x9E3779B1u + s); }\n" + " out[g] = x;\n" + "}\n" + "extern \"C\" __global__ void probe_stream(const uint4* ds, uint32_t perLane, uint32_t* out) {\n" + " uint32_t g = blockIdx.x * blockDim.x + threadIdx.x, n = gridDim.x * blockDim.x; uint4 acc = make_uint4(0u, 0u, 0u, 0u);\n" + " for (uint32_t s = 0u; s < perLane; ++s) { uint4 v = ds[s * n + g]; acc.x ^= v.x; acc.y ^= v.y; acc.z ^= v.z; acc.w ^= v.w; }\n" + " out[g] = acc.x ^ acc.y ^ acc.z ^ acc.w;\n" + "}\n" + "extern \"C\" __global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) {\n" + " uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n" + " for (uint32_t s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + ((y << 7u) | (y >> 25u)); y = (y ^ x) + s; }\n" + " out[g] = x ^ y;\n" + "}\n"; + +static double probeLaunch(Ctx& c, CUfunction f, size_t lanes, size_t local, int reps, int seedArg, uint32_t seed, void** args) { + double best = -1; + for (int r = 0; r < reps; ++r) { + uint32_t s = seed + (uint32_t)r * 0x9E3779B9u; + if (seedArg >= 0) args[seedArg] = &s; + double t0 = wallMs(); + if (c.drv.launchKernel(f, (unsigned)(lanes / local), 1, 1, (unsigned)local, 1, 1, 0, nullptr, args, nullptr) != CUDA_SUCCESS) return -1; + if (c.drv.streamSynchronize(nullptr) != CUDA_SUCCESS) return -1; + double ms = wallMs() - t0; + if (best < 0 || ms < best) best = ms; + } + return best; +} + +static int runMemprobe(Ctx& c, const Options& o) { + Compiled cp; + std::string err; + if (!rtcCompile(c, PROBE_CUDA, "probe.cu", "", "", {}, cp, err)) { std::printf("memprobe: build FAILED: %s\n", err.c_str()); return 2; } + CUmodule mod = nullptr; + if (c.drv.moduleLoadData(&mod, cp.image.data()) != CUDA_SUCCESS) { std::printf("memprobe: cuModuleLoadData failed\n"); return 2; } + CUfunction kFill, kChase, kIndep, kLine16, kLine, kStream, kAlu; + const char* names[7] = { "probe_fill", "probe_chase", "probe_indep", "probe_line16", "probe_line", "probe_stream", "probe_alu" }; + CUfunction* fns[7] = { &kFill, &kChase, &kIndep, &kLine16, &kLine, &kStream, &kAlu }; + for (int i = 0; i < 7; ++i) if (c.drv.moduleGetFunction(fns[i], mod, names[i]) != CUDA_SUCCESS) { std::printf("memprobe: %s not in the module\n", names[i]); return 2; } + int sizes[3] = { 4, 64, 1024 }, nSizes = 3; + if (o.probeMib > 0) { sizes[0] = o.probeMib; nSizes = 1; } + const size_t lanesList[8] = { 256, 1024, 1u << 12, 1u << 14, 1u << 16, 1u << 18, 1u << 20, 1u << 22 }; + const size_t groups[2] = { 32, 256 }; + const uint32_t STEPS = 256u, ALU_STEPS = 4096u; + const size_t maxLanes = 1u << 22; + CUdeviceptr dOut = 0; + if (c.drv.memAlloc(&dOut, maxLanes * 4u) != CUDA_SUCCESS) { std::printf("memprobe: cuMemAlloc out\n"); return 2; } + std::printf("memprobe on %s (sm_%d%d, %d SMs, driver %d.%d, NVRTC %d.%d), wall time around cuStreamSynchronize\n", c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor); + std::printf("| probe | MiB | block | lanes in flight | steps per lane | best ms | G loads/s | ns per dependent load |\n|---|---|---|---|---|---|---|---|\n"); + for (int si = 0; si < nSizes; ++si) { + int mib = sizes[si]; + uint64_t bytes = (uint64_t)mib << 20; + uint32_t words = (uint32_t)(bytes / 4ull), mask = words - 1u, n = words; + CUdeviceptr dDs = 0; + if (c.drv.memAlloc(&dDs, (size_t)bytes) != CUDA_SUCCESS) { std::printf("| chase | %d | skipped: cuMemAlloc failed | | | | | |\n", mib); continue; } + { void* a[2] = { &dDs, &n }; probeLaunch(c, kFill, ((size_t)words + 255) / 256 * 256, 256, 1, -1, 0, a); } + for (int gi = 0; gi < 2; ++gi) { + size_t local = groups[gi]; + for (int li = 0; li < 8; ++li) { + size_t lanes = lanesList[li]; + if (lanes < local) continue; + uint32_t seed = 0x1234567u + (uint32_t)li * 977u, steps = STEPS; + void* a[5] = { &dDs, &mask, &steps, &seed, &dOut }; + double ms = probeLaunch(c, kChase, lanes, local, 3, 3, seed, a); + std::printf("| chase | %d | %zu | %zu | %u | %.3f | %.3f | %.0f |\n", mib, local, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, ms * 1e6 / STEPS); + std::fflush(stdout); + } + } + for (size_t lanes = 1u << 16; lanes <= maxLanes; lanes <<= 2) { + uint32_t seed = 0x7654321u, steps = STEPS; + void* a[5] = { &dDs, &mask, &steps, &seed, &dOut }; + double ms = probeLaunch(c, kIndep, lanes, 256, 3, 3, seed, a); + std::printf("| indep x8 | %d | 256 | %zu | %u | %.3f | %.3f | (8 loads in flight per lane) |\n", mib, lanes, STEPS, ms, (double)lanes * 8.0 * STEPS / (ms / 1000.0) / 1e9); + } + for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) { + uint32_t vecMask = (words / 4u) - 1u, seed = 0x2718281u, steps = STEPS; + void* a[5] = { &dDs, &vecMask, &steps, &seed, &dOut }; + double ms = probeLaunch(c, kLine16, lanes, 256, 3, 3, seed, a); + std::printf("| line 16 B | %d | 256 | %zu | %u | %.3f | %.3f G reads/s | %.1f GB/s in 16 B reads |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 16.0 / (ms / 1000.0) / 1e9); + } + for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) { + uint32_t lineMask = (words / 16u) - 1u, seed = 0x3141592u, steps = STEPS; + void* a[5] = { &dDs, &lineMask, &steps, &seed, &dOut }; + double ms = probeLaunch(c, kLine, lanes, 256, 3, 3, seed, a); + std::printf("| line 64 B | %d | 256 | %zu | %u | %.3f | %.3f G lines/s | %.1f GB/s in lines |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 64.0 / (ms / 1000.0) / 1e9); + } + { + size_t lanes = 1u << 20; + uint32_t perLane = (uint32_t)((uint64_t)words / 4ull / (uint64_t)lanes); + if (perLane == 0) { perLane = 1; lanes = (size_t)words / 4u; } + double bytesRead = (double)perLane * (double)lanes * 16.0; + void* a[3] = { &dDs, &perLane, &dOut }; + double ms = probeLaunch(c, kStream, lanes, 256, 3, -1, 0, a); + std::printf("| stream | %d | 256 | %zu | %u | %.3f | %.1f GB/s coalesced | (%.0f MiB read once) |\n", mib, lanes, perLane, ms, bytesRead / (ms / 1000.0) / 1e9, bytesRead / 1048576.0); + } + c.drv.memFree(dDs); + std::fflush(stdout); + } + { + size_t lanes = 1u << 20; + uint32_t seed = 0x2468aceu, steps = ALU_STEPS; + void* a[3] = { &steps, &seed, &dOut }; + double ms = probeLaunch(c, kAlu, lanes, 256, 3, 1, seed, a); + double ops = (double)lanes * ALU_STEPS * 5.0; + std::printf("| alu | 0 | 256 | %zu | %u | %.3f | %.1f G int ops/s | %.3f G steps/s per SM (approximate: 5 ops per step counted) |\n", lanes, ALU_STEPS, ms, ops / (ms / 1000.0) / 1e9, (double)lanes * ALU_STEPS / (ms / 1000.0) / 1e9 / (c.sms ? c.sms : 1)); + } + c.drv.memFree(dOut); + c.drv.moduleUnload(mod); + std::printf("memprobe: done\n"); + return 0; +} + int main(int argc, char** argv) { Options o = parseArgs(argc, argv); Ctx c; c.blockWarps = o.blockWarps; c.race = o.race; c.raceBenchMs = o.raceBenchMs; c.raceBudgetS = o.raceBudgetS; c.raceRounds = o.raceRounds; c.batchLog2 = o.batchLog2; c.pinned = o.pinned; + c.warps = o.warps; c.batches = o.batches; + if (o.bench || o.memprobe) c.race = "off"; if (!o.tuningPath.empty()) { bool ok = false; c.tuning = readText(o.tuningPath, ok); if (!ok) c.tuning.clear(); } std::string err, drvLib, rtcLib; if (!loadDriver(c.drv, err, drvLib)) { emit("error 0 " + err); return 2; } @@ -1140,9 +1390,17 @@ int main(int argc, char** argv) { info(fmt("igneum-worker-cuda %s: device %d %s (sm_%d%d, %d SMs), driver %d.%d from %s, NVRTC %d.%d from %s, target %s (%s)", WORKER_VERSION, o.device, c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, drvLib.c_str(), c.rtcMajor, c.rtcMinor, rtcLib.c_str(), c.archOpt.c_str(), c.why.c_str())); if (!c.tuning.empty()) info(fmt("tuning file %s (%zu bytes): %s", o.tuningPath.c_str(), c.tuning.size(), readTuning(c.tuning, c.name).found ? "has an entry for this card" : "no entry for this card")); + if (o.memprobe) { int rc = runMemprobe(c, o); c.drv.primaryCtxRelease(c.dev); return rc; } double t0 = wallMs(); - Pair* cur = buildPair(c, o.pack, nullptr, err, !o.check); + Pair* cur = buildPair(c, o.pack, nullptr, err, !o.check && !o.bench); if (!cur) { emit("error 0 " + err); return 1; } + if (o.bench) { + std::printf("pack %s on %s: %s\n", o.pack.c_str(), c.name.c_str(), pairSummary(cur).c_str()); + int rc = runBench(c, o, cur); + releasePair(c, cur); + c.drv.primaryCtxRelease(c.dev); + return rc; + } if (o.raceOnly) { std::printf("race %s on %s (%s, %d SMs, driver %d.%d, NVRTC %d.%d, %s): %s\n", o.pack.c_str(), c.name.c_str(), c.archOpt.c_str(), c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor, c.why.c_str(), pairSummary(cur).c_str()); std::printf("%s\n", cur->raceLine.c_str()); diff --git a/proto-metal/packbench.swift b/proto-metal/packbench.swift index 75b6fc4f3..3110bbe12 100644 --- a/proto-metal/packbench.swift +++ b/proto-metal/packbench.swift @@ -205,5 +205,5 @@ print("device \(device.name); compile \(String(format: "%.0f", compileMs)) ms; c print("cache FNV-1a 64 \(String(format: "%016llx", cacheFnv)) \(cacheOk ? "PASS" : "FAIL"); dataset head and last \(dsOk ? "PASS" : "FAIL"); vectors standalone \(vecPass)/\(vecBases.count), in batch \(batchVecPass)/\(batchVecN)") print("warm-up batch \(nonces) hashes: \(String(format: "%.1f", warmGpu)) ms GPU, \(String(format: "%.1f", warmWall)) ms wall") let overall = cacheOk && dsOk && vecPass == vecBases.count && batchVecPass == batchVecN -print("RESULT pack=\(packName) class=\(className) device=\(device.name.replacingOccurrences(of: " ", with: "_")) group=\(opts.group) warps=\(warpsN) arena_mib=\(persistent ? warpsN : 0) nonces=\(nonces) batches=\(opts.batches) vectors=\(vecPass)/\(vecBases.count) batch_vectors=\(batchVecPass)/\(batchVecN) cache=\(cacheOk ? "PASS" : "FAIL") dataset=\(dsOk ? "PASS" : "FAIL") fingerprint=\(String(format: "%016llx", fingerprint)) mhs_gpu=\(String(format: "%.3f", mhsGpu)) mhs_wall=\(String(format: "%.3f", mhsWall)) loads=\(loadsPerHash) bytes=\(bytesPerHash) scratch_ops=\(scratchOps * 8) overall=\(overall ? "PASS" : "FAIL")") +print("RESULT pack=\(packName) class=\(className) device=\(device.name.replacingOccurrences(of: " ", with: "_")) group=\(opts.group) warps=\(warpsN) arena_mib=\(persistent ? warpsN * 32 * scratchWordsPerLane * 4 / 1048576 : 0) nonces=\(nonces) batches=\(opts.batches) vectors=\(vecPass)/\(vecBases.count) batch_vectors=\(batchVecPass)/\(batchVecN) cache=\(cacheOk ? "PASS" : "FAIL") dataset=\(dsOk ? "PASS" : "FAIL") fingerprint=\(String(format: "%016llx", fingerprint)) mhs_gpu=\(String(format: "%.3f", mhsGpu)) mhs_wall=\(String(format: "%.3f", mhsWall)) loads=\(loadsPerHash) bytes=\(bytesPerHash) scratch_ops=\(scratchOps * 8) overall=\(overall ? "PASS" : "FAIL")") exit(overall ? 0 : 1) diff --git a/relay/playbooks/readwidth-5090.ps1 b/relay/playbooks/readwidth-5090.ps1 new file mode 100644 index 000000000..c9384f214 --- /dev/null +++ b/relay/playbooks/readwidth-5090.ps1 @@ -0,0 +1,67 @@ +# Igneum run job: the read-width experiment on PC 2's RTX 5090 (machine 1ccfe586), 5 October 2026 (docs/plans/read-width.md). +# Published as a plain `run` job (NOT --stop-miners): the installed app keeps every other card mining; this script switches +# off ONLY the NVIDIA card in the app through POST api/cards, waits for its worker to stop, runs the fetched +# igneum-worker-cuda.exe (--memprobe, then --bench on every pack of the fetched packs folder), and switches the card back +# on with the settings it had. Every result line starts with RESULT so `node tools/jobs.mjs ` shows them. +$ErrorActionPreference = 'Continue' +function Say([string] $m) { Write-Host ("[" + (Get-Date -Format 'HH:mm:ss') + "] " + $m) } +$jobs = Split-Path $env:IGNEUM_JOB_DIR +$fetched = Join-Path $jobs 'fetch-readwidth-20261005' +$exe = Join-Path $fetched 'igneum-worker-cuda.exe' +$packs = Join-Path $fetched 'packs-readwidth' +if (-not (Test-Path $exe)) { Write-Output "RESULT error worker missing at $exe (the fetch job runs first)"; exit 2 } +if (-not (Test-Path $packs)) { Write-Output "RESULT error packs missing at $packs"; exit 2 } +$inst = @("$env:LOCALAPPDATA\Programs\Igneum Miner", "$env:ProgramFiles\Igneum Miner") | Where-Object { Test-Path (Join-Path $_ 'igneum-worker-cuda.exe') } | Select-Object -First 1 +if (-not $inst) { Write-Output 'RESULT error no installed igneum-worker-cuda.exe (the NVRTC DLLs come from there)'; exit 2 } +Get-ChildItem $inst -Filter 'nvrtc*.dll' | Copy-Item -Destination $fetched -Force +Write-Output "RESULT worker $exe sha256 $((Get-FileHash -Algorithm SHA256 $exe).Hash.ToLower()) with $((Get-ChildItem $fetched -Filter 'nvrtc*.dll').Count) NVRTC DLL(s) from $inst" + +# the app: switch off the NVIDIA card only, remember its settings +$appDir = $env:IGNEUM_APP_DIR +if (-not $appDir) { $appDir = Join-Path $env:LOCALAPPDATA 'igneum\app' } +$urlFile = Join-Path $appDir 'app.url' +$url = $null +if (Test-Path $urlFile) { $url = (Get-Content -LiteralPath $urlFile -Raw).Trim() } +$card = $null +if ($url) { + try { + $st = Invoke-RestMethod -Uri ($url + 'api/state') -Method GET -TimeoutSec 10 + $card = $st.mining.cards | Where-Object { $_.vendor -eq 'nvidia' } | Select-Object -First 1 + if (-not $card) { $card = $st.cards | Where-Object { $_.vendor -eq 'nvidia' } | Select-Object -First 1 } + } catch { Say ("api/state: " + $_.Exception.Message) } +} +if ($card) { + Write-Output ("RESULT card " + $card.key + " enabled=" + $card.enabled + " identities=" + $card.identities + " power_pct=" + $card.power_pct + " state=" + $card.state) + $body = @{ cards = @(@{ key = $card.key; enabled = $false; identities = [int]$card.identities; power_pct = [int]$card.power_pct }) } | ConvertTo-Json -Depth 5 + try { Invoke-RestMethod -Uri ($url + 'api/cards') -Method POST -Body $body -ContentType 'application/json' -TimeoutSec 10 | Out-Null; Say "card off requested" } catch { Say ("api/cards off: " + $_.Exception.Message) } + $t = 0 + while ($t -lt 90) { + Start-Sleep -Seconds 5; $t += 5 + try { $st = Invoke-RestMethod -Uri ($url + 'api/state') -Method GET -TimeoutSec 10; $c2 = $st.mining.cards | Where-Object { $_.key -eq $card.key }; if (-not $c2) { $c2 = $st.cards | Where-Object { $_.key -eq $card.key } }; if ($c2 -and $c2.state -eq 'off' -and $c2.pid -eq 0) { break } } catch { } + } + Write-Output ("RESULT card-off after " + $t + " s") + Start-Sleep -Seconds 5 +} else { Write-Output 'RESULT card none-found (the app is not running or has no NVIDIA card); measuring with whatever else runs on the GPU' } + +& nvidia-smi --query-gpu=name,driver_version,power.limit,clocks.sm,clocks.mem,memory.used,temperature.gpu --format=csv,noheader 2>&1 | ForEach-Object { "RESULT gpu-before $_" } +Write-Output "RESULT memprobe start $(Get-Date -Format HH:mm:ss)" +& $exe --memprobe 2>&1 | ForEach-Object { "RESULT $_" } +foreach ($pk in @('w4', 'w16', 'w64', 'w64x4', 'mixA-0', 'mixA-1', 'mixA-2', 'mixA-3', 'mixA-4', 'mixA-5', 'mixB-0', 'mixB-1', 'mixB-2', 'mixB-3', 'mixB-4', 'mixB-5')) { + $d = Join-Path $packs $pk + Write-Output "RESULT bench $pk start $(Get-Date -Format HH:mm:ss)" + & $exe --bench --pack $d --batches 5 --batch-log2 24 --block-warps 1 2>&1 | ForEach-Object { "RESULT $_" } + & $exe --bench --pack $d --batches 5 --batch-log2 24 --block-warps 8 2>&1 | Where-Object { $_ -match '^RESULT|error|FAIL' } | ForEach-Object { "RESULT $_" } +} +foreach ($pk in @('scr0k32', 'scr2k32', 'scr4k32', 'scr8k32', 'scr2k128', 'scr4k128', 'scr8k128')) { + $d = Join-Path $packs $pk + Write-Output "RESULT bench $pk start $(Get-Date -Format HH:mm:ss)" + & $exe --bench --pack $d --batches 5 --batch-log2 24 --block-warps 1 2>&1 | ForEach-Object { "RESULT $_" } + & $exe --bench --pack $d --batches 5 --batch-log2 24 --block-warps 1 --warps 4096 2>&1 | Where-Object { $_ -match '^RESULT|error|FAIL|variant 5' } | ForEach-Object { "RESULT $_" } +} +& nvidia-smi --query-gpu=power.draw,clocks.sm,clocks.mem,memory.used,temperature.gpu --format=csv,noheader 2>&1 | ForEach-Object { "RESULT gpu-after $_" } + +if ($card) { + $body = @{ cards = @(@{ key = $card.key; enabled = [bool]$card.enabled; identities = [int]$card.identities; power_pct = [int]$card.power_pct }) } | ConvertTo-Json -Depth 5 + try { Invoke-RestMethod -Uri ($url + 'api/cards') -Method POST -Body $body -ContentType 'application/json' -TimeoutSec 10 | Out-Null; Write-Output ("RESULT card restored enabled=" + $card.enabled) } catch { Write-Output ("RESULT error card restore: " + $_.Exception.Message) } +} +exit 0 diff --git a/relay/playbooks/readwidth-9070.ps1 b/relay/playbooks/readwidth-9070.ps1 new file mode 100644 index 000000000..4e44dea09 --- /dev/null +++ b/relay/playbooks/readwidth-9070.ps1 @@ -0,0 +1,69 @@ +# Igneum run job: the read-width experiment on PC 1's RX 9070 XT on the eGPU (machine ae432dc7), 5 October 2026 (docs/plans/read-width.md). +# Published as a plain `run` job (NOT --stop-miners): the installed app keeps every other card mining; this script switches +# off ONLY the NVIDIA card in the app through POST api/cards, waits for its worker to stop, runs the fetched +# igneum-worker-opencl.exe (--memprobe, then --bench on every pack of the fetched packs folder), and switches the card back +# on with the settings it had. Every result line starts with RESULT so `node tools/jobs.mjs ` shows them. +$ErrorActionPreference = 'Continue' +function Say([string] $m) { Write-Host ("[" + (Get-Date -Format 'HH:mm:ss') + "] " + $m) } +$jobs = Split-Path $env:IGNEUM_JOB_DIR +$fetched = Join-Path $jobs 'fetch-readwidth-20261005' +$exe = Join-Path $fetched 'igneum-worker-opencl.exe' +$packs = Join-Path $fetched 'packs-readwidth' +if (-not (Test-Path $exe)) { Write-Output "RESULT error worker missing at $exe (the fetch job runs first)"; exit 2 } +if (-not (Test-Path $packs)) { Write-Output "RESULT error packs missing at $packs"; exit 2 } +Write-Output "RESULT worker $exe sha256 $((Get-FileHash -Algorithm SHA256 $exe).Hash.ToLower())" +# the card's OpenCL device index on the current (3683.0) platform, from the worker's own list (the older platform's duplicate is marked dup) +$list = & $exe --list 2>&1 +$list | ForEach-Object { "RESULT list $_" } +$dev = $null +foreach ($l in $list) { if ($l -match '^\s*\[(\d+)\].*gfx1201' -and $l -notmatch 'dup') { $dev = [int]$Matches[1]; break } } +if ($null -eq $dev) { Write-Output 'RESULT error no gfx1201 device in --list'; exit 2 } +Write-Output "RESULT device $dev" + +# the app: switch off the NVIDIA card only, remember its settings +$appDir = $env:IGNEUM_APP_DIR +if (-not $appDir) { $appDir = Join-Path $env:LOCALAPPDATA 'igneum\app' } +$urlFile = Join-Path $appDir 'app.url' +$url = $null +if (Test-Path $urlFile) { $url = (Get-Content -LiteralPath $urlFile -Raw).Trim() } +$card = $null +if ($url) { + try { + $st = Invoke-RestMethod -Uri ($url + 'api/state') -Method GET -TimeoutSec 10 + $card = $st.mining.cards | Where-Object { ($_.vendor -eq 'amd' -and $_.key -match 'gfx1201') } | Select-Object -First 1 + if (-not $card) { $card = $st.cards | Where-Object { ($_.vendor -eq 'amd' -and $_.key -match 'gfx1201') } | Select-Object -First 1 } + } catch { Say ("api/state: " + $_.Exception.Message) } +} +if ($card) { + Write-Output ("RESULT card " + $card.key + " enabled=" + $card.enabled + " identities=" + $card.identities + " power_pct=" + $card.power_pct + " state=" + $card.state) + $body = @{ cards = @(@{ key = $card.key; enabled = $false; identities = [int]$card.identities; power_pct = [int]$card.power_pct }) } | ConvertTo-Json -Depth 5 + try { Invoke-RestMethod -Uri ($url + 'api/cards') -Method POST -Body $body -ContentType 'application/json' -TimeoutSec 10 | Out-Null; Say "card off requested" } catch { Say ("api/cards off: " + $_.Exception.Message) } + $t = 0 + while ($t -lt 90) { + Start-Sleep -Seconds 5; $t += 5 + try { $st = Invoke-RestMethod -Uri ($url + 'api/state') -Method GET -TimeoutSec 10; $c2 = $st.mining.cards | Where-Object { $_.key -eq $card.key }; if (-not $c2) { $c2 = $st.cards | Where-Object { $_.key -eq $card.key } }; if ($c2 -and $c2.state -eq 'off' -and $c2.pid -eq 0) { break } } catch { } + } + Write-Output ("RESULT card-off after " + $t + " s") + Start-Sleep -Seconds 5 +} else { Write-Output 'RESULT card none-found (the app is not running or has no NVIDIA card); measuring with whatever else runs on the GPU' } + +Write-Output "RESULT memprobe start $(Get-Date -Format HH:mm:ss)" +& $exe --device $dev --memprobe 2>&1 | ForEach-Object { "RESULT $_" } +foreach ($pk in @('w4', 'w16', 'w64', 'w64x4', 'mixA-0', 'mixA-1', 'mixA-2', 'mixA-3', 'mixA-4', 'mixA-5', 'mixB-0', 'mixB-1', 'mixB-2', 'mixB-3', 'mixB-4', 'mixB-5')) { + $d = Join-Path $packs $pk + Write-Output "RESULT bench $pk start $(Get-Date -Format HH:mm:ss)" + & $exe --bench-pack --pack $d --batches 5 --batch-log2 24 --device $dev 2>&1 | ForEach-Object { "RESULT $_" } + & $exe --bench-pack --pack $d --batches 5 --batch-log2 24 --device $dev --group-warps 8 2>&1 | Where-Object { $_ -match '^RESULT|error|FAIL' } | ForEach-Object { "RESULT $_" } +} +foreach ($pk in @('scr0k32', 'scr2k32', 'scr4k32', 'scr8k32', 'scr2k128', 'scr4k128', 'scr8k128')) { + $d = Join-Path $packs $pk + Write-Output "RESULT bench $pk start $(Get-Date -Format HH:mm:ss)" + & $exe --bench-pack --pack $d --batches 5 --batch-log2 24 --device $dev 2>&1 | ForEach-Object { "RESULT $_" } + & $exe --bench-pack --pack $d --batches 5 --batch-log2 24 --device $dev --warps 4096 2>&1 | Where-Object { $_ -match '^RESULT|error|FAIL|variant 5' } | ForEach-Object { "RESULT $_" } +} + +if ($card) { + $body = @{ cards = @(@{ key = $card.key; enabled = [bool]$card.enabled; identities = [int]$card.identities; power_pct = [int]$card.power_pct }) } | ConvertTo-Json -Depth 5 + try { Invoke-RestMethod -Uri ($url + 'api/cards') -Method POST -Body $body -ContentType 'application/json' -TimeoutSec 10 | Out-Null; Write-Output ("RESULT card restored enabled=" + $card.enabled) } catch { Write-Output ("RESULT error card restore: " + $_.Exception.Message) } +} +exit 0