diff --git a/proto-cuda/emu/cuda_runtime.h b/proto-cuda/emu/cuda_runtime.h index fd4d82e04..995d5c509 100644 --- a/proto-cuda/emu/cuda_runtime.h +++ b/proto-cuda/emu/cuda_runtime.h @@ -61,6 +61,7 @@ template cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessor(int* // Device intrinsics inline unsigned int __umulhi(unsigned int a, unsigned int b) { return (unsigned int)(((uint64_t)a * (uint64_t)b) >> 32); } unsigned int __shfl_xor_sync(unsigned int mask, unsigned int v, int laneMask, int width = 32); +unsigned int __shfl_sync(unsigned int mask, unsigned int v, int srcLane, int width = 32); // wide loads (lever b) broadcast lane 0 // Warp emulation struct EmuWarp { diff --git a/proto-cuda/emu/shim.cpp b/proto-cuda/emu/shim.cpp index f83baf1e2..f646f4902 100644 --- a/proto-cuda/emu/shim.cpp +++ b/proto-cuda/emu/shim.cpp @@ -27,3 +27,13 @@ unsigned int __shfl_xor_sync(unsigned int, unsigned int v, int laneMask, int) { warp_barrier(w); return r; } + +unsigned int __shfl_sync(unsigned int, unsigned int v, int srcLane, int) { + EmuWarp* w = emu_current_warp; + unsigned lane = threadIdx.x & 31u; + w->slot[lane] = v; + warp_barrier(w); + unsigned int r = w->slot[(unsigned)srcLane & 31u]; + warp_barrier(w); + return r; +} diff --git a/proto-cuda/host.cu b/proto-cuda/host.cu index 903a9a5f1..c501e5555 100644 --- a/proto-cuda/host.cu +++ b/proto-cuda/host.cu @@ -287,15 +287,36 @@ static SizeResult runSize(const Options& o, int mib, uint64_t* dOut, uint32_t no uint32_t idx = (uint32_t)z & mask; uint32_t v = 0; CUDA_CHECK(cudaMemcpy(&v, dDs + idx, sizeof(v), cudaMemcpyDeviceToHost)); +#if IGNEUM_DATASET_MODE == 1 + uint32_t want = host_ds_word(idx); +#else uint32_t want = host_ds_elem(idx, IGNEUM_DAY0, IGNEUM_DAY1); +#endif if (v != want) { - if (badRnd == 0) std::printf(" dataset[%u] = 0x%08x, host formula 0x%08x\n", idx, v, want); + if (badRnd == 0) std::printf(" dataset[%u] = 0x%08x, host %s 0x%08x\n", idx, v, IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", want); ++badRnd; } } - r.dsPass = (badHead == 0 && lastOk && badRnd == 0); - std::printf("dataset self-test: %s (head 16 vs Mac %s, element [MASK] vs Mac %s, 64 random points vs host formula %s)\n", - r.dsPass ? "PASS" : "FAIL", badHead == 0 ? "PASS" : "FAIL", lastText, badRnd == 0 ? "PASS" : "FAIL"); + // The Mac's sampled words: every sample whose index lies inside this dataset size (items are the same + // at every size, the smaller dataset is a prefix of the larger one). + int badSample = 0, nSample = 0; +#ifdef IGNEUM_DS_SAMPLES + for (int k = 0; k < IGNEUM_DS_SAMPLES; ++k) { + if (IGNEUM_DS_SAMPLE_INDEX[k] > mask) continue; + ++nSample; + uint32_t v = 0; + CUDA_CHECK(cudaMemcpy(&v, dDs + IGNEUM_DS_SAMPLE_INDEX[k], sizeof(v), cudaMemcpyDeviceToHost)); + if (v != IGNEUM_DS_SAMPLE_VALUE[k]) { + if (badSample == 0) std::printf(" dataset[%u] = 0x%08x, Mac 0x%08x\n", IGNEUM_DS_SAMPLE_INDEX[k], v, IGNEUM_DS_SAMPLE_VALUE[k]); + ++badSample; + } + } +#endif + r.dsPass = (badHead == 0 && lastOk && badRnd == 0 && badSample == 0); + std::printf("dataset self-test: %s (head 16 vs Mac %s, element [MASK] vs Mac %s, 64 random points vs host %s %s, %d Mac samples %s)\n", + r.dsPass ? "PASS" : "FAIL", badHead == 0 ? "PASS" : "FAIL", lastText, + IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", badRnd == 0 ? "PASS" : "FAIL", + nSample, nSample == 0 ? "none in pack" : (badSample == 0 ? "PASS" : "FAIL")); } // Vectors, standalone: one 32-thread block per base nonce, exactly like the Mac cross-check. @@ -410,6 +431,13 @@ int main(int argc, char** argv) { SEEDW[0], SEEDW[1], SEEDW[2], SEEDW[3], SEEDW[4], SEEDW[5], SEEDW[6], SEEDW[7]); std::printf("day \"%s\" (d0 0x%08x, d1 0x%08x), pack dataset 2^%d words = %d MiB\n", IGNEUM_DAY_STRING, IGNEUM_DAY0, IGNEUM_DAY1, IGNEUM_DATASET_LOG2, packMib()); +#if IGNEUM_DATASET_MODE == 1 + std::printf("dataset construction: memory-hard (256 MiB ChaCha cache, %d dependent cache reads per 64-byte item; proto-metal/MEMHARD.md)\n", IGNEUM_ITEM_ROUNDS); + bool cachePass = setupCache(); +#else + std::printf("dataset construction: closed-form ds_elem (the original prototype dataset, not memory-hard)\n"); + bool cachePass = true; +#endif uint32_t nonces = 1u << o.batchLog2; if (nonces % (32u * (uint32_t)o.blockWarps) != 0u) { @@ -429,9 +457,10 @@ int main(int argc, char** argv) { std::printf("\n=== summary (%s, pack %s, batch 2^%d x %d, %d warp(s)/block, GPU-event time) ===\n", prop.name, IGNEUM_SEED_STRING, o.batchLog2, o.batches, o.blockWarps); - std::printf("| dataset MiB | fill ms (second) | Mhash/s | GB/s useful | random loads/s (G) | loads/hash | dataset self-test | vectors |\n"); + std::printf("| dataset MiB | %s ms (second) | Mhash/s | GB/s useful | random loads/s (G) | loads/hash | dataset self-test | vectors |\n", + IGNEUM_DATASET_MODE == 1 ? "build" : "fill"); std::printf("|---|---|---|---|---|---|---|---|\n"); - bool overall = true, anyVec = false; + bool overall = cachePass, anyVec = false; for (size_t i = 0; i < results.size(); ++i) { const SizeResult& r = results[i]; overall = overall && r.dsPass && (!r.vecChecked || r.vecPass); @@ -442,6 +471,10 @@ int main(int argc, char** argv) { r.dsPass ? "PASS" : "FAIL", r.vecChecked ? (r.vecPass ? "PASS (3 warps, standalone and in batch)" : "FAIL") : "skipped (not pack size)"); } +#if IGNEUM_DATASET_MODE == 1 + std::printf("cache: GPU fill %.2f ms (second), host fill %.1f ms one thread, cache check %s\n", gCacheFillSecondMs, gCacheHostMs, cachePass ? "PASS" : "FAIL"); + CUDA_CHECK(cudaFree(gCache)); +#endif if (!anyVec) std::printf("NOTE: no vectors were checked. Run at %d MiB (the default) to verify against the Mac.\n", packMib()); std::printf("OVERALL: %s\n", overall ? "PASS" : "FAIL"); return overall ? 0 : 1; diff --git a/site/index.html b/site/index.html index 04b6e9ea7..16262ed13 100644 --- a/site/index.html +++ b/site/index.html @@ -163,8 +163,16 @@ pre{margin:0;font-family:'IBM Plex Mono',monospace;font-size:13px;line-height:1. .band .wrap{display:flex;flex-wrap:wrap;align-items:center;justify-content:space-between;gap:24px;padding-block:56px} .band h2{font-weight:900;font-size:clamp(26px,3.5vw,36px)} .band p{font-size:16px;max-width:56ch} -footer .wrap{display:flex;flex-wrap:wrap;justify-content:space-between;align-items:center;gap:20px;padding-block:40px 48px;font-size:14px;color:var(--ash)} -footer .fl{display:flex;flex-wrap:wrap;gap:20px} +footer{border-top:1px solid var(--line)} +footer .wrap{padding-block:48px 32px} +.foot-grid{display:grid;grid-template-columns:repeat(auto-fit,minmax(min(100%,200px),1fr));gap:32px} +.foot-brand{display:flex;flex-direction:column;gap:12px;max-width:34ch} +.foot-brand p{color:var(--ash);font-size:14px} +.foot-col{display:flex;flex-direction:column;gap:10px;font-size:15px} +.foot-col .eyebrow{margin-bottom:4px} +.foot-col a{color:var(--ink-2)} +.foot-col a:hover{color:var(--molten)} +.foot-base{display:flex;flex-wrap:wrap;justify-content:space-between;gap:8px 24px;margin-top:36px;padding-top:20px;border-top:1px solid var(--line);font-size:13px;color:var(--ash)} @@ -393,11 +401,27 @@ footer .fl{display:flex;flex-wrap:wrap;gap:20px} diff --git a/site/litepaper.html b/site/litepaper.html index eb27c0e02..3e3112bb2 100644 --- a/site/litepaper.html +++ b/site/litepaper.html @@ -79,6 +79,30 @@ td.num{font-family:'IBM Plex Mono',monospace;font-variant-numeric:tabular-nums;w pre{margin:0 0 16px;padding:16px 18px;background:var(--code);color:var(--code-ink);border-radius:12px;font-family:'IBM Plex Mono',monospace;font-size:13px;line-height:1.6;overflow-x:auto} .foot{padding-block:28px 48px;border-top:1px solid var(--line);display:flex;flex-wrap:wrap;gap:16px 32px;justify-content:space-between;font-size:14px;color:var(--quiet)} .mark{flex:0 0 auto} +article section{display:none;border-top:0;padding-top:0} +article section.active{display:block} +body.all article section{display:block;border-top:1px solid var(--line);padding-top:34px} +body.all article section:first-of-type{border-top:0;padding-top:0} +.toc li a.active{background:var(--tint);color:var(--ink);font-weight:600} +.toc li a.active::before{color:var(--ember)} +.toc .mode{margin-top:14px;display:flex;gap:8px} +.toc .mode button{flex:1;min-height:40px;border-radius:8px;border:1px solid var(--line);background:var(--surface);color:var(--ink-2);font:inherit;font-size:13px;cursor:pointer} +.toc .mode button.on{background:var(--ink);color:var(--ground);border-color:var(--ink)} +.pager{display:flex;justify-content:space-between;gap:12px;margin-top:36px;padding-top:22px;border-top:1px solid var(--line)} +.pager a{display:inline-flex;align-items:center;gap:8px;min-height:44px;padding:10px 16px;border:1px solid var(--line);border-radius:10px;color:var(--ink);font-size:14px;max-width:48%} +.pager a:hover{text-decoration:none;border-color:var(--ember)} +.pager a span{color:var(--quiet);font-family:'IBM Plex Mono',monospace;font-size:11px;letter-spacing:.12em;text-transform:uppercase} +.pager a.next{margin-left:auto;text-align:right} +body.all .pager{display:none} +@media (max-width:959px){ + .toc{position:sticky;top:0;z-index:5;background:var(--ground);margin-inline:-16px;padding:10px 16px;border-bottom:1px solid var(--line)} + .toc ol{flex-direction:row;overflow-x:auto;gap:6px;scrollbar-width:none;-webkit-overflow-scrolling:touch;padding-bottom:2px} + .toc ol::-webkit-scrollbar{display:none} + .toc li a{white-space:nowrap;border:1px solid var(--line);border-radius:999px;padding:8px 12px;font-size:13px;gap:6px} + .toc li a::before{min-width:0} + .toc .dl{display:none} + .toc .mode{margin-top:8px} +} @media (prefers-reduced-motion:no-preference){.pull{transition:transform .2s}} @@ -126,6 +150,7 @@ pre{margin:0 0 16px;padding:16px 18px;background:var(--code);color:var(--code-in
  • Questions builders ask
  • What Igneum does not claim
    • +
    Latin igneum: fiery. A cupel is the vessel in which metal is proven by fire.
    @@ -449,5 +474,32 @@ pre{margin:0 0 16px;padding:16px 18px;background:var(--code);color:var(--code-in +