From 192a683d0822e9e69a3f9049a8fd28467a35804f Mon Sep 17 00:00:00 2001 From: igneum-josh <337424239+igneum-josh@users.noreply.github.com> Date: Tue, 6 Oct 2026 08:36:12 +0100 Subject: [PATCH] Counter ASIC 3.0 items 6 and 7: family step-cost probes (Metal, CUDA), the PC 2 playbook, the M5 Max rows in the bench-log --- docs/bench-log.md | 40 +++++ proto-cuda/family-probe.cu | 179 ++++++++++++++++++++++ proto-metal/family-probe.swift | 199 +++++++++++++++++++++++++ tools/ca3-reserve/make-pc2-playbook.sh | 20 +++ tools/ca3-reserve/pc2-family-probe.ps1 | 63 ++++++++ 5 files changed, 501 insertions(+) create mode 100644 proto-cuda/family-probe.cu create mode 100644 proto-metal/family-probe.swift create mode 100755 tools/ca3-reserve/make-pc2-playbook.sh create mode 100644 tools/ca3-reserve/pc2-family-probe.ps1 diff --git a/docs/bench-log.md b/docs/bench-log.md index cc01b90ec..a0412f266 100644 --- a/docs/bench-log.md +++ b/docs/bench-log.md @@ -1961,3 +1961,43 @@ Ten distinct programs (seed strings `igneum-devnet-v4-epoch0`, `/epoch1` .. `/ep Cache fill 1.95 ms GPU (192.4 ms one core), dataset build 20.8 ms GPU for 1 GiB. The devnet pack three times through `packbench --pack ../proto-cuda/packs/igneum-devnet-v4-epoch0 --batches 1 --batch-log2 20 --group 256` (the pack's two libraries, `memhard.metal` and `program.metal`): compile 79 ms, 1 ms, 1 ms (the system shader cache answers the identical source from the second run); cache fill 0.6 to 0.7 ms GPU, dataset build 20.7 to 20.8 ms GPU. Reading: a fresh program compiles in about 18 ms on this card with the Metal compiler service warm, 79 ms for a pack with its dataset kernels, up to 1.8 s cold (the variant-racing entry's first seed), 0 to 444 ms at the fleet's live boundaries (M11). The hot table fill of layer 5 is 0.07 to 0.22 ms (ca2-cache). So the Mac's per-epoch compile-ahead is under 2 s without the race and about 38 s with it (M11: 34.0 / 34.9 / 37.8 s), and the race is the only item visible against the 600-s window in which the program is known (lead 1,200 s minus the 600-s VDF, fixed at every epoch length). PC cards, cited in the plan: RTX 5090 NVRTC 151 to 180 ms, prepare 0.5 to 1.0 s without the dataset (M11), race one round about 37 s; RX 9070 XT OpenCL compile NOT MEASURED at the current worker (owed: `host.c` times `clBuildProgram` only in the `prepare` path and no `prepared` line from gfx1201 is in any upload); Intel UHD build 3.0 to 6.4 s (M11). Floor by the rule (slowest compile-ahead under 10% of the epoch and inside the window, dataset excluded): 600 DAA s, carried by the race at 6.3% of 600 s; with the race off (M11 found base wins on both the 5090 and the Mac) the slowest measured row is the Intel iGPU at 1.1%. Consequences per tier and the difficulty-settle constraint (24% of a 600-s epoch in settle at the measured 144 s) are in the plan. + +## 6 October 2026, Counter ASIC 3.0 item 6: the reserve families' step costs + +Branch `ca3-reserve`, worker "reserve" (`docs/plans/counter-asic-3-reserve.md` carries the proposed order and spec text; this entry carries the measurements). Method: the dot4 probe's dependent chain (`docs/analysis/int8-matrix-family.md` section 4), one op of the family per step per lane, 1,048,576 lanes x 4,096 steps, best of 3 dispatches per run, three runs, bit-exact against a CPU reference on two whole 32-lane warps (the shuffle rows need the whole warp). Every chain has the same glue (`acc = OP(acc, x, y); x = x * K + acc; y = rotl(y, 7) ^ (acc + s)`), so the "step cost" is the family's one op plus four glue ops against the add-xor-rotate chain of the 9070 XT bench-log entry (`alu`: `x = x * K + rotl(y, 7); y = (y ^ x) + s`, 5 ops per step counted, no `acc`). Reference rows are live families (`alu`; `rotr` = the live `rotr_var` text; `shflx` = the live `shfl`, lane XOR 8). Candidate rows are the seven families of spec 1.13.2 (`shl`, `shr`, `bfe` with the vendor's extract function and `bfec` the C form `(y >> 7) & 0x1fff`, `andn`, `perm` = bytes (b1, b3, b0, b2), `popc` and `clz` folded by add, `sel` on bit 5, `shfla` = lane + 3 mod 32). Comparison rows: `dot4u` and `dot4s` (Apple, emulated), `dot4i` (`__dp4a`) and `mm8` (one `mma.sync.m8n8k16` u8 per step per warp, inline PTX) on CUDA. Sources: `proto-metal/family-probe.swift`, `proto-cuda/family-probe.cu`, the PC 2 job `tools/ca3-reserve/pc2-family-probe.ps1` (made by `make-pc2-playbook.sh`). G steps/s is the whole card's dependent-step throughput; ops per step counted = the family's op plus 4 glue (`alu` 5, `shfla` and `shflx` 6: shuffle plus xor, `mm8` 1 mma plus 4). + +**Apple M5 Max, Metal** (`swiftc -O -o family-probe family-probe.swift -framework Metal` under `with-lock.sh build`; three runs of `with-lock.sh measure ./family-probe --reps 3`, 07:29:05 to 07:29:08 UTC, load average 7.64 / 7.59 / 7.14 before and after every run (the Mac was loaded by other agents' builds the whole morning; the measure lock held, the GPU idle: the Mac mines nothing), GPU start-to-end time): + +| kernel | best ms, runs 1 / 2 / 3 | best of the three, ms | G steps/s (best) | ns per step (best) | ops per step counted | step cost (ratio to `alu`, best) | bit-exact, 3 runs | +|---|---|---|---|---|---|---|---| +| alu | 5.039 / 4.918 / 4.873 | 4.873 | 881 | 1,190 | 5 | 1.00 | yes | +| rotr (live) | 5.610 / 5.533 / 5.492 | 5.492 | 782 | 1,341 | 5 | 1.13 | yes | +| shflx (live) | 4.196 / 4.259 / 4.223 | 4.196 | 1,024 | 1,024 | 6 | 0.86 | yes | +| shl | 4.120 / 4.202 / 4.119 | 4.119 | 1,043 | 1,006 | 5 | 0.85 | yes | +| shr | 4.296 / 4.194 / 4.319 | 4.194 | 1,024 | 1,024 | 5 | 0.86 | yes | +| bfe (`extract_bits`) | 3.731 / 3.805 / 3.816 | 3.731 | 1,151 | 911 | 5 | 0.77 | yes | +| bfec (C form) | 3.762 / 3.818 / 3.646 | 3.646 | 1,178 | 890 | 5 | 0.75 | yes | +| andn | 3.680 / 3.732 / 3.676 | 3.676 | 1,168 | 898 | 5 | 0.75 | yes | +| perm | 5.500 / 5.548 / 5.524 | 5.500 | 781 | 1,343 | 5 | 1.13 | yes | +| popc | 4.261 / 4.223 / 4.262 | 4.223 | 1,017 | 1,031 | 5 | 0.87 | yes | +| clz | 4.918 / 4.916 / 4.914 | 4.914 | 874 | 1,200 | 5 | 1.01 | yes | +| sel | 3.718 / 3.831 / 3.718 | 3.718 | 1,155 | 908 | 5 | 0.76 | yes | +| shfla (lane + 3) | 9.337 / 9.323 / 9.287 | 9.287 | 462 | 2,267 | 6 | 1.91 | yes | +| dot4u (emulated) | 7.818 / 8.044 / 7.925 | 7.818 | 549 | 1,909 | 5 | 1.60 | yes | +| dot4s (emulated) | 23.264 / 23.254 / 23.045 | 23.045 | 186 | 5,626 | 5 | 4.73 | yes | + +Reading of the Mac rows. The run-to-run spread is under 4% on every row. The dot4 rows reproduce the 5 October figures (1.6x unsigned, 4.7x signed), which is the check on the method. A step cost under 1.00 means the family's op plus the glue is cheaper than the five-op reference chain: the reference's two registers are a tighter dependency than the three-register candidate chains, and Apple's shifts, extract, andn and select each cost about what an xor costs. Three rows cost more than the reference: `perm` (1.13: no byte-permute function in MSL; the `uchar4` swizzle compiles to shifts and masks, so a byte permute is emulated on Apple at about the price of the live `rotr`), `clz` (1.01) and `shfla` (1.91: a shuffle by a computed lane index costs 2.2x the live xor shuffle on Apple, `simd_shuffle` against `simd_shuffle_xor`; the second shuffle form is the one candidate Apple pays for). `mm8` as a chain on Apple is owed (Metal 4 `matmul2d`; this toolchain is Swift 5.8 without the tensor API). + +**RTX 5090 (PC 2, 1ccfe586), CUDA**: PENDING the PC 2 job (`run-ca3-family-pc2-20261006`, published only after `/tmp/igneum-devnet/pc2-ca3.clear` and under the mkdir lock; the card to itself: `--stop-miners`, prover off for the run). The rows land in this entry when the closing report is read. + +**RX 9070 XT (PC 1, ae432dc7), OpenCL**: OWED. PC 1 is Josh's desk and not released today (the brief's rule); the OpenCL twin of the probe (`__builtin_amdgcn_*` paths for `v_bfe_u32`, `v_perm_b32`, `v_bcnt_u32_b32`, `v_cndmask_b32`, `ds_bpermute_b32`) is the next job on that card. + +Consequences per tier, Mac rows (the hash is latency-bound by 128 dependent DRAM reads; a family at `W_new` = 4 points is about 4% of the 64 instructions, so these per-op costs bound a family's hash-rate cost and are not hash rates; the 5% rule of 1.13.2 is argued from them, not measured, until a family is live): + +| Tier | What the rows mean | What is being done | +|---|---|---| +| Apple user (M-series laptop or desktop, the app's Metal worker) | six of the seven candidates (`shl`, `shr`, `bfe`, `andn`, `popc`, `sel`) cost at most the live `rotr` step; `clz` the same as the reference; `perm` 1.13x (emulated); `shfla` 1.91x, the only candidate over the live `shfl`'s cost by more than 2x on this card. At 4 points of 64 a 1.91x op costs under 1% of the program's ALU time, itself a small share of a latency-bound hash (approximate: argued, measured when live) | the proposed order puts `shfla` after the plain datapath families (R6), so Apple pays it last; `mm8` stays last | +| NVIDIA user (8 to 32 GB card) | pending the 5090 rows above | the PC 2 job | +| AMD user (RX 9070 XT, 16 GB) | owed: no row today | the PC 1 job when the desk is free | +| A rig or a pool user | the same per-card figures; no family changes the dependent-read bound | nothing until a family is live | +| A chip | every candidate but `mm8` is a 32-bit datapath structure (barrel shifter, byte crossbar, popcount tree, 32-lane shuffle crossbar: `docs/plans/counter-asic-3-reserve.md` section 3 names them with approximate areas); none is licensable as a block the way an int8 matrix unit is | the reserve order of that document | diff --git a/proto-cuda/family-probe.cu b/proto-cuda/family-probe.cu new file mode 100644 index 000000000..ed4519b4e --- /dev/null +++ b/proto-cuda/family-probe.cu @@ -0,0 +1,179 @@ +// family-probe (CUDA): the step cost of every reserve candidate family of spec 1.13.2 on NVIDIA, standalone (no pack). +// Counter ASIC 3.0 item 6 (docs/plans/counter-asic-3-reserve.md), 6 October 2026. PC job; not run on the Mac. Same +// method as dot4-probe.cu (docs/analysis/int8-matrix-family.md section 4): a dependent chain of one op per step per +// lane, 1,048,576 lanes x 4,096 steps, best of N, event time, bit-exact against a CPU reference on two whole warps. +// +// Every chain has the dot4 probe's glue: acc = OP(acc, x, y); x = x * K + acc; y = rotl(y, 7) ^ (acc + s). Reference +// rows (live families): alu (add-xor-rotate, 5 ops per step counted, no acc), rotr (the live rotr_var text), shflx +// (the live shfl: dst ^= src of lane (lane XOR 8)). Candidate rows, the seven families of 1.13.2: +// shl acc = y << (x & 31) shr acc = y >> (x & 31) +// bfe acc = bfe.u32(y, 7, 13) (inline PTX) bfec acc = (y >> 7) & 0x1fff (the C form, what nvcc emits) +// andn acc = y & ~x perm acc = __byte_perm(y, y, 0x2031): bytes (y.b1, y.b3, y.b0, y.b2) +// popc acc = acc + __popc(x) clz acc = acc + __clz(x) (clz(0) = 32) +// sel acc = bit 5 of y ? x : acc shfla acc ^= __shfl_sync(x, (lane + 3) & 31) +// Comparison rows: dot4i (__dp4a, PTX dp4a.s32.s32, the dot4 probe's intrinsic row) and mm8 (one mma.sync +// m8n8k16 u8 x u8 -> s32 per step per warp, inline PTX, sm_75+; PTX ISA 9.4 section 9.7.16.5.3 fragment layout: +// lane i holds A[i/4][4(i%4)..+3] as the bytes of x, B[4(i%4)..+3][i/4] as the bytes of y, C and D [i/4][2(i%4)+j] +// as acc (j = 0) and acc2 (j = 1); the chain carries d0 forward as acc, d1 as acc2). The shifted and extracted forms +// take y (fresh every step) as the shifted register so the chain never drains to 0. +// +// Build (Linux or WSL2, CUDA Toolkit): nvcc -O2 -arch=sm_120 -o family-probe family-probe.cu +// (older toolkits: nvcc -O2 -gencode arch=compute_89,code=compute_89 ...) +// Run: family-probe [--lanes N] [--steps N] [--reps N] [--device N] +#include +#include +#include +#include +#include + +__host__ __device__ inline uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; } +__host__ __device__ inline uint32_t rotl32(uint32_t v, uint32_t n) { return (v << n) | (v >> (32u - n)); } +__host__ __device__ inline uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); } + +#define K32 0x9E3779B1u +#define CHAIN(NAME, OP) \ +__global__ void NAME(uint32_t steps, uint32_t seed, uint32_t* out) { \ + uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t lane = threadIdx.x & 31u; (void)lane; \ + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; uint32_t acc = pm_mix(x); \ + for (uint32_t s = 0; s < steps; ++s) { OP; x = x * K32 + acc; y = rotl32(y, 7u) ^ (acc + s); } \ + out[g] = acc ^ x ^ y; \ +} + +__global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) { + uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + for (uint32_t s = 0; s < steps; ++s) { x = x * K32 + rotl32(y, 7u); y = (y ^ x) + s; } + out[g] = x ^ y; +} +__device__ inline uint32_t bfe_ptx(uint32_t v) { uint32_t r; asm("bfe.u32 %0, %1, 7, 13;" : "=r"(r) : "r"(v)); return r; } +CHAIN(probe_rotr, acc = rotr_var(y, x)) +CHAIN(probe_shflx, acc = acc ^ __shfl_xor_sync(0xffffffffu, x, 8)) +CHAIN(probe_shl, acc = y << (x & 31u)) +CHAIN(probe_shr, acc = y >> (x & 31u)) +CHAIN(probe_bfe, acc = bfe_ptx(y)) +CHAIN(probe_bfec, acc = (y >> 7u) & 0x1fffu) +CHAIN(probe_andn, acc = y & ~x) +CHAIN(probe_perm, acc = __byte_perm(y, y, 0x2031)) +CHAIN(probe_popc, acc = acc + (uint32_t)__popc((int)x)) +CHAIN(probe_clz, acc = acc + (uint32_t)__clz((int)x)) +CHAIN(probe_sel, acc = ((y >> 5u) & 1u) ? x : acc) +CHAIN(probe_shfla, acc = acc ^ __shfl_sync(0xffffffffu, x, (int)((lane + 3u) & 31u))) +CHAIN(probe_dot4i, acc = (uint32_t)__dp4a((int)x, (int)y, (int)acc)) + +// one mma.sync m8n8k16 per step per warp; the chain carries d0 as acc, d1 as acc2 +__global__ void probe_mm8(uint32_t steps, uint32_t seed, uint32_t* out) { + uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; uint32_t acc = pm_mix(x), acc2 = pm_mix(acc); + for (uint32_t s = 0; s < steps; ++s) { + uint32_t d0, d1; + asm volatile("mma.sync.aligned.m8n8k16.row.col.s32.u8.u8.s32 {%0,%1}, {%2}, {%3}, {%4,%5};" + : "=r"(d0), "=r"(d1) : "r"(x), "r"(y), "r"(acc), "r"(acc2)); + acc = d0; acc2 = d1; + x = x * K32 + acc; y = rotl32(y, 7u) ^ (acc + s); + } + out[g] = acc ^ acc2 ^ x ^ y; +} + +// ---- CPU reference: one whole warp (lanes g0 .. g0+31, lane = g AND 31) ---- +static uint32_t clz32(uint32_t x) { if (!x) return 32; uint32_t n = 0; while (!(x & 0x80000000u)) { x <<= 1; ++n; } return n; } +static uint32_t popc32(uint32_t x) { uint32_t n = 0; while (x) { n += x & 1u; x >>= 1; } return n; } +static uint32_t perm_ref(uint32_t y) { return ((y >> 8) & 0xffu) | (((y >> 24) & 0xffu) << 8) | ((y & 0xffu) << 16) | (((y >> 16) & 0xffu) << 24); } +static uint32_t dot4s_ref(uint32_t a, uint32_t b, uint32_t acc) { + int32_t r = (int32_t)acc; + for (int i = 0; i < 4; ++i) { int32_t ba = (int32_t)(int8_t)((a >> (8 * i)) & 0xffu); int32_t bb = (int32_t)(int8_t)((b >> (8 * i)) & 0xffu); r = (int32_t)((uint32_t)r + (uint32_t)(ba * bb)); } + return (uint32_t)r; +} +static void warp_ref(const char* name, uint32_t g0, uint32_t seed, uint32_t steps, uint32_t* res) { + uint32_t x[32], y[32], acc[32], acc2[32], xs[32]; + for (int l = 0; l < 32; ++l) { x[l] = pm_mix((g0 + l) ^ seed); y[l] = x[l] ^ 0x5bd1e995u; acc[l] = pm_mix(x[l]); acc2[l] = pm_mix(acc[l]); } + if (!strcmp(name, "alu")) { + for (uint32_t s = 0; s < steps; ++s) for (int l = 0; l < 32; ++l) { x[l] = x[l] * K32 + rotl32(y[l], 7u); y[l] = (y[l] ^ x[l]) + s; } + for (int l = 0; l < 32; ++l) res[l] = x[l] ^ y[l]; + return; + } + for (uint32_t s = 0; s < steps; ++s) { + memcpy(xs, x, sizeof xs); + if (!strcmp(name, "mm8")) { + uint32_t D[8][8]; + for (int r = 0; r < 8; ++r) for (int n = 0; n < 8; ++n) { + uint32_t c = (n & 1) ? acc2[r * 4 + n / 2] : acc[r * 4 + n / 2]; + for (int k = 0; k < 16; ++k) { + uint32_t a = (xs[r * 4 + k / 4] >> (8 * (k % 4))) & 0xffu; + uint32_t b = (y[n * 4 + k / 4] >> (8 * (k % 4))) & 0xffu; + c += a * b; + } + D[r][n] = c; + } + for (int l = 0; l < 32; ++l) { acc[l] = D[l / 4][2 * (l % 4)]; acc2[l] = D[l / 4][2 * (l % 4) + 1]; } + for (int l = 0; l < 32; ++l) { x[l] = xs[l] * K32 + acc[l]; y[l] = rotl32(y[l], 7u) ^ (acc[l] + s); } + continue; + } + for (int l = 0; l < 32; ++l) { + uint32_t xv = xs[l], yv = y[l]; + if (!strcmp(name, "rotr")) acc[l] = rotr_var(yv, xv); + else if (!strcmp(name, "shflx")) acc[l] = acc[l] ^ xs[l ^ 8]; + else if (!strcmp(name, "shl")) acc[l] = yv << (xv & 31u); + else if (!strcmp(name, "shr")) acc[l] = yv >> (xv & 31u); + else if (!strcmp(name, "bfe") || !strcmp(name, "bfec")) acc[l] = (yv >> 7u) & 0x1fffu; + else if (!strcmp(name, "andn")) acc[l] = yv & ~xv; + else if (!strcmp(name, "perm")) acc[l] = perm_ref(yv); + else if (!strcmp(name, "popc")) acc[l] = acc[l] + popc32(xv); + else if (!strcmp(name, "clz")) acc[l] = acc[l] + clz32(xv); + else if (!strcmp(name, "sel")) acc[l] = ((yv >> 5u) & 1u) ? xv : acc[l]; + else if (!strcmp(name, "shfla")) acc[l] = acc[l] ^ xs[(l + 3) & 31]; + else if (!strcmp(name, "dot4i")) acc[l] = dot4s_ref(xv, yv, acc[l]); + else { printf("no reference for %s\n", name); exit(3); } + x[l] = xv * K32 + acc[l]; + y[l] = rotl32(yv, 7u) ^ (acc[l] + s); + } + } + for (int l = 0; l < 32; ++l) res[l] = acc[l] ^ x[l] ^ y[l] ^ (!strcmp(name, "mm8") ? acc2[l] : 0u); +} + +#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { printf("CUDA error %s at %s:%d\n", cudaGetErrorString(e), __FILE__, __LINE__); return 1; } } while (0) +typedef void (*kernel_t)(uint32_t, uint32_t, uint32_t*); + +int main(int argc, char** argv) { + uint32_t lanes = 1u << 20, steps = 4096u; int reps = 3, device = 0; + for (int i = 1; i < argc; ++i) { + if (!strcmp(argv[i], "--lanes") && i + 1 < argc) lanes = (uint32_t)strtoul(argv[++i], 0, 10); + else if (!strcmp(argv[i], "--steps") && i + 1 < argc) steps = (uint32_t)strtoul(argv[++i], 0, 10); + else if (!strcmp(argv[i], "--reps") && i + 1 < argc) reps = atoi(argv[++i]); + else if (!strcmp(argv[i], "--device") && i + 1 < argc) device = atoi(argv[++i]); + else { printf("unknown argument %s\n", argv[i]); return 2; } + } + CK(cudaSetDevice(device)); + cudaDeviceProp p; CK(cudaGetDeviceProperties(&p, device)); + printf("family-probe (CUDA) on %s, sm_%d%d, %d SMs, %d MHz, lanes %u, steps %u, best of %d, event time\n", p.name, p.major, p.minor, p.multiProcessorCount, p.clockRate / 1000, lanes, steps, reps); + uint32_t* d_out; CK(cudaMalloc(&d_out, (size_t)lanes * 4)); + uint32_t* h_out = (uint32_t*)malloc((size_t)lanes * 4); + cudaEvent_t e0, e1; CK(cudaEventCreate(&e0)); CK(cudaEventCreate(&e1)); + const char* names[] = { "alu", "rotr", "shflx", "shl", "shr", "bfe", "bfec", "andn", "perm", "popc", "clz", "sel", "shfla", "dot4i", "mm8" }; + kernel_t kernels[] = { probe_alu, probe_rotr, probe_shflx, probe_shl, probe_shr, probe_bfe, probe_bfec, probe_andn, probe_perm, probe_popc, probe_clz, probe_sel, probe_shfla, probe_dot4i, probe_mm8 }; + const int nk = (int)(sizeof(names) / sizeof(names[0])); + printf("| kernel | lanes | steps | best ms | G steps/s | ns per step | ratio to alu | warps 0 and last ok |\n|---|---|---|---|---|---|---|---|\n"); + float alu_best = 0; + for (int k = 0; k < nk; ++k) { + float best = 1e30f; int ok = 1; + for (int r = 0; r < reps; ++r) { + uint32_t seed = 0x2468aceu + (uint32_t)r * 0x9E3779B9u; + CK(cudaEventRecord(e0)); + kernels[k]<<>>(steps, seed, d_out); + CK(cudaEventRecord(e1)); CK(cudaEventSynchronize(e1)); CK(cudaGetLastError()); + float ms = 0; CK(cudaEventElapsedTime(&ms, e0, e1)); if (ms < best) best = ms; + CK(cudaMemcpy(h_out, d_out, (size_t)lanes * 4, cudaMemcpyDeviceToHost)); + uint32_t g0s[2] = { 0u, lanes - 32u }; + for (int j = 0; j < 2; ++j) { + uint32_t want[32]; warp_ref(names[k], g0s[j], seed, steps, want); + for (int l = 0; l < 32; ++l) if (h_out[g0s[j] + l] != want[l]) { ok = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", names[k], g0s[j] + l, h_out[g0s[j] + l], want[l]); break; } + } + } + if (k == 0) alu_best = best; + double sps = (double)lanes * (double)steps / (best / 1000.0); + double ratio = alu_best > 0 ? best / alu_best : 0; + printf("| %s | %u | %u | %.3f | %.2f | %.3f | %.2f | %s |\n", names[k], lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ratio, ok ? "yes" : "NO"); + printf("RESULT FAMILY vendor=nvidia device=\"%s\" kernel=%s lanes=%u steps=%u best_ms=%.3f gsteps_per_s=%.2f ns_per_step=%.3f ratio_alu=%.3f ok=%d\n", p.name, names[k], lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ratio, ok); + } + printf("family-probe: done\n"); + return 0; +} diff --git a/proto-metal/family-probe.swift b/proto-metal/family-probe.swift new file mode 100644 index 000000000..5dfb2c47c --- /dev/null +++ b/proto-metal/family-probe.swift @@ -0,0 +1,199 @@ +// family-probe: the step cost of every reserve candidate family of spec 1.13.2 on Apple silicon, standalone (no pack). +// Counter ASIC 3.0 item 6 (docs/plans/counter-asic-3-reserve.md), 6 October 2026. Same method as dot4-probe.swift +// (docs/analysis/int8-matrix-family.md section 4): a dependent chain of one op per step per lane, 1,048,576 lanes x +// 4,096 steps, best of N, GPU start-to-end time, bit-exact against a CPU reference on two whole 32-lane SIMD groups. +// +// Every chain has the dot4 probe's glue: acc = OP(acc, x, y); x = x * K + acc; y = rotate(y, 7) ^ (acc + s). The +// reference rows are the live families: alu (the add-xor-rotate chain of the 9070 XT bench-log entry, 5 ops per step +// counted, no acc), rotr (dst = rotr(src, src2 AND 31), the live rotr_var text), shflx (dst = dst XOR src of lane +// (lane XOR 8), the live shfl). The candidate rows are the seven families of 1.13.2: +// shl acc = y << (x AND 31) variable left shift +// shr acc = y >> (x AND 31) variable logical right shift +// bfe acc = extract_bits(y, 7, 13) bit-field extract, immediate offset and width (bfec = the C form) +// andn acc = y AND NOT x +// perm acc = bytes (y.b1, y.b3, y.b0, y.b2) byte permute by an immediate selector +// popc acc = acc + popcount(x) popcount folded by add +// clz acc = acc + clz(x) count-leading-zeros folded by add (clz(0) = 32) +// sel acc = bit 5 of y ? x : acc three-register select +// shfla acc = acc XOR x of lane ((lane + 3) mod 32) the second shuffle form +// and for comparison the dp4a-class rows of the dot4 probe (dot4u unsigned, dot4s signed, both emulated on Apple). +// The shifted and extracted forms take y (fresh every step) as the shifted register so the chain never drains to 0; +// the op count per step is the family's one op plus the same glue in every row. mm8 as a chain (Metal 4 matmul2d) is +// owed: this toolchain (Swift 5.8) has no Metal 4 tensor API. +// +// Build: swiftc -O -o family-probe family-probe.swift -framework Metal +// Run: ./family-probe [--lanes N] [--steps N] [--reps N] +import Foundation +import Metal + +let source = """ +#include +using namespace metal; + +inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; } +inline uint rotr_var(uint x, uint n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); } +inline int dot4_s(uint a, uint b, int acc) { + int4 va = int4(as_type(a)); int4 vb = int4(as_type(b)); + return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w; +} +inline uint dot4_u(uint a, uint b, uint acc) { + uint4 va = uint4(as_type(a)); uint4 vb = uint4(as_type(b)); + return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w; +} + +#define CHAIN(NAME, OP) \\ +kernel void NAME(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]], device uint* out [[buffer(2)]], \\ + uint g [[thread_position_in_grid]], ushort lane [[thread_index_in_simdgroup]]) { \\ + uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; uint acc = pm_mix(x); \\ + for (uint s = 0u; s < steps; ++s) { OP; x = x * 0x9E3779B1u + acc; y = rotate(y, 7u) ^ (acc + s); } \\ + out[g] = acc ^ x ^ y; \\ +} + +kernel void probe_alu(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]], device uint* out [[buffer(2)]], + uint g [[thread_position_in_grid]]) { + uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + for (uint s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; } + out[g] = x ^ y; +} +CHAIN(probe_rotr, acc = rotr_var(y, x)) +CHAIN(probe_shflx, acc = acc ^ simd_shuffle_xor(x, (ushort)8)) +CHAIN(probe_shl, acc = y << (x & 31u)) +CHAIN(probe_shr, acc = y >> (x & 31u)) +CHAIN(probe_bfe, acc = extract_bits(y, 7u, 13u)) +CHAIN(probe_bfec, acc = (y >> 7u) & 0x1fffu) +CHAIN(probe_andn, acc = y & ~x) +CHAIN(probe_perm, uchar4 b_ = as_type(y); acc = as_type(uchar4(b_.y, b_.w, b_.x, b_.z))) +CHAIN(probe_popc, acc = acc + popcount(x)) +CHAIN(probe_clz, acc = acc + clz(x)) +CHAIN(probe_sel, acc = select(acc, x, ((y >> 5u) & 1u) != 0u)) +CHAIN(probe_shfla, acc = acc ^ simd_shuffle(x, (ushort)((lane + 3u) & 31u))) +CHAIN(probe_dot4u, acc = dot4_u(x, y, acc)) +CHAIN(probe_dot4s, acc = uint(dot4_s(x, y, int(acc)))) +""" + +// ---- CPU reference: one whole 32-lane SIMD group (lanes g0 .. g0+31, lane = g AND 31), bit-exact ---- +func pmMix(_ v: UInt32) -> UInt32 { + var x = v + x ^= x >> 16; x = x &* 0x7feb352d; x ^= x >> 15; x = x &* 0x846ca68b; x ^= x >> 16 + return x +} +func rotl(_ v: UInt32, _ n: UInt32) -> UInt32 { (v << n) | (v >> (32 - n)) } +func rotrVar(_ x: UInt32, _ nIn: UInt32) -> UInt32 { let n = nIn & 31; return (x >> n) | (x << ((32 - n) & 31)) } +func clz32(_ x: UInt32) -> UInt32 { x == 0 ? 32 : UInt32(x.leadingZeroBitCount) } +func perm(_ y: UInt32) -> UInt32 { + ((y >> 8) & 0xff) | (((y >> 24) & 0xff) << 8) | ((y & 0xff) << 16) | (((y >> 16) & 0xff) << 24) +} +func dot4sRef(_ a: UInt32, _ b: UInt32, _ acc: UInt32) -> UInt32 { + var r = Int32(bitPattern: acc) + for i in 0..<4 { + let ba = Int32(Int8(truncatingIfNeeded: a >> (8 * UInt32(i)))) + let bb = Int32(Int8(truncatingIfNeeded: b >> (8 * UInt32(i)))) + r = r &+ ba &* bb + } + return UInt32(bitPattern: r) +} +func dot4uRef(_ a: UInt32, _ b: UInt32, _ acc: UInt32) -> UInt32 { + var r = acc + for i in 0..<4 { + let ba = UInt32(UInt8(truncatingIfNeeded: a >> (8 * UInt32(i)))) + let bb = UInt32(UInt8(truncatingIfNeeded: b >> (8 * UInt32(i)))) + r = r &+ ba &* bb + } + return r +} +// the 32 outputs of the SIMD group whose first lane is g0 +func warpRef(kernel: String, g0: UInt32, seed: UInt32, steps: UInt32) -> [UInt32] { + var x = [UInt32](repeating: 0, count: 32), y = [UInt32](repeating: 0, count: 32), acc = [UInt32](repeating: 0, count: 32) + for l in 0..<32 { x[l] = pmMix((g0 + UInt32(l)) ^ seed); y[l] = x[l] ^ 0x5bd1e995; acc[l] = pmMix(x[l]) } + if kernel == "probe_alu" { + for s in 0..> (xv & 31) + case "probe_bfe", "probe_bfec": acc[l] = (yv >> 7) & 0x1fff + case "probe_andn": acc[l] = yv & ~xv + case "probe_perm": acc[l] = perm(yv) + case "probe_popc": acc[l] = acc[l] &+ UInt32(xv.nonzeroBitCount) + case "probe_clz": acc[l] = acc[l] &+ clz32(xv) + case "probe_sel": acc[l] = ((yv >> 5) & 1) != 0 ? xv : acc[l] + case "probe_shfla": acc[l] = acc[l] ^ xs[(l + 3) & 31] + case "probe_dot4u": acc[l] = dot4uRef(xv, yv, acc[l]) + case "probe_dot4s": acc[l] = dot4sRef(xv, yv, acc[l]) + default: fatalError("no reference for \(kernel)") + } + x[l] = xv &* 0x9E3779B1 &+ acc[l] + y[l] = rotl(yv, 7) ^ (acc[l] &+ s) + } + } + return (0..<32).map { acc[$0] ^ x[$0] ^ y[$0] } +} + +var lanes = 1 << 20, steps: UInt32 = 4096, reps = 3 +var args = Array(CommandLine.arguments.dropFirst()) +while !args.isEmpty { + let a = args.removeFirst() + switch a { + case "--lanes": lanes = Int(args.removeFirst())! + case "--steps": steps = UInt32(args.removeFirst())! + case "--reps": reps = Int(args.removeFirst())! + default: print("unknown argument \(a)"); exit(2) + } +} + +guard let dev = MTLCreateSystemDefaultDevice() else { print("no Metal device"); exit(1) } +let lib: MTLLibrary +do { lib = try dev.makeLibrary(source: source, options: nil) } catch { print("compile failed: \(error)"); exit(1) } +let queue = dev.makeCommandQueue()! +let outBuf = dev.makeBuffer(length: lanes * 4, options: .storageModeShared)! +let names = ["probe_alu", "probe_rotr", "probe_shflx", "probe_shl", "probe_shr", "probe_bfe", "probe_bfec", "probe_andn", "probe_perm", + "probe_popc", "probe_clz", "probe_sel", "probe_shfla", "probe_dot4u", "probe_dot4s"] +print("family-probe on \(dev.name), lanes \(lanes), steps \(steps), best of \(reps), GPU start-to-end time") +print("| kernel | lanes | steps | best ms | G steps/s | ns per step | ratio to alu | warps 0 and last ok |") +print("|---|---|---|---|---|---|---|---|") +var aluBest = 0.0 +for name in names { + let fn = lib.makeFunction(name: name)! + let pso = try! dev.makeComputePipelineState(function: fn) + let tg = min(256, pso.maxTotalThreadsPerThreadgroup) + var best = Double.infinity + var okAll = true + for r in 0.. 0 ? best / aluBest : 0 + print(String(format: "| %@ | %d | %u | %.3f | %.2f | %.3f | %.2f | %@ |", name, lanes, steps, best, stepsPerS / 1e9, best * 1e6 / Double(steps), ratio, okAll ? "yes" : "NO")) + print(String(format: "RESULT FAMILY vendor=apple device=\"%@\" kernel=%@ lanes=%d steps=%u best_ms=%.3f gsteps_per_s=%.2f ns_per_step=%.3f ratio_alu=%.3f ok=%d", dev.name, name, lanes, steps, best, stepsPerS / 1e9, best * 1e6 / Double(steps), ratio, okAll ? 1 : 0)) +} +print("family-probe: done") diff --git a/tools/ca3-reserve/make-pc2-playbook.sh b/tools/ca3-reserve/make-pc2-playbook.sh new file mode 100755 index 000000000..7d0d434a1 --- /dev/null +++ b/tools/ca3-reserve/make-pc2-playbook.sh @@ -0,0 +1,20 @@ +#!/usr/bin/env bash +# Writes the PC 2 family-probe playbook: pc2-family-probe.ps1 with proto-cuda/family-probe.cu inlined at +# CU_PLACEHOLDER, into $1 (default: $TMPDIR/pc2-family-probe.ps1). Prints the sha256 of the .cu so the job's +# "RESULT source ... sha256" line can be checked against it. The shape follows tools/prover-floor/make-measure-playbook.sh. +set -euo pipefail +HERE="$(cd "$(dirname "$0")" && pwd)" +ROOT="$(cd "$HERE/../.." && pwd)" +OUT="${1:-${TMPDIR:-/tmp}/pc2-family-probe.ps1}" +CU="$ROOT/proto-cuda/family-probe.cu" +grep -q "^'@" "$CU" && { echo "the .cu has a line starting with '@ (would end the here-string)" >&2; exit 1; } +python3 - "$HERE/pc2-family-probe.ps1" "$CU" "$OUT" <<'PY' +import sys +tpl, cu, out = sys.argv[1], sys.argv[2], sys.argv[3] +s = open(tpl).read() +c = open(cu).read().rstrip('\n') +assert 'CU_PLACEHOLDER' in s +open(out, 'w').write(s.replace('CU_PLACEHOLDER', c)) +print("wrote", out) +PY +echo "family-probe.cu sha256 $(shasum -a 256 "$CU" | cut -c1-64) bytes $(wc -c < "$CU" | tr -d ' ')" diff --git a/tools/ca3-reserve/pc2-family-probe.ps1 b/tools/ca3-reserve/pc2-family-probe.ps1 new file mode 100644 index 000000000..1c2a78c88 --- /dev/null +++ b/tools/ca3-reserve/pc2-family-probe.ps1 @@ -0,0 +1,63 @@ +# Counter ASIC 3.0 item 6 (6 October 2026): the step cost of every reserve candidate family of spec 1.13.2 on PC 2's +# RTX 5090 (docs/plans/counter-asic-3-reserve.md). A signed `run` job published with --stop-miners: the app stops its +# miners before this script starts and restarts them when it ends (app/igneum-app/src/jobrun.rs, stop_miners_first); +# the script never quits, restarts or updates the app. It switches the prover off for the run and back on at the end +# (as tools/prover-floor/pc2-floor-measure.ps1 does), confirms the card is empty by nvidia-smi's compute-apps list +# (never api/state, which answers {} on this machine), writes proto-cuda/family-probe.cu (inlined below by +# make-pc2-playbook.sh, sha256 printed on both sides) into the job dir, compiles it with nvcc inside WSL2 +# Ubuntu-24.04 (/usr/local/cuda-12.*) and runs it three times. Every number is a RESULT line; the probe's own +# "RESULT FAMILY ..." lines carry kernel, best ms, G steps/s, ns per step, the ratio to the alu chain and ok=1 for +# bit-exact against the CPU reference on two whole warps. Read back with `node tools/jobs.mjs `. +$ErrorActionPreference = 'Continue' +$urlFile = if ($env:IGNEUM_APP_DIR) { Join-Path $env:IGNEUM_APP_DIR 'app.url' } else { Join-Path $env:LOCALAPPDATA 'igneum\app\app.url' } +if (-not (Test-Path $urlFile)) { $urlFile = Join-Path $env:LOCALAPPDATA 'igneum\app\app.url' } +$base = (Get-Content $urlFile -Raw).Trim().TrimEnd('/') +function Stamp { (Get-Date).ToUniversalTime().ToString('yyyy-MM-ddTHH:mm:ssZ') } +function Prove($on) { try { (Invoke-RestMethod -Method Post -Uri "$base/api/prove" -ContentType 'application/json' -Body (@{on=$on} | ConvertTo-Json -Compress) -TimeoutSec 10) | ConvertTo-Json -Compress } catch { "error: $_" } } +function Smi($q) { try { (& nvidia-smi --query-gpu=$q --format=csv,noheader,nounits 2>$null) -join ' | ' } catch { 'nvidia-smi failed' } } +function CudaApps { @(& nvidia-smi --query-compute-apps=pid,process_name,used_memory --format=csv,noheader 2>$null | ForEach-Object { "$_" } | Where-Object { $_ -match 'igneum|sp1|prove' }) } +"RESULT start $(Stamp) prover off for the run: $(Prove $false)" +Start-Sleep -Seconds 30 +$t = 0; $apps = @(CudaApps) +while ($t -lt 120 -and $apps.Count -gt 0) { Start-Sleep -Seconds 10; $t += 10; $apps = @(CudaApps) } +if ($apps.Count -eq 0) { "RESULT card $(Stamp) empty after $t s (no igneum, sp1 or prove compute app on the card)" } else { "RESULT card $(Stamp) UNCONFIRMED after $t s: the numbers below are beside " + ($apps -join '; ') } +"RESULT gpus $(Stamp) $(Smi 'index,name,driver_version,memory.used,clocks.sm,clocks.mem,power.draw,temperature.gpu')" +$job = $env:IGNEUM_JOB_DIR; if (-not $job) { $job = Join-Path $env:TEMP 'igneum-ca3-family' }; New-Item -ItemType Directory -Force -Path $job | Out-Null +$cu = @' +CU_PLACEHOLDER +'@ +$cuFile = Join-Path $job 'family-probe.cu' +[IO.File]::WriteAllText($cuFile, (($cu -replace "`r`n", "`n") + "`n"), (New-Object System.Text.UTF8Encoding $false)) +"RESULT source family-probe.cu sha256 $((Get-FileHash -Algorithm SHA256 $cuFile).Hash.ToLower()) bytes $((Get-Item $cuFile).Length)" +function WslPath($p) { $w = (& wsl.exe -d Ubuntu-24.04 -u root -- wslpath -a ($p -replace '\\', '/') 2>$null); if ($w) { ($w -replace "`0", '').Trim() } else { '/mnt/c' + ($p.Substring(2) -replace '\\', '/') } } +$bash = @' +set -uo pipefail +CUDA_DIR="$(ls -d /usr/local/cuda-12.* 2>/dev/null | sort -V | tail -1 || true)" +[ -n "$CUDA_DIR" ] || { echo "RESULT build_failed no /usr/local/cuda-12.*"; exit 2; } +export PATH="$CUDA_DIR/bin:$PATH" LD_LIBRARY_PATH="$CUDA_DIR/lib64:/usr/lib/wsl/lib:${LD_LIBRARY_PATH:-}" +stamp() { date -u +%Y-%m-%dT%H:%M:%SZ; } +SRC='SRC_PLACEHOLDER' +W=/tmp/igneum-ca3-family; rm -rf "$W"; mkdir -p "$W"; cp "$SRC" "$W/family-probe.cu"; cd "$W" +echo "RESULT nvcc $(nvcc --version | grep -o 'release [0-9.]*' | head -1) cuda_dir=$CUDA_DIR" +echo "RESULT source_wsl sha256=$(sha256sum family-probe.cu | cut -c1-64) bytes=$(stat -c %s family-probe.cu)" +if nvcc -O2 -arch=sm_120 -o family-probe family-probe.cu > build.log 2>&1; then echo "RESULT build ok arch=sm_120" +else + echo "RESULT build sm_120 failed: $(grep -m1 -i 'error' build.log | cut -c1-200)" + if nvcc -O2 -gencode arch=compute_89,code=compute_89 -o family-probe family-probe.cu > build2.log 2>&1; then echo "RESULT build ok arch=compute_89 (PTX, driver JIT)" + else echo "RESULT build_failed $(grep -m1 -i 'error' build2.log | cut -c1-200)"; exit 2; fi +fi +echo "RESULT binary sha256=$(sha256sum family-probe | cut -c1-16)" +for i in 1 2 3; do + echo "RESULT run $i start $(stamp) gpu_before $(nvidia-smi --query-gpu=clocks.sm,clocks.mem,power.draw,temperature.gpu --format=csv,noheader,nounits 2>/dev/null | head -1)" + ./family-probe --reps 3 2>&1 | sed "s/^RESULT /RESULT run=$i /; t; s/^/RESULT run=$i text /" + echo "RESULT run $i end $(stamp) exit=${PIPESTATUS[0]} gpu_after $(nvidia-smi --query-gpu=clocks.sm,clocks.mem,power.draw,temperature.gpu --format=csv,noheader,nounits 2>/dev/null | head -1)" +done +echo "RESULT measure_end $(stamp)" +'@ +$bash = $bash.Replace('SRC_PLACEHOLDER', (WslPath $cuFile)) +$bashFile = Join-Path $job 'measure.sh' +[IO.File]::WriteAllText($bashFile, ($bash -replace "`r`n", "`n"), (New-Object System.Text.UTF8Encoding $false)) +& wsl.exe -d Ubuntu-24.04 -u root -- bash (WslPath $bashFile) 2>&1 | ForEach-Object { ($_ -replace "`0", '') } +"RESULT gpus_after $(Stamp) $(Smi 'clocks.sm,clocks.mem,power.draw,temperature.gpu')" +"RESULT end $(Stamp) prover back on: $(Prove $true)" +exit 0