igneum/proto-cuda/CHECKLIST.md

# Metal to CUDA equivalence checklist

Written 3 October 2026 for the program packs in `packs/`. Everything below is about the kernel text that
`proto-metal/igneum-bench --export-pack` emits, compared with the Metal text the same run emits and with the
CPU interpreter `cpuWarp` in `proto-metal/main.swift`. There is no CUDA toolchain on the Mac, so "verified"
means one of the three methods in the last section, never an nvcc build.

## Types and arithmetic

| Item | Metal | CUDA | Status |
|---|---|---|---|
| Register type | `uint` (32-bit unsigned) | `uint32_t` (`unsigned int`, 32-bit on Linux and Windows) | same width, same wraparound |
| Output type | `ulong` | `uint64_t` | 64-bit on both; `(uint64_t)hi << 32 \| (uint64_t)lo` |
| Literals | `0x...u` | `0x...u` | emitted by the same `hex()` helper |
| Floats | none | none | nothing for fast-math or FMA contraction to touch |
| Shift amounts | always 0..31 (see rotates) | always 0..31 | no out-of-range shift anywhere |
| Evaluation order | each instruction is one statement with one assignment | same statement text | no sequence-point question arises |

## Register init and framing

| Item | Metal | CUDA | Status |
|---|---|---|---|
| Nonce | `baseNonce + gid`, gid = thread_position_in_grid | `baseNonce + blockIdx.x * blockDim.x + threadIdx.x` | identical for blockDim 32; also identical for blockDim 32 x W because the nonce is still base + global index |
| Register init | `x = nonce ^ SEEDW[i]; x += 0x9e3779b9u * (i+1)u; x = splitmix32(x); r_i = x ^ SEEDW[(i+1)&7]` | same, with `SEEDW[i]` and the product `0x9e3779b9 * (i+1)` written as literals (computed by the exporter mod 2^32) | verified by vectors |
| splitmix32 | 3 shift-xor, 2 multiply | identical text | verified by vectors |
| Iteration | `for it in 0..<8 { sel = r0; 64 instructions }` | identical | verified by vectors |
| Output mix | `lo = r0 ^ rotl(r1,7) ^ rotl(r2,14) ^ rotl(r3,21)`, `hi = r4 ^ rotl(r5,9) ^ rotl(r6,18) ^ rotl(r7,27)` | identical | verified by vectors |

## Instructions

| Op | Metal | CUDA | CPU interpreter | Notes |
|---|---|---|---|---|
| add | `d = d + a + select(imm, imm2, ((sel >> bit) & 1u) != 0u)` | `d = d + a + ((((sel >> bit) & 1u) != 0u) ? imm2 : imm)` | `d &+ a &+ (s != 0 ? imm2 : imm)` | Metal `select(A, B, c)` is `c ? B : A`, so the true branch is `imm2` in all three. `sel` is r0 sampled at the top of the iteration in all three. |
| sub | `d = d - a` | same | `d &- a` | wraparound |
| mul | `d = d * a` | same | `d &* a` | low 32 bits |
| mulhi | `mulhi(d, a)` | `__umulhi(d, a)` | `(UInt64(d) * UInt64(a)) >> 32` | high 32 bits of the unsigned 64-bit product on all three |
| xor | `d = d ^ a` | same | same | |
| or | `d = d \| a` | same | same | |
| rotl (immediate) | `rotl_imm(d, n)` = `(x << n) \| (x >> (32u - n))`, n literal 1..31 | identical text | `rotl32` with n in 1..31 | the generator draws `rot` from 1..31, so neither shift amount is ever 0 or 32 |
| rotr (register) | `rotr_var(d, a)` = `n &= 31u; (x >> n) \| (x << ((32u - n) & 31u))` | identical text | `rotr32`: `n & 31`, `n == 0 ? x : ...` | for n = 0 both GPU forms give `x \| x = x`; for n in 1..31 both shifts are in range |
| mad | `d = a * b + d` | same | `(a &* r[b]) &+ d` | `b` may equal `d` or `a`; all three read every operand before the single write |
| shfl | `d = d ^ simd_shuffle_xor(a, (ushort)mask)` | `d = d ^ __shfl_xor_sync(0xffffffffu, a, mask)` | `r[lane][dst] ^= tmp[lane ^ mask]` where `tmp` is a copy of register `a` across the warp taken before any lane writes | mask in {1,2,4,8,16} so the partner lane is always inside the same 32-lane group. `a != dst` by construction, so there is no read-after-write hazard inside the instruction. Control flow is uniform (no branches at all), so the full member mask is valid and no lane is missing from the shuffle. |
| load | `d = d ^ dataset[a & MASK]`, MASK a compile-time literal | `d = d ^ ds[a & mask]`, mask a kernel argument | `d ^ datasetElem(a & mask, d0, d1)` | the index is a 32-bit unsigned value below 2^28, so pointer arithmetic needs no 64-bit care |

No instruction had to be removed or changed. Every op is an unsigned 32-bit integer operation whose result is
defined identically in Metal Shading Language, CUDA C++ and Swift's wrapping operators.

## Lane mapping (the one place where the two platforms differ in guarantees)

| Item | Metal | CUDA |
|---|---|---|
| Group that `shuffle_xor` spans | the SIMD group, width `threadExecutionWidth` (32 on this M5 Max; the tool warns if not) | the warp, width 32 on every NVIDIA GPU (`host.cu` prints `cudaDevAttrWarpSize` and warns if not 32) |
| Lane id of a thread | `thread_index_in_simdgroup`; for a 32-wide threadgroup this equalled `thread_position_in_threadgroup`, as shown by the bit-exact match with the CPU model across 21 warps on 3 Oct 2026 and the 3 vector warps per pack today | `threadIdx.x % 32`, guaranteed by the CUDA programming model |
| Partner lane | `lane ^ mask` | `lane ^ mask` |
| Warps per block | 1 (threadgroup of 32) | `--block-warps W`, default 1. Any W is bit-exact because each warp is an aligned run of 32 consecutive nonces either way. The emulation passed with W = 1 and W = 2. |

## Dataset

| Item | Metal | CUDA |
|---|---|---|
| Element | `ds_elem(i, d0, d1)`: xor, mul 0x9E3779B1, xor-shift 15, add d1, mul 0x85EBCA77, xor-shift 13, mul 0xC2B2AE3D, xor-shift 16 | identical text in `kernel.cu`; a third copy `host_ds_elem` in `host.cu` |
| Day words | `seedWords("day/" + day)[0..1]` on the Mac | written into `program.h` as `IGNEUM_DAY0`, `IGNEUM_DAY1` |
| Fill | one thread per word, threadgroup 256 | one thread per word, block 256, with an `i < n` guard (a no-op for power-of-two sizes) |
| Self-test | not needed on the Mac (CPU and GPU share one process) | head 16 words and word `[MASK]` against values the Mac wrote into `vectors.h`; 64 pseudo-random words against `host_ds_elem` |

## Dataset, memory-hard pack (added later on 3 October 2026)

Pack `igneum-genesis-mh`, `IGNEUM_DATASET_MODE 1`. Construction in `../proto-metal/MEMHARD.md`.

| Item | Metal (Mac) | CUDA pack | Host reference (host.cu) | CPU verifier (main.swift) |
|---|---|---|---|---|
| Core text (ChaCha12 block, cache segment chain, mixer M_r, item derivation) | `emitMemhardCore(cuda: false)`, constants as literals | `memhard.h`: `emitMemhardCore(cuda: true)`, the same emitter, `IGNEUM_HD` = `__host__ __device__` | the same `memhard.h`, compiled as plain C++ (`IGNEUM_HD` = `static inline`) | separate Swift implementation (`chachaBlock`, `cpuFillSegment`, `mixer`, `deriveItems`) |
| Cache fill | `igneum_cache_fill`, one thread per segment | `igneum_cache_fill<<<256, 256>>>`, `if (seg < nSegments)` | `mh_cache_segment` for all 65,536 segments on one thread | `cpuFillCache`, one core |
| Dataset build | `igneum_build`, one thread per item, writes 16 words | `igneum_build<<<items/256, 256>>>`, `if (t < nItems)`, writes through a `d` pointer (no `ds[` write, so the static mask check expects 0 fill writes) | not built; words derived on demand with `mh_word` | not built; `MemhardCPU.fetch` |
| Agreement | GPU cache == CPU cache on all 2^26 words in every run; 1,024 sampled GPU dataset words == CPU derivation; 21 bench warps, 800 fuzz warps, edge, determinism, memcheck PASS | emulation: emulated-GPU cache == host cache on all words, host FNV == Mac FNV, dataset self-test PASS at 1 GiB, 3/3 vectors standalone, 2/2 in batch at 2 warps per block | | |

Residual risk specific to this pack: `mh_item` reads a cache line as 16 scalar `uint32_t` loads through a `const uint32_t*`;
nvcc may or may not vectorise them. That affects speed on NVIDIA, not results. The 256 MiB cache plus 1 GiB dataset
need 1.25 GiB of device memory plus the output buffer.

## How each claim above was verified

1. Mac, Metal GPU against the CPU interpreter: the regular bench (`proto-metal/igneum-bench`) passed 3 warps per
   program on both seeds again after the exporter was added, and `--export-pack` itself runs the Metal kernel for
   the three vector warps (base nonces 0, 4096, 1000000) and refuses to write a pack unless all 96 outputs match
   the interpreter. Both packs were written with "Metal GPU cross-check PASS 3/3 warps".
2. Mac, CUDA text executed as C++: `emu/emu.sh` compiles `host.cu` and the generated `kernel.cu` with clang
   against a shim `cuda_runtime.h` (the launch syntax is the only text rewritten) and runs the kernels on host
   threads, 32 per warp, with a barrier inside `__shfl_xor_sync`. Result on 3 Oct 2026, both packs, 1 GiB
   dataset: dataset self-test PASS, 3 of 3 warps PASS standalone, warps 0 and 4096 PASS inside a batch with
   1 and with 2 warps per block. Sweep sizes 4, 64, 256, 512 MiB: dataset self-test PASS. `-Wall -Wextra` clean.
3. Reading: every emitted line was compared by eye against the Metal line for the same instruction index
   (the comment at the end of each CUDA line carries the index and the op).

## Residual risks that the Mac cannot remove

- nvcc never ran on this code. Syntax that clang accepted could still trip nvcc's front end, and the shim's
  idea of the runtime API could differ from the real header in a detail. The six `cudaDevAttr*` enumerators
  and the template signatures of `cudaFuncGetAttributes(cudaFuncAttributes*, T* entry)` and
  `cudaOccupancyMaxActiveBlocksPerMultiprocessor(int*, T func, int blockSize, size_t dynamicSMemSize)` were
  checked against the NVIDIA CUDA Runtime API documentation on 3 Oct 2026 and match what the code uses.
  Anything left would be a compile error, not a silent output difference, and a one-line fix in `host.cu` or
  the three wrapper functions at the bottom of `kernel.cu`.
- `__shfl_xor_sync` and `__umulhi` semantics were emulated, not executed on NVIDIA hardware. Both are
  documented as lane ^ laneMask within a 32-wide warp and the high 32 bits of the unsigned product; the vectors
  will settle it on the 5090.
- `-arch=sm_120` requires CUDA 12.8 or newer. On an older toolkit the build fails at the flag, not at the code.
- Performance figures from the emulation are meaningless and were not recorded. No NVIDIA hash rate exists yet.