1// Search coordinator: per-dispatch random chunk start (CSPRNG), CPU base-pubkey 2// computation, dispatch loop with TDR-aware sizing, device-loss recovery, 3// localStorage checkpoint of the calibrated dispatch size. 4 5import { getDevice, initWebGPU } from "./webgpu.js"; 6import { searchWGSL, BATCH_HALF, BATCH_KEYS, KG_ENTRIES } from "./shaders/search.js"; 7import { scalarMul, kgTablePoints, bigIntToLimbsLE } from "./reference/secp.js"; 8import { privToPkh, privToAddress } from "./reference/address.js"; 9import { decodeP2PKH } from "./reference/base58.js"; 10 11const WORKGROUP_SIZE = 64; // matches @workgroup_size in search.js 12const MAX_HITS = 16; 13// Two affine batches per thread by default (BATCH_KEYS in search.js is 257 = 2 * 128 + 1). The 14// per-thread start point (about 12 mixed adds plus one Fermat inverse for the affine conversion) 15// is paid once per dispatch per thread, so more keys per thread amortise it: one batch spends 16// about 3 field ops per key on it, two batches about 1.5. The literal is kept so 17// scripts/test-offset-table.mjs can scrape it; the check below pins it to a whole number of 18// batches. A thread can be handed fewer keys (small dispatches, watchdog recovery, the smoke 19// tests); the kernel then spends a batch's arithmetic on fewer hashes. 20const KEYS_PER_THREAD = 514; 21if (KEYS_PER_THREAD % BATCH_KEYS !== 0) console.error("KEYS_PER_THREAD must be a multiple of BATCH_KEYS from search.js"); 22const OFFSET_TABLE_BITS = 22; // matches OFFSET_TABLE_BITS in search.js 23const HEADER_BYTES = 32 + 32 + 20 + 4; // base_x, base_y, target_pkh, keys_per_thread 24const TABLE_BYTES = OFFSET_TABLE_BITS * 64; // 22 entries à (X + Y) = 1408 bytes 25const KG_TABLE_BYTES = KG_ENTRIES * 64; // 129 entries à (X + Y) = 8256 bytes, after the offset table 26const CHUNK_FLOOR = BigInt(BATCH_KEYS + 1); // smallest chunk start the kernel's dx != 0 argument allows 27const PARAMS_BYTES = HEADER_BYTES + TABLE_BYTES + KG_TABLE_BYTES; 28const OUTPUT_BYTES = 4 + MAX_HITS * 8; // count + 16 * (thread, offset) 29 30// Rolling window for the reported keys/sec. Long enough to smooth out 31// dispatch-to-dispatch jitter, short enough that the number tracks reality 32// within a few seconds of resuming from a locked screen. 33const RATE_WINDOW_MS = 5000; 34const TARGET_DISPATCH_MS = 250; // safe distance below the OS TDR ~2s threshold 35const MIN_DISPATCH_MS = 50; 36const MAX_DISPATCH_MS = 500; 37const CALIBRATION_KEY = "puzzlecrack.dispatchKeys"; 38 39// Smallest and default per-dispatch sizes. We START at one workgroup's worth so 40// the very first dispatch on an uncalibrated device can't blow the OS GPU 41// watchdog (TDR) before recalibrate() has a chance to run â critical on slower 42// mobile GPUs. recalibrate() then ramps up toward TARGET_DISPATCH_MS. A device 43// with a persisted calibration (desktop) keeps its larger value. 44const MIN_DISPATCH_KEYS = WORKGROUP_SIZE; // 64 (one thread per lane, one key each) 45const DEFAULT_DISPATCH_KEYS = WORKGROUP_SIZE * KEYS_PER_THREAD; // 32896 (one workgroup, two batches per thread) 46const MAX_RECOVERY_ATTEMPTS = 4; // consecutive device losses before giving up 47// Re-acquire schedule after a device loss: one immediate attempt, then these waits between the 48// retries, so five attempts and about 15 s of waiting (plus up to ACQUIRE_BUDGET_MS per attempt 49// if the browser hangs on the request, so the worst case is about 65 s, with Stop live). The waits exist because `requestDevice` on a D3D12 50// adapter whose device was just removed fails with DXGI_ERROR_DEVICE_REMOVED until the browser's 51// GPU process has re-enumerated, which takes on the order of a second or a few, and a single 52// immediate attempt landed inside that window every time (Intel Xe-LP, 38 minutes in, 2026-09-16). 53// Backoff rather than polling, because the driver is already unhappy. The page cannot observe a 54// GPU-process restart, so this is best effort: past Chromium's crash limit no retry succeeds, and 55// the banner then tells the user to quit and reopen the browser. docs/device-loss-recovery.md. 56const RECOVERY_DELAYS_MS = [1000, 2000, 4000, 8000]; 57// A hung requestDevice is not bounded by the backoff, so each attempt is. A timed-out attempt 58// counts as a failed one; the promise itself cannot be cancelled, and if it resolves late 59// initWebGPU's generation check refuses to install its device (and this file destroys it too, as 60// a belt to that brace). 61const ACQUIRE_BUDGET_MS = 10_000; 62 63// Ceiling on the calibration ramp. Raised from 1<<25 on 2026-08-19 because two cards 64// were measured PINNED here and therefore under-reporting: an RTX PRO 6000 at 381 MK/s 65// (88 ms per dispatch) and an RTX 5090 at 343 (98 ms), both wanting to grow toward the 66// 250 ms target and unable to. The tell that the cap was the cause rather than the 67// silicon: per 1000 CUDA cores those two read 15.8 and 15.8, BELOW the uncapped RTX 5050 68// and 5070 at 16.4 and 16.7. Fastest cards, worst per core, which is what fixed
69// per-dispatch overhead looks like. At 1<<26 the PRO 6000 lands near 176 ms. 70// 71// THE BINDING LIMIT IS OFFSET_TABLE_BITS, NOT u32 AND NOT THE WORKGROUP COUNT, and it 72// fails SILENTLY. The shader places thread t at (t * KEYS_PER_THREAD) * G by walking the 73// bits of t across offset_table[0..OFFSET_TABLE_BITS-1], so max t must stay under 74// 2^OFFSET_TABLE_BITS. Past that the high bits are never read, the thread starts from the 75// wrong pubkey and scans the wrong keys, and nothing raises an error. 76// Both ceilings scale with KEYS_PER_THREAD. The rows below were written for 48 keys per 77// thread; at 258 (the affine batch kernel) the table allows 2^22 * 258 = 1.08e9 keys and 78// 65,535 workgroups allow 65,535 * 64 * 258 = 1.08e9, so: 79// 1<<26 -> 4,065 workgroups, max t 260,159 = 18 of 22 bits (v81 to v102) 80// 1<<28 -> 16,257 workgroups, max t 1,040,447 = 20 of 22 bits, two spare <- here 81// 1<<29 -> 32,514 workgroups, max t 2,080,895 = 21 bits 82// 1<<30 -> 65,028 workgroups, max t 4,161,791 = fits both, with 507 workgroups to spare 83// 1<<31 breaks both. Do NOT go past 1<<29 without raising OFFSET_TABLE_BITS in lockstep 84// (here and in search.js; the table costs 64 bytes per bit). partition() asserts the table 85// bound so a future change cannot cross it silently. The workgroup limit binds one 86// workgroup before the table does, but maxComputeWorkgroupsPerDimension is only a spec 87// MINIMUM of 65,535 and devices may report more, so that validation error cannot be relied 88// on to fire first. 89// 90// RAISED TO 1<<28 in v103 (2026-09-21). The v102 kernel made mid-range Blackwell exceed 91// the 1<<26 bind point (268 MK/s at the 250 ms threshold): an RTX 5080, a 5070 Ti and a 92// 5060 Ti all reported dispatch 67,108,864 at 111, 144 and 220 ms, i.e. pinned and 93// under-reporting (606.8, 467.2 and 305.6 MK/s as FLOORS). At 1<<28 the cap binds near 94// 1.07 GK/s. Growth is 1.25x from a dispatch that ran under 250 ms, so no growth step can 95// exceed 312.5 ms. The one exposure the raise does change is a persisted calibration 96// restored onto a slower GPU on the same machine (hybrid graphics falling back to the iGPU): 97// a saved size is always about 250 to 312 ms of the GPU that saved it, so that hazard already 98// applies to every non-pinned device and the /8 shrink on device loss exists for it. Fast 99// cards used to be accidentally protected by being pinned far below the band; they now carry 100// the same exposure as everyone else, and in the worst swap take one more recovery step, 101// still inside MAX_RECOVERY_ATTEMPTS. The restore path is clamped to the cap. 102const MAX_KEYS_PER_DISPATCH = 1 << 28; // 268,435,456 103 104const cryptoRandomBigInt = (lo, hi) => { 105 // Sample uniformly in [lo, hi]. hi-lo+1 must fit in some byte length. 106 const range = hi - lo + 1n; 107 if (range <= 0n) return lo; 108 const bits = range.toString(2).length; 109 const bytes = Math.ceil(bits / 8); 110 const buf = new Uint8Array(bytes); 111 while (true) { 112 crypto.getRandomValues(buf); 113 let n = 0n; 114 for (const b of buf) n = (n << 8n) | BigInt(b); 115 // Trim to bit length, reject samples ⥠range to keep the distribution uniform. 116 n &= (1n << BigInt(bits)) - 1n; 117 if (n < range) return lo + n; 118 } 119}; 120 121const packU256LE = (n) => { 122 const limbs = bigIntToLimbsLE(n); 123 const out = new Uint8Array(32); 124 for (let i = 0; i < 8; i++) { 125 out[i*4] = limbs[i] & 0xFF; 126 out[i*4 + 1] = (limbs[i] >>> 8) & 0xFF; 127 out[i*4 + 2] = (limbs[i] >>> 16) & 0xFF; 128 out[i*4 + 3] = (limbs[i] >>> 24) & 0xFF; 129 } 130 return out; 131}; 132 133// kg_table for the shader: (i + 1) * G for i < BATCH_HALF, then BATCH_KEYS * G. Stride-independent. 134const buildKgTable = () => { 135 const bytes = new Uint8Array(KG_TABLE_BYTES); 136 kgTablePoints(BATCH_HALF).forEach((point, i) => { 137 bytes.set(packU256LE(point[0]), i * 64); 138 bytes.set(packU256LE(point[1]), i * 64 + 32); 139 }); 140 return bytes; 141}; 142 143export class SearchCoordinator { 144 constructor() { 145 this.device = null; 146 this.pipeline = null; 147 this.paramsBuf = null; 148 this.outputBuf = null; 149 this.outputReadBuf = null; 150 this.bindGroup = null; 151 this.running = false;
152 this.totalChecked = 0n; 153 this.startTime = 0; 154 this.lastReportTime = 0; 155 // Rolling window of recent dispatches, for the reported keys/sec. See 156 // recordSample()/currentRate(): the rate must describe the last few seconds, 157 // NOT the average since Start. 158 this.samples = []; 159 // Calibrated keys-per-dispatch (loaded from localStorage if present, else conservative default). 160 // Clamp the restored value to the same ceiling the ramp obeys. The old bound was 161 // `saved < 1e9`, which is 15x the previous cap and 30x this one, so a stored value 162 // could bypass the ramp's limit entirely and go straight to a size that overflows 163 // the offset table (silent wrong keys) or the workgroup limit. Nothing WRITES such a 164 // value today, since both writers are the capped ramp and the shrink path, but a 165 // ceiling that only applies to growth is not a ceiling. 166 const saved = parseInt(localStorage.getItem(CALIBRATION_KEY) ?? "0", 10); 167 this.keysPerDispatch = (saved > 0 && saved <= MAX_KEYS_PER_DISPATCH) 168 ? saved 169 : DEFAULT_DISPATCH_KEYS; 170 this.consecutiveLosses = 0; 171 // Recovery hooks, overridable per instance so the retry loop is unit-testable without a GPU: 172 // the acquisition (initWebGPU), the wait between attempts, and the per-attempt budget. 173 this._acquire = initWebGPU; 174 this._sleep = (ms) => new Promise((r) => setTimeout(r, ms)); 175 this._acquireBudgetMs = ACQUIRE_BUDGET_MS; 176 // The error that ended the last failed recovery, and which half failed ("acquire": the browser 177 // would not hand back a GPU; "setup": it did, and the search could not be rebuilt on it). Two 178 // different banners, because they need opposite advice. Cleared on a successful recovery. 179 this.lastRecoveryError = null; 180 this.lastRecoveryFailure = null;
181 // Offset tables, keyed by stride. See ensureOffsetTable(): the table encodes the 182 // per-thread STRIDE, and partition() can change that stride between dispatches, so 183 // one table built at setup() is not enough. Cached because the recovery ramp walks 184 // back up through a handful of strides and each build costs 22 CPU scalarMuls. 185 this.tableCache = new Map(); 186 // Stride of the table currently sitting in paramsBuf, or null when unknown. MUST be 187 // reset whenever paramsBuf is recreated (setup() after a device loss makes a NEW 188 // buffer, and a cached "already uploaded" would then skip writing the table at all). 189 this.uploadedStride = null; 190 } 191 192 async setup() { 193 // Acquire the device. If it was never created or was lost â e.g. the 194 // off-path diagnostics crashed it during the capability check on a fragile 195 // driver â request a fresh one so the search can still run. 196 let device; 197 try { 198 device = getDevice(); 199 } catch { 200 ({ device } = await initWebGPU()); 201 } 202 this.device = device; 203 204 const module = device.createShaderModule({ code: searchWGSL }); 205 const info = await module.getCompilationInfo(); 206 for (const m of info.messages) { 207 if (m.type === "error") throw new Error(`search shader compile error: ${m.message} at line ${m.lineNum}`); 208 } 209 this.pipeline = await device.createComputePipelineAsync({ 210 layout: "auto", 211 compute: { module, entryPoint: "search" }, 212 }); 213 214 this.paramsBuf = device.createBuffer({ 215 size: PARAMS_BYTES, 216 usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, 217 }); 218 219 // paramsBuf is brand new, so whatever table used to be uploaded is gone. Forcing a 220 // re-upload on the next dispatch is the whole reason this is tracked: setup() runs 221 // again after a device loss, and a stale "already uploaded" would leave the table 222 // region zeroed, which the shader would read as the point at infinity. 223 this.uploadedStride = null; 224 // The kG table does not depend on the stride, so it is written once per paramsBuf. 225 device.queue.writeBuffer(this.paramsBuf, HEADER_BYTES + TABLE_BYTES, buildKgTable()); 226 this.outputBuf = device.createBuffer({ 227 size: OUTPUT_BYTES, 228 usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, 229 }); 230 this.outputReadBuf = device.createBuffer({ 231 size: OUTPUT_BYTES, 232 usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, 233 }); 234 this.bindGroup = device.createBindGroup({ 235 layout: this.pipeline.getBindGroupLayout(0), 236 entries: [ 237 { binding: 0, resource: { buffer: this.paramsBuf } }, 238 { binding: 1, resource: { buffer: this.outputBuf } }, 239 ], 240 }); 241 242 // Note: we do NOT flip `this.running` here. Device loss is handled in the 243 // dispatch loop's catch (which can shrink + recover), and flipping running 244 // from this async callback would race with an in-progress recovery. 245 device.lost.then((info) => { 246 console.warn("device lost during search:", info.reason, info.message); 247 }); 248 } 249 250 // Race an acquisition against the per-attempt budget. If the acquisition resolves AFTER the 251 // budget fired, whatever device it produced is destroyed. initWebGPU's own generation check is 252 // what actually keeps a late device from being installed (a superseded call destroys its device 253 // and throws); this destroy is the belt to that brace, for the case where no later call has 254 // started yet and the late device DID get installed: destroying it fires `lost`, and the 255 // identity-guarded handler nulls the module handle only if it still points at this device. 256 _acquireWithinBudget() { 257 return new Promise((resolve, reject) => { 258 let timedOut = false; 259 const t = setTimeout(() => { 260 timedOut = true; 261 reject(new Error(`device re-acquire timed out after ${this._acquireBudgetMs} ms`)); 262 }, this._acquireBudgetMs); 263 Promise.resolve().then(() => this._acquire()).then( 264 (got) => { 265 if (timedOut) { try { got?.device?.destroy(); } catch { /* already gone */ } return; } 266 clearTimeout(t); 267 resolve(got); 268 }, 269 (e) => { clearTimeout(t); if (!timedOut) reject(e); }, 270 ); 271 }); 272 } 273 274 // Re-acquire the WebGPU device and rebuild the pipeline/buffers after a loss. Returns true if 275 // the device is usable again. 276 // 277 // ONLY THE ACQUISITION IS RETRIED. It is the transient part: a just-removed D3D12 adapter 278 // refuses requestDevice until the GPU process re-enumerates. setup() is deterministic (shader 279 // compile, buffer creation), so a failure there will not get better with waiting, and retrying 280 // it five times would misreport a compile error as an unrecoverable device loss. 281 // 282 // `running` is checked at the top of EVERY iteration, so Stop pressed during an 8 s wait ends 283 // the loop before the next attempt, and Stop pressed before recovery starts skips it entirely. 284 // The caller tells the two apart by reading `running` before it flips it. 285 // 286 // A fallback (software) adapter is refused. After enough GPU-process crashes Chromium can stop 287 // offering the hardware adapter, and a search that "recovers" onto CPU emulation runs honestly 288 // at thousands of keys per second with nothing to flag it, which is worse than stopping. 289 // 290 // `onAttempt({ attempt, of })` fires before each attempt so the banner can show the retry 291 // happening; only a Windows reporter can prove this path, so the number has to be visible. 292 // Stop pressed DURING setup() takes effect the moment setup returns, via the dispatch loop's 293 // `while (this.running)`; a shader compile cannot be cancelled, so there is nothing earlier to do. 294 async recover(onAttempt) { 295 const attempts = RECOVERY_DELAYS_MS.length + 1; 296 for (let i = 0; i < attempts; i++) { 297 if (!this.running) return false;
298 onAttempt?.({ attempt: i + 1, of: attempts }); 299 let acquired; 300 try { 301 acquired = await this._acquireWithinBudget(); 302 if (acquired?.adapter?.isFallbackAdapter) { 303 try { acquired.device?.destroy(); } catch { /* nothing to release */ } 304 throw new Error("re-acquired a software (fallback) adapter; refusing to search on it"); 305 } 306 } catch (e) { 307 this.lastRecoveryError = e; 308 this.lastRecoveryFailure = "acquire"; 309 console.warn(`device re-acquire attempt ${i + 1}/${attempts} failed:`, e); 310 if (e?.initPhase === "gpu-missing") return false; // navigator.gpu is gone: nothing to wait for 311 if (i === attempts - 1) return false; 312 await this._sleep(RECOVERY_DELAYS_MS[i]); 313 continue; 314 } 315 try { 316 await this.setup(); // rebuilds pipeline, buffers, offset table, bind group on the new device 317 this.lastRecoveryError = null; 318 this.lastRecoveryFailure = null; 319 return true; 320 } catch (e) { 321 this.lastRecoveryError = e; 322 this.lastRecoveryFailure = "setup"; 323 console.warn("setup failed on the re-acquired device:", e); 324 // The device we just acquired, by reference: getDevice() could name a different one. 325 try { acquired?.device?.destroy(); } catch { /* the device may already be gone */ } 326 return false; 327 } 328 } 329 return false; 330 } 331 332 // Record a finished dispatch for the rolling rate window. `t1` is when it 333 // ended, `elapsedMs` how long it took, so a sample owns a real time span 334 // rather than an instant. Samples that end before the window are dropped. 335 recordSample(t1, elapsedMs, keys) { 336 this.samples.push({ t0: t1 - elapsedMs, t1, keys }); 337 const cutoff = t1 - RATE_WINDOW_MS; 338 while (this.samples.length > 1 && this.samples[0].t1 < cutoff) this.samples.shift(); 339 } 340 341 // Keys/sec over the last RATE_WINDOW_MS of ACTUAL dispatch time. 342 // 343 // Deliberately divides by the summed dispatch durations, not by wall clock. 344 // Idle gaps (a locked screen or a backgrounded tab stops requestAnimationFrame, 345 // so the loop simply stalls) are therefore excluded instead of being averaged 346 // in. The previous implementation divided lifetime keys by wall clock since 347 // Start, so hours of a locked screen crushed the figure and it needed roughly 348 // the length of the idle period to recover, which read as the GPU slowly 349 // warming up when it was actually at full speed the whole time. 350 // 351 // Returns null until there is a sample, so the UI can keep showing a placeholder. 352 currentRate() { 353 if (this.samples.length === 0) return null; 354 let keys = 0, ms = 0; 355 for (const s of this.samples) { keys += s.keys; ms += s.t1 - s.t0; } 356 return ms > 0 ? (keys * 1000) / ms : null; 357 } 358 359 // offset_table[b] = (2^b * stride) * G in affine, packed for upload. 360 // 361 // The table encodes the per-thread STRIDE: the shader places thread t at 362 // base + (t * stride) * G by summing the entries for the set bits of t. So the stride 363 // baked in here MUST equal the keys_per_thread in the header and the stride the CPU 364 // reconstructs a hit with. Those three drifting apart is exactly the bug this method 365 // exists to make impossible; before 2026-08-19 the table was built once with 366 // KEYS_PER_THREAD while partition() was free to hand the shader a smaller value. 367 // 368 // Cached because a build is 22 CPU scalarMuls (~100ms) and the recovery ramp walks back 369 // up through several strides, so an uncached build would run repeatedly on the one path 370 // where the device is already struggling. 371 buildOffsetTable(stride) { 372 const cached = this.tableCache.get(stride); 373 if (cached) return cached; 374 const bytes = new Uint8Array(TABLE_BYTES); 375 let multiplier = BigInt(stride); 376 for (let k = 0; k < OFFSET_TABLE_BITS; k++) { 377 const point = scalarMul(multiplier); 378 bytes.set(packU256LE(point[0]), k * 64); 379 bytes.set(packU256LE(point[1]), k * 64 + 32); 380 multiplier *= 2n; 381 } 382 this.tableCache.set(stride, bytes); 383 return bytes; 384 } 385 386 // Make the table in paramsBuf match `stride`, uploading only when it has to change. 387 // Returns the stride now in the buffer, which callers treat as the authority.
388 ensureOffsetTable(stride) { 389 if (this.uploadedStride !== stride) { 390 this.device.queue.writeBuffer(this.paramsBuf, HEADER_BYTES, this.buildOffsetTable(stride)); 391 this.uploadedStride = stride; 392 } 393 return this.uploadedStride; 394 } 395 396 // Run one dispatch. Returns { hits: [...], elapsedMs, stride }. 397 async runDispatch({ basePriv, basePubAffine, targetPkh, dispatchWorkgroups, keysPerThread }) { 398 const device = this.device; 399 400 // Align the table with this dispatch's stride. Deliberately before t0: a cold build 401 // is ~100ms of CPU work, and inside the timed section it would read as a slow 402 // dispatch, making recalibrate() shrink. On the recovery path that is a feedback loop 403 // straight into the 64-key floor. 404 const stride = this.ensureOffsetTable(keysPerThread); 405 406 // Pack the per-dispatch header. 407 const header = new Uint8Array(HEADER_BYTES); 408 header.set(packU256LE(basePubAffine[0]), 0); // base_x 409 header.set(packU256LE(basePubAffine[1]), 32); // base_y 410 const dv = new DataView(header.buffer); 411 for (let i = 0; i < 5; i++) { 412 const w = (targetPkh[i*4] | (targetPkh[i*4+1] << 8) | (targetPkh[i*4+2] << 16) | (targetPkh[i*4+3] << 24)) >>> 0; 413 dv.setUint32(64 + i*4, w, true); 414 } 415 dv.setUint32(84, keysPerThread, true); 416 device.queue.writeBuffer(this.paramsBuf, 0, header); 417 418 // Clear output buffer (zero out the atomic counter and result slots). 419 device.queue.writeBuffer(this.outputBuf, 0, new Uint8Array(OUTPUT_BYTES)); 420 421 const t0 = performance.now(); 422 const cmd = device.createCommandEncoder(); 423 const pass = cmd.beginComputePass(); 424 pass.setPipeline(this.pipeline); 425 pass.setBindGroup(0, this.bindGroup); 426 pass.dispatchWorkgroups(dispatchWorkgroups, 1, 1); 427 pass.end(); 428 cmd.copyBufferToBuffer(this.outputBuf, 0, this.outputReadBuf, 0, OUTPUT_BYTES); 429 device.queue.submit([cmd.finish()]); 430 431 await this.outputReadBuf.mapAsync(GPUMapMode.READ); 432 const out = new Uint8Array(this.outputReadBuf.getMappedRange().slice(0)); 433 this.outputReadBuf.unmap(); 434 const elapsedMs = performance.now() - t0; 435 436 const odv = new DataView(out.buffer); 437 const count = odv.getUint32(0, true); 438 const hits = []; 439 for (let i = 0; i < Math.min(count, MAX_HITS); i++) { 440 hits.push({ 441 threadId: odv.getUint32(4 + i*8, true), 442 keyOffset: odv.getUint32(4 + i*8 + 4, true), 443 }); 444 } 445 return { hits, elapsedMs, stride }; 446 } 447 448 // Adjust keysPerDispatch toward the TARGET_DISPATCH_MS sweet spot. 449 // Halve immediately on overshoot (the next dispatch could be the one that hits TDR). 450 // 451 // Grow while a dispatch reads under TARGET_DISPATCH_MS itself, not under 0.7x of it (v99, 452 // 2026-09-14). The 175 ms threshold parked a smooth engine in [175, 218.75) ms, which was 453 // fine, but Firefox on macOS completes mapAsync on a ~100 ms grid: timestamp queries put the 454 // GPU pass at 121 ms for 172k keys with the page hearing back at 200, and Chrome/Safari hear 455 // within 2 ms of GPU end. The first reading over 175 on that grid is 200, taken the moment 456 // the work crosses one tick, so the ramp parked with the GPU idle ~45% of every dispatch and 457 // the tile showed about half the real rate. Growing to the 300 ms reading cuts the idle share 458 // to ~30%: Firefox on an M1 went from 1.02 to 1.37 MK/s on the app's own meter with only this 459 // line changed. Smooth engines keep their rate and settle in [250, 312.5) ms, still 1.6x 460 // under the 500 ms halving line, which covers the 36% throttling measured on an M4 Max 461 // laptop (312.5 x 1.36 = 425). Do NOT push the threshold past ~290: at 310 the settled band 462 // reaches 387.5 ms and the same laptop would cross 500 and cycle. No engine detection, on 463 // purpose; this is the one loop that must never misbehave. docs/dispatch-ramp-threshold.md. 464 recalibrate(elapsedMs) { 465 let next = this.keysPerDispatch; 466 if (elapsedMs > MAX_DISPATCH_MS) { 467 next = Math.max(WORKGROUP_SIZE, Math.floor(next / 2)); 468 } else if (elapsedMs < MIN_DISPATCH_MS) { 469 next = Math.min(MAX_KEYS_PER_DISPATCH, next * 2); 470 } else if (elapsedMs < TARGET_DISPATCH_MS) { 471 next = Math.min(MAX_KEYS_PER_DISPATCH, Math.floor(next * 1.25)); 472 } 473 if (next !== this.keysPerDispatch) { 474 this.keysPerDispatch = next;
475 localStorage.setItem(CALIBRATION_KEY, String(next)); 476 } 477 } 478 479 // Compute a (workgroups, keysPerThread) split that produces approximately 480 // `desired` keys per dispatch. Constraints: 481 // - keysPerThread >= 1 482 // - workgroups * WORKGROUP_SIZE * keysPerThread >= desired 483 // - keep workgroups * WORKGROUP_SIZE under 2^OFFSET_TABLE_BITS, which is the real 484 // ceiling (see MAX_KEYS_PER_DISPATCH). The old comment here named u32 as the 485 // constraint; u32 is ~64x away and was never the thing that binds. 486 partition(desired) { 487 let keysPerThread = KEYS_PER_THREAD; 488 let workgroups = Math.max(1, Math.ceil(desired / (WORKGROUP_SIZE * keysPerThread))); 489 // The shader reads only OFFSET_TABLE_BITS bits of the thread index; a thread above that 490 // would start from the wrong pubkey and report wrong work as done, with no error raised. 491 if (workgroups * WORKGROUP_SIZE > (1 << OFFSET_TABLE_BITS)) { 492 throw new Error(`dispatch of ${desired} keys needs ${workgroups * WORKGROUP_SIZE} threads, above the 2^${OFFSET_TABLE_BITS} offset-table bound`); 493 } 494 // For a single-workgroup dispatch (the small end, incl. watchdog recovery), 495 // also shrink keys-per-thread so the per-thread wall-time can drop below the 496 // OS watchdog. Costs batch-inversion efficiency, but a completed small 497 // dispatch beats a device-killing large one. 498 if (workgroups === 1) { 499 keysPerThread = Math.max(1, Math.min(KEYS_PER_THREAD, Math.ceil(desired / WORKGROUP_SIZE))); 500 } 501 return { workgroups, keysPerThread }; 502 } 503 504 async start(opts) { 505 if (this.running) return; 506 this.running = true; 507 this.totalChecked = 0n; 508 this.consecutiveLosses = 0; 509 this.samples = []; 510 this.startTime = performance.now(); 511 this.lastReportTime = this.startTime; 512 513 const targetPkh = await decodeP2PKH(opts.targetAddress); 514 const targetPkhArr = Array.from(targetPkh); 515 516 const onHit = async (priv) => { 517 const addr = await privToAddress(priv); 518 const ok = addr === opts.targetAddress; 519 opts.onHit({ priv, addr, verified: ok }); 520 this.running = false; 521 }; 522 523 while (this.running) { 524 const rangeSize = opts.rangeEnd - opts.rangeStart + 1n; 525 const dispatchSize = BigInt(this.keysPerDispatch); 526 let chunkPriv; 527 if (rangeSize <= dispatchSize) { 528 // The entire range fits in one dispatch â sweep from the start. 529 chunkPriv = opts.rangeStart; 530 } else { 531 // Sample chunk_start in [rangeStart - dispatchSize + 1, rangeEnd] so 532 // every key in [rangeStart, rangeEnd] has equal probability of falling 533 // inside any dispatch â including keys near the low edge. Negative 534 // chunk_start values are fine: scalarMul reduces mod N internally, and 535 // hits whose recovered priv lands outside the range are filtered below. 536 const lo = opts.rangeStart - dispatchSize + 1n; 537 const hi = opts.rangeEnd; 538 chunkPriv = cryptoRandomBigInt(lo, hi); 539 } 540 // Floor: scalarMul(0) is the point at infinity, and the affine batch kernel needs every 541 // centre key above BATCH_KEYS so that no kG table entry can coincide with a centre (a zero 542 // dx would corrupt the batch; see the header comment in search.js). Never fires for the 543 // puzzle range, which starts near 2^70; it exists so the kernel's precondition is a 544 // guarantee rather than a probability on small test ranges. 545 if (chunkPriv <= CHUNK_FLOOR) chunkPriv = CHUNK_FLOOR; 546 let basePubAffine; 547 try { 548 basePubAffine = scalarMul(chunkPriv); // CPU JS reference (handles negative & wraps mod N) 549 if (basePubAffine === null) continue; 550 } catch (e) { 551 console.warn("CPU scalarMul failed for chunk:", e); 552 continue; 553 } 554 555 const { workgroups, keysPerThread } = this.partition(this.keysPerDispatch); 556 const totalKeysThisDispatch = workgroups * WORKGROUP_SIZE * keysPerThread; 557 558 let hits, elapsedMs, stride; 559 try { 560 ({ hits, elapsedMs, stride } = await this.runDispatch({ 561 basePriv: chunkPriv, 562 basePubAffine, 563 targetPkh: targetPkhArr, 564 dispatchWorkgroups: workgroups, 565 keysPerThread, 566 })); 567 this.consecutiveLosses = 0; // a completed dispatch clears the recovery counter 568 } catch (e) { 569 // Almost always the OS GPU watchdog (TDR) killing an over-long dispatch. 570 // Shrink hard and PERSIST so we converge instead of re-crashing on the 571 // same size, then try to re-acquire the device and continue â no manual 572 // reload. Give up only after repeated losses (a GPU that can't run even 573 // a minimal dispatch within its watchdog window). 574 console.warn("dispatch failed (likely device lost / TDR):", e); 575 this.consecutiveLosses += 1; 576 const shrunk = Math.max(MIN_DISPATCH_KEYS, Math.floor(this.keysPerDispatch / 8)); 577 this.keysPerDispatch = shrunk;
578 localStorage.setItem(CALIBRATION_KEY, String(shrunk)); 579 580 if (this.consecutiveLosses > MAX_RECOVERY_ATTEMPTS) { 581 this.running = false; 582 opts.onDeviceLost?.({ recovered: false, reason: "watchdog", of: MAX_RECOVERY_ATTEMPTS, keysPerDispatch: shrunk }); 583 return; 584 } 585 opts.onDeviceLost?.({ 586 recovering: true, keysPerDispatch: shrunk, 587 attempt: this.consecutiveLosses, of: MAX_RECOVERY_ATTEMPTS, 588 }); 589 const ok = await this.recover((a) => opts.onDeviceLost?.({ 590 recovering: true, reacquiring: true, keysPerDispatch: shrunk, 591 attempt: a.attempt, of: a.of, 592 })); 593 if (!ok) { 594 // recover() also returns false when Stop was pressed during a wait. That is the user's 595 // choice, not a failure, and gets no failure banner. 596 const stopped = !this.running; 597 this.running = false; 598 if (!stopped) { 599 opts.onDeviceLost?.({ 600 recovered: false, keysPerDispatch: shrunk, 601 // "acquire": no GPU came back, quit and reopen the browser. "setup": one did, and 602 // the search could not be rebuilt on it, which is a compile or allocation failure 603 // and needs different words. 604 reason: this.lastRecoveryFailure === "setup" ? "setup" : "acquire", 605 error: this.lastRecoveryError?.message || String(this.lastRecoveryError || ""), 606 }); 607 } 608 return; 609 } 610 continue; // retry the loop with the smaller, persisted dispatch size 611 } 612 613 for (const hit of hits) { 614 // `stride`, not `keysPerThread`: it is the stride actually baked into the 615 // uploaded table, which is what the shader stepped by. Equal today because 616 // ensureOffsetTable() makes them equal, and reading the returned value keeps 617 // that true even if the table logic changes later. 618 const foundPriv = chunkPriv + BigInt(hit.threadId * stride + hit.keyOffset); 619 // Drop hits whose privkey falls outside the configured range. These come 620 // from threads scanning the "negative tail" when chunkPriv < rangeStart; 621 // their pubkey collision with the target is astronomically unlikely but 622 // wouldn't be a Puzzle solution anyway. 623 if (foundPriv < opts.rangeStart || foundPriv > opts.rangeEnd) continue; 624 // Independently re-derive the PKH on the CPU before acting on this key. The GPU 625 // reports only (thread_id, key_offset), so a reconstruction bug yields a WRONG 626 // key that looks entirely valid, and the right one is unrecoverable because 627 // nothing else recorded it. One scalarMul on a once-in-the-universe event costs 628 // nothing; being silently wrong here costs the prize. Loud and skipped rather 629 // than claimed, because a key we cannot confirm is worse than no key. 630 const checkPkh = await privToPkh(foundPriv); 631 if (!checkPkh.every((b, k) => b === targetPkhArr[k])) { 632 console.error( 633 "REJECTED a GPU hit: CPU re-derivation disagrees. This is a bug, not bad luck.", 634 { thread: hit.threadId, keyOffset: hit.keyOffset, stride, 635 keysPerThread, dispatchKeys: this.keysPerDispatch, 636 priv: foundPriv.toString(16) }, 637 ); 638 continue; 639 } 640 await onHit(foundPriv); 641 if (!this.running) return; 642 } 643 644 this.totalChecked += BigInt(totalKeysThisDispatch); 645 this.recalibrate(elapsedMs); 646 647 const now = performance.now(); 648 this.recordSample(now, elapsedMs, totalKeysThisDispatch); 649 if (now - this.lastReportTime > 200) { 650 const rate = this.currentRate(); 651 opts.onProgress({ 652 totalChecked: this.totalChecked, 653 rate, 654 dispatchKeys: this.keysPerDispatch, 655 chunkPriv, 656 chunkSpan: BigInt(totalKeysThisDispatch), 657 }); 658 this.lastReportTime = now; 659 } 660 661 // Yield to keep the page responsive. 662 await new Promise((r) => requestAnimationFrame(r)); 663 } 664 } 665 666 stop() { 667 this.running = false;
668 } 669}
Line numbers count LF bytes from the start of the resource, as the search results do. Vendor segments are library code the classifier recognised; they are stored but not indexed. Bytes are shown as Latin1 characters, one per byte.