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 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
ea863d6d4b
commit
b4207ab2d0
4 changed files with 400 additions and 6 deletions
|
|
@ -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 <dir> 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 <dir> 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 <dir> is required (igneum-miner export-pack <node> <dir> writes one)\n"); std::exit(2); }
|
||||
if (o.pack.empty() && !o.memprobe) { std::printf("--pack <dir> is required (igneum-miner export-pack <node> <dir> 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<uint64_t> 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 <cstdint>\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());
|
||||
|
|
|
|||
|
|
@ -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)
|
||||
|
|
|
|||
67
relay/playbooks/readwidth-5090.ps1
Normal file
67
relay/playbooks/readwidth-5090.ps1
Normal file
|
|
@ -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 <app.url>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 <job id>` 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
|
||||
69
relay/playbooks/readwidth-9070.ps1
Normal file
69
relay/playbooks/readwidth-9070.ps1
Normal file
|
|
@ -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 <app.url>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 <job id>` 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
|
||||
Loading…
Reference in a new issue