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
This commit is contained in:
igneum-josh
parent
dda9fa3fcf
commit
192a683d08
5 changed files with 501 additions and 0 deletions
|
|
@ -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.
|
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.
|
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 |
|
||||||
|
|
|
||||||
179
proto-cuda/family-probe.cu
Normal file
179
proto-cuda/family-probe.cu
Normal file
|
|
@ -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 <cuda_runtime.h>
|
||||||
|
#include <cstdio>
|
||||||
|
#include <cstdlib>
|
||||||
|
#include <cstring>
|
||||||
|
#include <cstdint>
|
||||||
|
|
||||||
|
__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]<<<lanes / 256, 256>>>(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;
|
||||||
|
}
|
||||||
199
proto-metal/family-probe.swift
Normal file
199
proto-metal/family-probe.swift
Normal file
|
|
@ -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 <metal_stdlib>
|
||||||
|
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<char4>(a)); int4 vb = int4(as_type<char4>(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<uchar4>(a)); uint4 vb = uint4(as_type<uchar4>(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<uchar4>(y); acc = as_type<uint>(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..<steps { for l in 0..<32 { x[l] = x[l] &* 0x9E3779B1 &+ rotl(y[l], 7); y[l] = (y[l] ^ x[l]) &+ s } }
|
||||||
|
return (0..<32).map { x[$0] ^ y[$0] }
|
||||||
|
}
|
||||||
|
for s in 0..<steps {
|
||||||
|
let xs = x // shuffles read the other lanes' x of this step
|
||||||
|
for l in 0..<32 {
|
||||||
|
let xv = x[l], yv = y[l]
|
||||||
|
switch kernel {
|
||||||
|
case "probe_rotr": acc[l] = rotrVar(yv, xv)
|
||||||
|
case "probe_shflx": acc[l] = acc[l] ^ xs[l ^ 8]
|
||||||
|
case "probe_shl": acc[l] = yv << (xv & 31)
|
||||||
|
case "probe_shr": acc[l] = yv >> (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..<reps {
|
||||||
|
var st = steps
|
||||||
|
var seed = UInt32(0x2468ace) &+ UInt32(r) &* 0x9E3779B9
|
||||||
|
let cb = queue.makeCommandBuffer()!
|
||||||
|
let enc = cb.makeComputeCommandEncoder()!
|
||||||
|
enc.setComputePipelineState(pso)
|
||||||
|
enc.setBytes(&st, length: 4, index: 0)
|
||||||
|
enc.setBytes(&seed, length: 4, index: 1)
|
||||||
|
enc.setBuffer(outBuf, offset: 0, index: 2)
|
||||||
|
enc.dispatchThreads(MTLSize(width: lanes, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: tg, height: 1, depth: 1))
|
||||||
|
enc.endEncoding()
|
||||||
|
cb.commit()
|
||||||
|
cb.waitUntilCompleted()
|
||||||
|
let ms = (cb.gpuEndTime - cb.gpuStartTime) * 1000.0
|
||||||
|
if ms < best { best = ms }
|
||||||
|
let p = outBuf.contents().bindMemory(to: UInt32.self, capacity: lanes)
|
||||||
|
for g0 in [UInt32(0), UInt32(lanes - 32)] {
|
||||||
|
let want = warpRef(kernel: name, g0: g0, seed: seed, steps: steps)
|
||||||
|
for l in 0..<32 where p[Int(g0) + l] != want[l] {
|
||||||
|
okAll = false
|
||||||
|
print("MISMATCH \(name) lane \(g0 + UInt32(l)): gpu \(String(p[Int(g0) + l], radix: 16)) cpu \(String(want[l], radix: 16))")
|
||||||
|
break
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
if name == "probe_alu" { aluBest = best }
|
||||||
|
let stepsPerS = Double(lanes) * Double(steps) / (best / 1000.0)
|
||||||
|
let ratio = aluBest > 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")
|
||||||
20
tools/ca3-reserve/make-pc2-playbook.sh
Executable file
20
tools/ca3-reserve/make-pc2-playbook.sh
Executable file
|
|
@ -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 ' ')"
|
||||||
63
tools/ca3-reserve/pc2-family-probe.ps1
Normal file
63
tools/ca3-reserve/pc2-family-probe.ps1
Normal file
|
|
@ -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 <job id>`.
|
||||||
|
$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
|
||||||
Loading…
Reference in a new issue