OpenCL worker fault guard, miner-side fault state in the launcher, NVRTC annotation rule in the emulation
After PC 2's gfx1036 (prebuilt-generic path) completed 2,000 jobs a second with no hash from 600 s on: every OpenCL call in host.c's job path is now fatal on error (exit 3, the miner restarts the worker), the dispatch event must read CL_COMPLETE, a chunk 20x faster per nonce than the running mean or an output buffer unchanged since the previous dispatch is a fault, and a stats line every 200 jobs carries the live event and buffer counts (a leak over 16 events or 12 buffers is fatal too). IGNEUM_FAULT_TEST=N exercises the detectors on a healthy device (verified on Apple OpenCL: the stale-output guard fires on the chunk after the injected fault). The launcher shows 'worker fault' and 'restarting' for the card on the miner's WORKER FAULT line and drops the last rate. The emulation's NVRTC stand-in now applies NVRTC's execution-space rule (program.h(46) igneum_launch_* declarations are host code unless -default-device or -DIGNEUM_NO_CUDA), which is what the RTX 5090 reported; test.sh checks the rejection. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
4574890602
commit
4714b510f2
6 changed files with 163 additions and 29 deletions
|
|
@ -18,6 +18,8 @@
|
|||
#include <vector>
|
||||
#include <list>
|
||||
#include <map>
|
||||
#include <regex>
|
||||
#include <sstream>
|
||||
|
||||
#include "../cuda_api.h" // the real cuda.h / nvrtc.h types
|
||||
#include "cuda_runtime.h" // the shim (proto-cuda/emu): emu_launch and the kernel symbols' world
|
||||
|
|
@ -64,9 +66,44 @@ static const char* packDir(int which) {
|
|||
return e ? e : "";
|
||||
}
|
||||
|
||||
// NVRTC's rule, which bit on the RTX 5090 on 4 October 2026 ("A function without execution space annotations
|
||||
// (__host__/__device__/__global__) is considered a host function, and host functions are not allowed in JIT mode",
|
||||
// program.h line 46, the igneum_launch_* declarations): every function declaration or definition in the source and in
|
||||
// the pack headers must carry an annotation (IGNEUM_HD expands to one under __CUDACC__, which NVRTC defines), unless
|
||||
// the compile passes -default-device (then unannotated functions are device code) or the text is inside an
|
||||
// #ifndef IGNEUM_NO_CUDA block and -DIGNEUM_NO_CUDA was passed. Returns the first offender as "file(line): text" or "".
|
||||
// IGNEUM_EMU_NO_DEFAULT_DEVICE=1 makes the stand-in ignore -default-device, to test this check.
|
||||
static std::string firstUnannotated(const std::string& name, const std::string& text, bool noCuda) {
|
||||
std::istringstream in(text);
|
||||
std::string line;
|
||||
int ln = 0, skip = 0, depth = 0;
|
||||
std::vector<bool> skipping;
|
||||
static const std::regex fn(R"(^\s*(static\s+|inline\s+|extern\s+)*[A-Za-z_][A-Za-z_0-9:<>]*\s*\**\s+\**\s*[A-Za-z_][A-Za-z_0-9]*\s*\()");
|
||||
while (std::getline(in, line)) {
|
||||
++ln;
|
||||
std::string t = line;
|
||||
if (t.rfind("#ifndef IGNEUM_NO_CUDA", 0) == 0) { skipping.push_back(noCuda); if (noCuda) ++skip; ++depth; continue; }
|
||||
if (t.rfind("#if", 0) == 0) { skipping.push_back(false); ++depth; continue; }
|
||||
if (t.rfind("#endif", 0) == 0) { if (!skipping.empty()) { if (skipping.back()) --skip; skipping.pop_back(); --depth; } continue; }
|
||||
if (skip > 0) continue;
|
||||
if (t.empty() || t[0] == '#' || t.find("//") == 0 || t.find("/*") == 0 || t.find(" *") == 0) continue;
|
||||
if (t.find("__device__") != std::string::npos || t.find("__global__") != std::string::npos || t.find("__host__") != std::string::npos || t.find("IGNEUM_HD") != std::string::npos) continue;
|
||||
if (t.find("typedef") != std::string::npos || t.find("struct ") == 0 || t.find("return") != std::string::npos || t.find("if (") != std::string::npos || t.find("for (") != std::string::npos || t.find("while (") != std::string::npos) continue;
|
||||
if (std::regex_search(t, fn)) return name + "(" + std::to_string(ln) + "): " + t;
|
||||
}
|
||||
return "";
|
||||
}
|
||||
|
||||
// The source check. Fills p->pack, p->ok, p->log.
|
||||
static void checkSource(EmuProg* p) {
|
||||
static void checkSource(EmuProg* p, bool defaultDevice, bool noCuda) {
|
||||
p->checked = true;
|
||||
if (std::getenv("IGNEUM_EMU_NO_DEFAULT_DEVICE")) defaultDevice = false;
|
||||
if (!defaultDevice) {
|
||||
std::string bad = firstUnannotated(p->name, p->src, noCuda);
|
||||
if (bad.empty() && p->headers.count("program.h")) bad = firstUnannotated("program.h", p->headers["program.h"], noCuda);
|
||||
if (bad.empty() && p->headers.count("memhard.h")) bad = firstUnannotated("memhard.h", p->headers["memhard.h"], noCuda);
|
||||
if (!bad.empty()) { p->log = "emu-nvrtc: " + bad + ": A function without execution space annotations (__host__/__device__/__global__) is considered a host function, and host functions are not allowed in JIT mode (pass -default-device or -DIGNEUM_NO_CUDA)"; p->ok = false; return; }
|
||||
}
|
||||
for (int which = 1; which <= 2; ++which) {
|
||||
std::string dir = packDir(which);
|
||||
if (dir.empty()) continue;
|
||||
|
|
@ -104,10 +141,10 @@ static nvrtcResult e_create(nvrtcProgram* prog, const char* src, const char* nam
|
|||
static nvrtcResult e_destroy(nvrtcProgram* prog) { delete (EmuProg*)*prog; *prog = nullptr; return NVRTC_SUCCESS; }
|
||||
static nvrtcResult e_compile(nvrtcProgram prog, int numOptions, const char* const* options) {
|
||||
EmuProg* p = (EmuProg*)prog;
|
||||
bool arch = false, std17 = false;
|
||||
for (int i = 0; i < numOptions; ++i) { std::string o = options[i]; if (o.rfind("--gpu-architecture=", 0) == 0) arch = true; if (o == "--std=c++17") std17 = true; }
|
||||
bool arch = false, std17 = false, defaultDevice = false, noCuda = false;
|
||||
for (int i = 0; i < numOptions; ++i) { std::string o = options[i]; if (o.rfind("--gpu-architecture=", 0) == 0) arch = true; if (o == "--std=c++17") std17 = true; if (o == "-default-device" || o == "--device-as-default-execution-space") defaultDevice = true; if (o == "-DIGNEUM_NO_CUDA" || o == "-DIGNEUM_NO_CUDA=1") noCuda = true; }
|
||||
if (!arch || !std17) { p->log = "emu-nvrtc: expected --gpu-architecture=... and --std=c++17"; return NVRTC_ERROR_COMPILATION; }
|
||||
checkSource(p);
|
||||
checkSource(p, defaultDevice, noCuda);
|
||||
if (!p->ok) return NVRTC_ERROR_COMPILATION;
|
||||
p->compiled = true;
|
||||
return NVRTC_SUCCESS;
|
||||
|
|
|
|||
|
|
@ -56,6 +56,11 @@ echo "== --check pack A"
|
|||
grep -q '^check PASS' "$OUT/check-a.log"
|
||||
grep -c 'source check PASS' "$OUT/check-a.log" | grep -qx 2
|
||||
|
||||
echo "== the annotation rule: without -default-device the stand-in must reject program.h's igneum_launch_* declarations as host code (what the RTX 5090 reported on 4 October 2026)"
|
||||
if IGNEUM_EMU_NO_DEFAULT_DEVICE=1 "$OUT/igneum-worker-cuda-emu" --check --pack "$OUT/pack-a" > "$OUT/check-nodefault.log" 2>&1; then echo "FAIL: the unannotated program.h declarations were not flagged"; exit 1; fi
|
||||
grep -q 'program.h(4[0-9]): cudaError_t igneum_launch_cache_fill.*host functions are not allowed in JIT mode' "$OUT/check-nodefault.log" || { echo "FAIL: wrong rejection:"; cat "$OUT/check-nodefault.log"; exit 1; }
|
||||
echo "flagged as NVRTC does: $(grep -o 'program.h([0-9]*): [^:]*' "$OUT/check-nodefault.log" | head -1)"
|
||||
|
||||
echo "== --serve: two jobs on A, prepare B, a job on B (swap), a job on A after the swap (error), quit"
|
||||
"$HERE/serve-check.sh" "$OUT/pack-a" "$OUT/pack-b" "$OUT/serve.log" "$OUT/igneum-worker-cuda-emu" --serve --pack "$OUT/pack-a" --batch-log2 13
|
||||
echo "NVRTC source check PASS for $(grep -c 'source check PASS' "$OUT/serve.log") files in --serve, $(grep -c 'source check PASS' "$OUT/check-a.log") in --check"
|
||||
|
|
|
|||
|
|
@ -7,6 +7,7 @@ Devnet v4 is a new chain from genesis. The node keeps its database in %LOCALAPPD
|
|||
3. Ctrl+C in the window stops the miners first, then the node, then uploads the logs one last time; the window stays open with a summary. If the window was closed instead, STOP-IGNEUM.bat stops everything cleanly.
|
||||
Payout: block rewards go to an EVM address. PAYOUT_EVM at the top of the bat sets one for the whole PC; left empty, the launcher derives one address per GPU vendor from this PC's name (the same on every run; it is printed in the launcher log and the status block). Finality voting is on: every identity signs every checkpoint (VOTE=0 switches it off).
|
||||
The hourly program change: every hour the lottery program changes. About ten minutes before the boundary the node announces the next program; the miner writes its pack (proto-cuda\packs\prepare\<epoch>-<day>) and tells the worker, which compiles it on the card in the background, builds its cache and dataset, self-tests it and keeps it ready; at the boundary the worker swaps with no pause (the dashboard says "the next hourly program is compiled and resident"). If that ever fails, the CUDA worker finds the pack by itself when the first job on the new program arrives and compiles it then (a pause of a few seconds); and if a worker keeps erroring on the new seeds for 90 s, the launcher re-exports the pack and restarts that card's miner. --exit-on-seed-change stays on as the last fallback: a worker that cannot prepare at all exits with code 42 at the boundary and the launcher restarts it on a fresh pack.
|
||||
Worker faults (added after the first field run on 4 October 2026, when an integrated AMD card's OpenCL runtime stopped running kernels after 600 s and kept answering every call with success): the OpenCL worker now treats any OpenCL error in a job as fatal, checks that each dispatch really completed, that it did not finish 20x faster per nonce than before and that its output changed, and exits with code 3 so the miner restarts it; every 200 jobs it prints a stats line with the live event and buffer counts. The miner itself kills and restarts a worker whose jobs finish in under 1/20 of the mean time per hash or whose rate jumps over 10x, drops the fake numbers, and prints "WORKER FAULT ..."; the dashboard then shows "worker fault" and "restarting" for that card instead of a rate, and the status block counts faults.
|
||||
If a prebuilt worker never reports ready (150 s), the launcher says so in the events; when a CUDA Toolkit and Visual Studio happen to be installed it builds a worker from source with them instead, otherwise look at the miner log named in the event (the worker prints the reason: a missing DLL, a compile error, a driver too old for CUDA 12).
|
||||
A node crash is restarted after a random 5 to 60 s; the miners are restarted once the node is back and synced (their connection dies with the node). The PC is held awake while the window runs; for the night also set Sleep to Never in Settings > System > Power.
|
||||
Logs land next to the bat (igneum-<stamp>.log for the launcher, node-<stamp>.log for the node, nvidia-<stamp>.log and amd-<stamp>.log per miner process; the worker's ready, info and self-test lines are in the miner log) and are uploaded every 60 s with upload-log.bat under the labels igneum-<PC>, nodelog-<PC>, nvidia-<PC>, amd-<PC>.
|
||||
|
|
|
|||
|
|
@ -80,6 +80,27 @@ hashes it on the CPU, so a few hundred ms). These go into docs/bench-log.md with
|
|||
changes to "worker built here with nvcc + MSVC"); the miner log of the first attempt holds the reason.
|
||||
- `FORCE_BUILD=1` at the top of the bat forces the old nvcc / cl.exe path for a comparison run.
|
||||
|
||||
## The fault guard (second field run)
|
||||
|
||||
PC 2's gfx1036 ran the OpenCL worker correctly for 577 s, then every job "completed" in half a millisecond with no
|
||||
hash: the runtime answered every call with success without running the kernel. Three guards now sit on that path,
|
||||
and this is what they print:
|
||||
|
||||
- the worker (amd-<stamp>.log, from the worker): `error <job> worker fault: <what was seen>; exiting 3 so the miner
|
||||
restarts the worker`, then the miner's `worker exited (code Some(3)); restarting it in N s`. The "what was seen" is
|
||||
one of: an OpenCL call failing (name and code), the dispatch event not CL_COMPLETE, a chunk 20x faster per nonce
|
||||
than the running mean, or the output buffer unchanged since the previous dispatch. Every 200 jobs the worker prints
|
||||
`info stats jobs N ... events created X released Y live Z; buffers created ... live ...`: live should stay at 0 for
|
||||
events and 4 for buffers (cache, dataset, out, init words; 6 while a prepared pair is resident).
|
||||
- the miner (same log): `WORKER FAULT job ... reported done in ... ms, Nx faster than the running mean ...; restarting
|
||||
the worker (fault 1)` or the 10x interval form; the STATUS line carries `faults=N` and its rates are rolled back to
|
||||
the last report, so the dashboard never shows the fake GH/s again.
|
||||
- the launcher: the card's cell says `worker fault` in red and `restarting` where the rate was, the events list
|
||||
`<card> worker fault, restarting: ...`, and the 30-s status block says `worker faults N`.
|
||||
|
||||
If the fault returns on the gfx1036, the first `worker fault` line names which guard fired and that is the clue to
|
||||
the runtime's failure mode; please send the amd log around it.
|
||||
|
||||
## What to send back
|
||||
|
||||
The launcher log, the two miner logs (they are uploaded every minute as well) and, for the bench log, the three
|
||||
|
|
|
|||
|
|
@ -672,7 +672,7 @@ function Initialize-Vendors($plan) {
|
|||
runId = "$name-$machine$suffix-$stamp"; proc = $null; restarts = 0; starts = 0; startedAt = $null; restartAt = $null; exitCode = $null;
|
||||
logPos = [long]0; logRem = ''; errPos = [long]0; errRem = ''; accepted = 0; found = (New-Object System.Collections.ArrayList);
|
||||
status = $null; ready = $false; readyAt = $null; lastLineAt = $null; errors = 0; lastError = ''; lastErrorAt = $null; rejected = 0; lastNoNodeAt = $null;
|
||||
identities = $identitiesPerProc; prepare = $null; mismatchSince = $null; lastMismatchAt = $null }
|
||||
identities = $identitiesPerProc; prepare = $null; mismatchSince = $null; lastMismatchAt = $null; faults = 0; lastFaultAt = $null }
|
||||
}
|
||||
$script:vendors[$name] = @{ name = $name; card = (Get-CardName $name); exe = $exe; instances = $instances; rebuilds = 0; building = $false; evm = (Get-EvmAddress $name);
|
||||
path = $script:workerPath[$name]; pathText = (Get-PathText $name); fellBack = $false }
|
||||
|
|
@ -820,6 +820,11 @@ function Update-Instance($v, $inst) {
|
|||
# the boundary with no exit; "prepare 0" (no nvcc on PATH) means the old exit-42 rebuild path.
|
||||
if ($l -match ' prepare 1') { $inst.prepare = $true } elseif ($l -match ' prepare 0') { $inst.prepare = $false }
|
||||
if ($inst.starts -eq 1) { Add-Event ("$($v.card) worker ready ($($v.pathText))" + $(if ($inst.prepare -eq $true) { ', hot swap at the hour boundary' } elseif ($inst.prepare -eq $false) { ', no hot swap (nvcc missing): rebuild at the boundary' } else { '' })) 'ok' }
|
||||
} elseif ($l -match ' WORKER FAULT (.*)$') {
|
||||
# The miner's guard: the worker answered jobs without running its kernel (or 10x the rate); the miner kills and
|
||||
# restarts it. The last rate is dropped so the dashboard never shows the fake number.
|
||||
$inst.faults += 1; $inst.lastFaultAt = Get-Date; $inst.status = $null
|
||||
Add-Event ("$($v.card) worker fault, restarting: " + $Matches[1].Substring(0, [Math]::Min(110, $Matches[1].Length))) 'error'
|
||||
} elseif ($l -match ' worker: prepared ') {
|
||||
Add-Event "$($v.card): the next hourly program is compiled and resident (hot swap ready)" 'build'
|
||||
} elseif ($l -match ' worker: prepare-failed ') {
|
||||
|
|
@ -835,7 +840,7 @@ function Update-Instance($v, $inst) {
|
|||
}
|
||||
if ($l -match ' rejected nonce=') { $inst.rejected += 1 }
|
||||
elseif ($l -match '^template error') { $inst.lastNoNodeAt = Get-Date }
|
||||
elseif ($l -match 'submit error|template error|WORKER MISMATCH|worker error|worker exited|could not prepare|above target|panicked|CUDA error|FAIL') {
|
||||
elseif ($l -match 'submit error|template error|WORKER MISMATCH|WORKER FAULT|worker fault|worker error|worker exited|could not prepare|above target|panicked|CUDA error|FAIL') {
|
||||
$inst.errors += 1; $inst.lastError = $l; $inst.lastErrorAt = Get-Date
|
||||
}
|
||||
}
|
||||
|
|
@ -958,6 +963,7 @@ function Get-InstanceState($v, $inst) {
|
|||
return @{ text = 'starting'; color = 'Yellow' }
|
||||
}
|
||||
if ($inst.lastNoNodeAt -and ($now - $inst.lastNoNodeAt).TotalSeconds -lt 30) { return @{ text = 'no node'; color = 'Yellow' } }
|
||||
if ($inst.lastFaultAt -and -not $inst.status -and ($now - $inst.lastFaultAt).TotalSeconds -lt 180) { return @{ text = 'worker fault'; color = 'Red' } }
|
||||
if ($inst.lastErrorAt -and ($now - $inst.lastErrorAt).TotalSeconds -lt 60) { return @{ text = 'errors'; color = 'Red' } }
|
||||
if (-not $inst.status) { return @{ text = 'warming up'; color = 'Yellow' } }
|
||||
$age = ($now - $inst.status.at).TotalSeconds
|
||||
|
|
@ -996,7 +1002,7 @@ function Print-Status {
|
|||
if ($inst.status) {
|
||||
$total += $inst.status.wall
|
||||
$old = [int]((Get-Date) - $inst.status.at).TotalSeconds
|
||||
Log ("{0}: accepted {1}, {2:N2} MH/s (report {3} s old), template age {4:N2}s, rejected {5}, restarts {6}, {7}" -f $inst.label, $inst.accepted, $inst.status.wall, $old, $inst.status.age, $inst.rejected, $inst.restarts, $st.text)
|
||||
Log ("{0}: accepted {1}, {2:N2} MH/s (report {3} s old), template age {4:N2}s, rejected {5}, restarts {6}, worker faults {8}, {7}" -f $inst.label, $inst.accepted, $inst.status.wall, $old, $inst.status.age, $inst.rejected, $inst.restarts, $st.text, $inst.faults)
|
||||
} else {
|
||||
$since = if ($inst.readyAt) { "ready for $([int]((Get-Date) - $inst.readyAt).TotalSeconds) s, first job still running" } else { 'worker not ready yet' }
|
||||
Log ("{0}: accepted {1}, no STATUS line yet ({2}), restarts {3}, {4}" -f $inst.label, $inst.accepted, $since, $inst.restarts, $st.text)
|
||||
|
|
@ -1185,7 +1191,7 @@ function Build-Frame {
|
|||
foreach ($inst in $v.instances) {
|
||||
$st = Get-InstanceState $v $inst
|
||||
$m = switch ($st.text) { 'mining' { $mark } 'warming up' { $mark } 'starting' { $dots.PadRight(3) } 'rebuilding' { $dots.PadRight(3) } 'restarting' { $dots.PadRight(3) } 'stopped' { 'x' } default { ' ' } }
|
||||
$rate = if ($inst.status) { (Format-Rate $inst.status.wall) + ' MH/s' } else { '' }
|
||||
$rate = if ($inst.status) { (Format-Rate $inst.status.wall) + ' MH/s' } elseif ($st.text -eq 'worker fault') { 'restarting' } else { '' }
|
||||
Add-Seg $L (' {0,2} ' -f $inst.index) 'DarkGray'
|
||||
Add-Seg $L ('{0,-10}' -f $st.text) $st.color
|
||||
Add-Seg $L (' ' + ('{0,-3}' -f $m)) 'DarkGray'
|
||||
|
|
|
|||
|
|
@ -544,10 +544,17 @@ static size_t kernelMaxLocal(const Device* dv, cl_kernel k, const DeviceInfo* di
|
|||
return wg < want ? wg : want;
|
||||
}
|
||||
|
||||
/* Object accounting for the serve loop's stats line (4 October 2026, after the gfx1036 fault): every event and buffer
|
||||
* the serve path creates or releases is counted here, so a leak shows as a growing "live" count long before a runtime
|
||||
* limit is hit. The bench path does not count. */
|
||||
static unsigned long gEvCreated = 0, gEvReleased = 0, gMemCreated = 0, gMemReleased = 0;
|
||||
static void countRelease(cl_event ev) { clReleaseEvent(ev); ++gEvReleased; }
|
||||
|
||||
static cl_event launch1D(const Device* dv, cl_kernel k, size_t global, size_t local) {
|
||||
cl_event ev = NULL;
|
||||
size_t g = ((global + local - 1) / local) * local;
|
||||
CL_CHECK(clEnqueueNDRangeKernel(dv->q, k, 1, NULL, &g, &local, 0, NULL, &ev));
|
||||
++gEvCreated;
|
||||
return ev;
|
||||
}
|
||||
|
||||
|
|
@ -601,7 +608,7 @@ static int setupCache(Device* dv, const DeviceInfo* di) {
|
|||
ev = launch1D(dv, dv->kCacheFill, nSeg, local);
|
||||
CL_CHECK(clFinish(dv->q));
|
||||
if (pass == 0) gCacheFillFirstMs = eventMs(ev); else gCacheFillSecondMs = eventMs(ev);
|
||||
clReleaseEvent(ev);
|
||||
countRelease(ev);
|
||||
}
|
||||
printf("cache fill (device): %.2f ms first, %.2f ms second (%u chains x %u ChaCha blocks, %u MiB, work-group %llu)\n",
|
||||
gCacheFillFirstMs, gCacheFillSecondMs, (unsigned)IGNEUM_CACHE_SEGMENTS, 1u << IGNEUM_CACHE_SEGMENT_LOG2_LINES,
|
||||
|
|
@ -697,7 +704,7 @@ static SizeResult runSize(Device* dv, const DeviceInfo* di, const Options* o, in
|
|||
#endif
|
||||
CL_CHECK(clFinish(dv->q));
|
||||
if (pass == 0) r.fillFirstMs = eventMs(ev); else r.fillSecondMs = eventMs(ev);
|
||||
clReleaseEvent(ev);
|
||||
countRelease(ev);
|
||||
}
|
||||
#if IGNEUM_DATASET_MODE == 1
|
||||
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, device time)\n",
|
||||
|
|
@ -779,7 +786,7 @@ static SizeResult runSize(Device* dv, const DeviceInfo* di, const Options* o, in
|
|||
ev = launchHash(dv, dDs, dOut, IGNEUM_VEC_BASE[w], mask, 32u, 32u);
|
||||
}
|
||||
CL_CHECK(clFinish(dv->q));
|
||||
clReleaseEvent(ev);
|
||||
countRelease(ev);
|
||||
CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, sizeof(got), got, 0, NULL, NULL));
|
||||
r.vecPass = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], o->groupWarps == 1 ? "standalone, 1 unit/work-group" : "standalone, first unit of one work-group") && r.vecPass;
|
||||
}
|
||||
|
|
@ -793,7 +800,7 @@ static SizeResult runSize(Device* dv, const DeviceInfo* di, const Options* o, in
|
|||
{
|
||||
cl_event ev = launchHash(dv, dDs, dOut, 0u, mask, nonces, groupSize);
|
||||
CL_CHECK(clFinish(dv->q));
|
||||
clReleaseEvent(ev);
|
||||
countRelease(ev);
|
||||
}
|
||||
w1 = wallMs();
|
||||
printf("warm-up batch: %u hashes in %.2f ms wall\n", nonces, w1 - w0);
|
||||
|
|
@ -934,8 +941,8 @@ typedef struct {
|
|||
|
||||
static void releasePair(ServePair* p) {
|
||||
if (!p) return;
|
||||
if (p->ds) clReleaseMemObject(p->ds);
|
||||
if (p->cache) clReleaseMemObject(p->cache);
|
||||
if (p->ds) { clReleaseMemObject(p->ds); ++gMemReleased; }
|
||||
if (p->cache) { clReleaseMemObject(p->cache); ++gMemReleased; }
|
||||
if (p->kHashBound) clReleaseKernel(p->kHashBound);
|
||||
if (p->kCacheFill) clReleaseKernel(p->kCacheFill);
|
||||
if (p->kBuild) clReleaseKernel(p->kBuild);
|
||||
|
|
@ -998,7 +1005,8 @@ static int pairSelfTest(Device* dv, const DeviceInfo* di, cl_command_queue q, Se
|
|||
for (i = 0; i < pk.nSamples && e == CL_SUCCESS; ++i) { samples[i] = 0; if (pk.sampleIdx[i] < words) e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, (size_t)pk.sampleIdx[i] * 4u, 4, &samples[i], 0, NULL, NULL); }
|
||||
if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: dataset read-back (%s)", clErrName(e)); return 0; }
|
||||
out = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, g * sizeof(uint64_t), NULL, &e);
|
||||
if (e == CL_SUCCESS) init = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, pk.seedw, &e);
|
||||
if (e == CL_SUCCESS) { ++gMemCreated; init = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, pk.seedw, &e); }
|
||||
if (e == CL_SUCCESS) ++gMemCreated;
|
||||
for (w = 0; w < pk.vecWarps && e == CL_SUCCESS; ++w) {
|
||||
cl_uint base = pk.vecBase[w], mask = words - 1u;
|
||||
e = clSetKernelArg(p->kHashBound, 0, sizeof(cl_mem), &p->ds);
|
||||
|
|
@ -1009,8 +1017,8 @@ static int pairSelfTest(Device* dv, const DeviceInfo* di, cl_command_queue q, Se
|
|||
if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kHashBound, 1, NULL, &g, &local, 0, NULL, NULL);
|
||||
if (e == CL_SUCCESS) e = clEnqueueReadBuffer(q, out, CL_TRUE, 0, 32 * sizeof(uint64_t), &vec[w * 32], 0, NULL, NULL);
|
||||
}
|
||||
if (out) clReleaseMemObject(out);
|
||||
if (init) clReleaseMemObject(init);
|
||||
if (out) { clReleaseMemObject(out); ++gMemReleased; }
|
||||
if (init) { clReleaseMemObject(init); ++gMemReleased; }
|
||||
if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: vector warp (%s)", clErrName(e)); return 0; }
|
||||
ok = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples, vec, p->check, sizeof(p->check));
|
||||
p->checked = 1;
|
||||
|
|
@ -1028,6 +1036,7 @@ static int pairBuffers(Device* dv, const DeviceInfo* di, cl_command_queue q, Ser
|
|||
double tb = wallMs();
|
||||
p->cache = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &e);
|
||||
if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer cache (%s)", clErrName(e)); return 0; }
|
||||
++gMemCreated;
|
||||
local = kernelMaxLocal(dv, p->kCacheFill, di, 256);
|
||||
g = ((nSeg + local - 1) / local) * local;
|
||||
e = clSetKernelArg(p->kCacheFill, 0, sizeof(cl_mem), &p->cache);
|
||||
|
|
@ -1039,6 +1048,7 @@ static int pairBuffers(Device* dv, const DeviceInfo* di, cl_command_queue q, Ser
|
|||
tb = wallMs();
|
||||
p->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &e);
|
||||
if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer dataset (%s)", clErrName(e)); return 0; }
|
||||
++gMemCreated;
|
||||
local = kernelMaxLocal(dv, p->kBuild, di, 256);
|
||||
g = ((nItems + local - 1) / local) * local;
|
||||
e = clSetKernelArg(p->kBuild, 0, sizeof(cl_mem), &p->ds);
|
||||
|
|
@ -1147,11 +1157,12 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
|
|||
CL_CHECK(clSetKernelArg(cur->kBuild, 0, sizeof(cl_mem), &cur->ds));
|
||||
CL_CHECK(clSetKernelArg(cur->kBuild, 1, sizeof(cl_mem), &cur->cache));
|
||||
CL_CHECK(clSetKernelArg(cur->kBuild, 2, sizeof(cl_uint), &nItems));
|
||||
clReleaseEvent(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256)));
|
||||
countRelease(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256)));
|
||||
CL_CHECK(clFinish(dv->q));
|
||||
}
|
||||
dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)batch * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out");
|
||||
dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY, 32, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer init words");
|
||||
gMemCreated += 2;
|
||||
hOut = (uint64_t*)malloc((size_t)batch * sizeof(uint64_t));
|
||||
strncpy(devName, di->name, 255); devName[255] = 0;
|
||||
for (k = 0; devName[k]; ++k) if (devName[k] == ' ') devName[k] = '_';
|
||||
|
|
@ -1159,6 +1170,18 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
|
|||
gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, log2u32(words), batch, dv->exchange, o->noPrepare ? 0 : 1, gGeneric ? "prebuilt-generic" : "compiled-in");
|
||||
fflush(stdout);
|
||||
|
||||
/* Fault detection (4 October 2026, after the gfx1036 run that completed 2,000 jobs a second with no hash after
|
||||
* 600 s): any OpenCL error in the job path is fatal (exit 3, the miner restarts the worker), the dispatch event
|
||||
* must report CL_COMPLETE, a chunk that runs 20x faster per nonce than the running mean is a fault, and two
|
||||
* consecutive chunks with the same output signature (different nonces or init words give different outputs) mean
|
||||
* the kernel did not run. A stats line every 200 jobs shows the live event and buffer counts. */
|
||||
long gFaultTestChunk = getenv("IGNEUM_FAULT_TEST") ? atol(getenv("IGNEUM_FAULT_TEST")) : -1;
|
||||
double meanNsPerNonce = 0; /* running mean of wall ns per nonce over the chunks so far */
|
||||
unsigned long chunksSeen = 0, jobsSeen = 0;
|
||||
uint64_t prevSig = 0; int havePrevSig = 0;
|
||||
double jobMsSum = 0;
|
||||
#define SERVE_FATAL(jobIdStr, fmt, ...) do { printf("error %s worker fault: " fmt "; exiting 3 so the miner restarts the worker\n", jobIdStr, __VA_ARGS__); fflush(stdout); exit(3); } while (0)
|
||||
#define SERVE_CHECK(jobIdStr, call) do { cl_int e_ = (call); if (e_ != CL_SUCCESS) SERVE_FATAL(jobIdStr, "%s failed with %s (%d)", #call, clErrName(e_), (int)e_); } while (0)
|
||||
while (fgets(line, sizeof(line), stdin)) {
|
||||
char* f[9];
|
||||
int nf = 0;
|
||||
|
|
@ -1252,20 +1275,53 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
|
|||
memcpy(b + 13, prehash, 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);
|
||||
CL_CHECK(clEnqueueWriteBuffer(dv->q, dInit, CL_TRUE, 0, 32, iw, 0, NULL, NULL));
|
||||
{
|
||||
double c0 = wallMs(), cms;
|
||||
cl_int status = 0;
|
||||
size_t g = ((chunk + groupSize - 1) / groupSize) * groupSize;
|
||||
uint64_t sig;
|
||||
SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dInit, CL_TRUE, 0, 32, iw, 0, NULL, NULL));
|
||||
baseNonce = lo;
|
||||
CL_CHECK(clSetKernelArg(cur->kHashBound, 0, sizeof(cl_mem), &cur->ds));
|
||||
CL_CHECK(clSetKernelArg(cur->kHashBound, 1, sizeof(cl_mem), &dOut));
|
||||
CL_CHECK(clSetKernelArg(cur->kHashBound, 2, sizeof(cl_uint), &baseNonce));
|
||||
CL_CHECK(clSetKernelArg(cur->kHashBound, 3, sizeof(cl_uint), &maskArg));
|
||||
CL_CHECK(clSetKernelArg(cur->kHashBound, 4, sizeof(cl_mem), &dInit));
|
||||
ev = launch1D(dv, cur->kHashBound, chunk, groupSize);
|
||||
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 0, sizeof(cl_mem), &cur->ds));
|
||||
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 1, sizeof(cl_mem), &dOut));
|
||||
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 2, sizeof(cl_uint), &baseNonce));
|
||||
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 3, sizeof(cl_uint), &maskArg));
|
||||
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 4, sizeof(cl_mem), &dInit));
|
||||
ev = NULL;
|
||||
if (gFaultTestChunk >= 0 && (long)chunksSeen >= gFaultTestChunk) {
|
||||
/* IGNEUM_FAULT_TEST=N (test only): from chunk N on, behave like a runtime that answers every call with
|
||||
* success and runs nothing, so the detectors below are exercised on a healthy device. */
|
||||
} else {
|
||||
SERVE_CHECK(jobId, clEnqueueNDRangeKernel(dv->q, cur->kHashBound, 1, NULL, &g, &groupSize, 0, NULL, &ev));
|
||||
++gEvCreated;
|
||||
/* Wait on the dispatch event, not clFinish: the runtime can sleep the thread on an event, where clFinish
|
||||
* on some drivers spins one core for the whole dispatch. */
|
||||
* on some drivers spins one core for the whole dispatch. The event must then say CL_COMPLETE: a runtime
|
||||
* that has dropped its device answers the wait at once with a negative status. */
|
||||
err = clWaitForEvents(1, &ev);
|
||||
clReleaseEvent(ev);
|
||||
if (err != CL_SUCCESS) { printf("error %s dispatch failed: %s\n", jobId, clErrName(err)); fflush(stdout); failed = 1; break; }
|
||||
CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)chunk * sizeof(uint64_t), hOut, 0, NULL, NULL));
|
||||
if (err != CL_SUCCESS) { countRelease(ev); SERVE_FATAL(jobId, "clWaitForEvents on the dispatch returned %s (%d)", clErrName(err), (int)err); }
|
||||
err = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(status), &status, NULL);
|
||||
countRelease(ev);
|
||||
if (err != CL_SUCCESS) SERVE_FATAL(jobId, "clGetEventInfo on the dispatch returned %s (%d)", clErrName(err), (int)err);
|
||||
if (status != CL_COMPLETE) SERVE_FATAL(jobId, "the dispatch event ended with status %d, not CL_COMPLETE (a device reset or a lost context)", (int)status);
|
||||
}
|
||||
SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)chunk * sizeof(uint64_t), hOut, 0, NULL, NULL));
|
||||
cms = wallMs() - c0;
|
||||
/* Plausibility: wall time per nonce against the running mean (the first chunk sets it; a chunk is a full
|
||||
* batch except at the 32-bit boundary, so per nonce is the comparable unit). 20x faster = the kernel did not run. */
|
||||
{
|
||||
double ns = cms * 1e6 / (double)chunk;
|
||||
if (chunksSeen >= 4 && meanNsPerNonce > 0 && ns * 20.0 < meanNsPerNonce)
|
||||
SERVE_FATAL(jobId, "%u nonces reported complete in %.3f ms, %.0fx faster than the running mean of %.2f ms per million (the runtime is not running the kernel)", chunk, cms, meanNsPerNonce / ns, meanNsPerNonce / 1e3);
|
||||
meanNsPerNonce = chunksSeen == 0 ? ns : meanNsPerNonce + (ns - meanNsPerNonce) / (double)(chunksSeen + 1 < 64 ? chunksSeen + 1 : 64);
|
||||
++chunksSeen;
|
||||
}
|
||||
/* Stale output: the first, middle and last words plus an FNV of the first 32. Two chunks never share it. */
|
||||
sig = fnv1a64(hOut, 32 * sizeof(uint64_t)) ^ hOut[chunk / 2] ^ hOut[chunk - 1];
|
||||
if (chunk >= 64) {
|
||||
if (havePrevSig && sig == prevSig) SERVE_FATAL(jobId, "the output buffer is unchanged since the previous dispatch (signature %016llx): the kernel did not run", (unsigned long long)sig);
|
||||
prevSig = sig; havePrevSig = 1;
|
||||
}
|
||||
}
|
||||
for (i = 0; i < chunk; ++i) if (hOut[i] <= target) {
|
||||
unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + i);
|
||||
printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hOut[i]);
|
||||
|
|
@ -1277,12 +1333,20 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
|
|||
}
|
||||
if (failed) continue;
|
||||
printf("done %s %llu %.2f\n", jobId, (unsigned long long)hashes, wallMs() - t0);
|
||||
jobMsSum += wallMs() - t0;
|
||||
++jobsSeen;
|
||||
if (jobsSeen % 200 == 0) {
|
||||
printf("info stats jobs %lu chunks %lu mean job %.1f ms mean %.2f ms per million nonces; events created %lu released %lu live %lu; buffers created %lu released %lu live %lu\n",
|
||||
jobsSeen, chunksSeen, jobMsSum / (double)jobsSeen, meanNsPerNonce / 1e3, gEvCreated, gEvReleased, gEvCreated - gEvReleased, gMemCreated, gMemReleased, gMemCreated - gMemReleased);
|
||||
if (gEvCreated - gEvReleased > 16 || gMemCreated - gMemReleased > 12) SERVE_FATAL(jobId, "object leak: %lu events and %lu buffers live after %lu jobs", gEvCreated - gEvReleased, gMemCreated - gMemReleased, jobsSeen);
|
||||
}
|
||||
if (switched && old) { releasePair(old); old = NULL; printf("info dropped the previous pair (its program, cache and dataset)\n"); }
|
||||
fflush(stdout);
|
||||
}
|
||||
free(hOut);
|
||||
clReleaseMemObject(dInit);
|
||||
clReleaseMemObject(dOut);
|
||||
gMemReleased += 2;
|
||||
if (old) releasePair(old);
|
||||
if (prepared) releasePair(prepared);
|
||||
releasePair(cur);
|
||||
|
|
|
|||
Loading…
Reference in a new issue