From c3da619571f44174bbd48829218442e31b1aa10c Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Sat, 3 Oct 2026 15:49:46 +0000 Subject: [PATCH] Domains attached and nameservers moved to Vercel; project notes Co-Authored-By: Claude Fable 5.1 --- proto-metal/main.swift | 278 +++++++++++++++++++++++++++++++++++++---- 1 file changed, 252 insertions(+), 26 deletions(-) diff --git a/proto-metal/main.swift b/proto-metal/main.swift index 877489172..d8409cfb5 100644 --- a/proto-metal/main.swift +++ b/proto-metal/main.swift @@ -369,14 +369,38 @@ struct Program { let opWeights: [(Op, Int)] = [(.load, 25), (.add, 12), (.xor, 10), (.mul, 8), (.mad, 8), (.shfl, 8), (.rotl, 7), (.sub, 6), (.mulhi, 6), (.rotr, 6), (.or, 4)] +// Generator levers (3 October 2026). The defaults reproduce the original generator instruction for instruction: +// with loadWeight 25 the weight table above is used unchanged, and with wideFrac 0 no load becomes a wload. +// Neither lever consumes extra random draws, so a program differs from the default one only where the lever acts. +struct GeneratorConfig { + var loadWeight = 25 + var wideFrac = 0 + // Scaled weights: load gets loadWeight, the other ten ops share the rest in their original proportions, + // rounded by largest remainder so the table still sums to 100. + var weights: [(Op, Int)] { + if loadWeight == 25 { return opWeights } + let others = opWeights.dropFirst() + let total = others.reduce(0) { $0 + $1.1 } // 75 + let budget = 100 - loadWeight + var scaled = others.map { (op: $0.0, floor: ($0.1 * budget) / total, rem: ($0.1 * budget) % total) } + var sum = scaled.reduce(0) { $0 + $1.floor } + let order = scaled.indices.sorted { scaled[$0].rem != scaled[$1].rem ? scaled[$0].rem > scaled[$1].rem : $0 < $1 } + var k = 0 + while sum < budget { scaled[order[k]].floor += 1; sum += 1; k += 1 } + return [(.load, loadWeight)] + scaled.map { ($0.op, $0.floor) } + } +} +var generatorConfig = GeneratorConfig() + func generateProgram(seedString: String) -> Program { let sw = seedWords(seedString) var rng = SplitMix64(s: (UInt64(sw[0]) | (UInt64(sw[1]) << 32)) ^ ((UInt64(sw[2]) | (UInt64(sw[3]) << 32)) &* 0x9E3779B97F4A7C15)) var instrs = [Instr]() + let weights = generatorConfig.weights for _ in 0..= dst { a += 1 } let b = rng.below(8) @@ -385,6 +409,8 @@ func generateProgram(seedString: String) -> Program { let rot = UInt32(1 + rng.below(31)) let bit = rng.below(32) let mask = 1 << rng.below(5) + // Lever (b): the already-drawn selector bit decides whether a load is wide, so the draw stream is unchanged. + if op == .load && bit * 100 < generatorConfig.wideFrac * 32 { op = .wload } instrs.append(Instr(op: op, dst: dst, a: a, b: b, imm: imm, imm2: imm2, rot: rot, bit: bit, mask: mask)) } return Program(seedString: seedString, seed: sw, instrs: instrs) @@ -394,9 +420,119 @@ func generateProgram(seedString: String) -> Program { func hex(_ v: UInt32) -> String { String(format: "0x%08xu", v) } -// inlineDay: when set, every load computes ds_elem(index, day) in registers instead of reading the buffer. -// Same function, no memory traffic. Used only by --inline-dataset to measure the closed-form shortcut. -func generateMSL(_ p: Program, datasetLog2: Int, inlineDay: (UInt32, UInt32)? = nil) -> String { +// The memory-hard core as source text, in Metal (cuda false) or CUDA C++ (cuda true). Mixer parameters are +// literals so the GPU kernels, the CUDA pack and its host reference share one text. Names are prefixed mh_. +// In the CUDA dialect every function is IGNEUM_HD (host and device) so host.cu can derive items too. +func emitMemhardCore(_ mp: MixParams, cuda: Bool) -> String { + let U = cuda ? "uint32_t" : "uint" + let fn = cuda ? "IGNEUM_HD" : "inline" + let cptr = cuda ? "const uint32_t*" : "device const uint*" + let wptr = cuda ? "uint32_t*" : "device uint*" + let lptr = cuda ? "uint32_t*" : "thread uint*" + let lcptr = cuda ? "const uint32_t*" : "const thread uint*" + let K = mp.keyWords, R = mp.rotWords, M = mp.mulWords, C = mp.rcWords + var s = """ + // Memory-hard dataset core (MEMHARD.md). Cache: 2^\(cacheLog2Words) words in 2^\(Int(log2(Double(cacheSegments)))) segments of \(cacheLinesPerSegment) chained ChaCha\(chachaRounds) lines. + // Item: 8 rounds of seed-parameterised mixer + one 64-byte cache read, then a final mixer. All parameters are literals. + #define MH_CACHE_LINE_MASK \(hex(cacheLineMask)) + #define MH_SEGMENT_LINES \(cacheLinesPerSegment)u + #define MH_QR(a, b, c, d, r1, r2, r3, r4) { a += b; d ^= a; d = mh_rotl(d, r1); c += d; b ^= c; b = mh_rotl(b, r2); a += b; d ^= a; d = mh_rotl(d, r3); c += d; b ^= c; b = mh_rotl(b, r4); } + \(fn) \(U) mh_rotl(\(U) x, \(U) n) { return (x << n) | (x >> (32u - n)); } // n in 1..31 at every call site + + // y = ChaCha\(chachaRounds) core(x) + x + \(fn) void mh_chacha_block(\(lcptr) x, \(lptr) y) { + for (\(U) i = 0u; i < 16u; ++i) y[i] = x[i]; + for (\(U) r = 0u; r < \(chachaRounds / 2)u; ++r) { + MH_QR(y[0], y[4], y[8], y[12], 16u, 12u, 8u, 7u) MH_QR(y[1], y[5], y[9], y[13], 16u, 12u, 8u, 7u) + MH_QR(y[2], y[6], y[10], y[14], 16u, 12u, 8u, 7u) MH_QR(y[3], y[7], y[11], y[15], 16u, 12u, 8u, 7u) + MH_QR(y[0], y[5], y[10], y[15], 16u, 12u, 8u, 7u) MH_QR(y[1], y[6], y[11], y[12], 16u, 12u, 8u, 7u) + MH_QR(y[2], y[7], y[8], y[13], 16u, 12u, 8u, 7u) MH_QR(y[3], y[4], y[9], y[14], 16u, 12u, 8u, 7u) + } + for (\(U) i = 0u; i < 16u; ++i) y[i] += x[i]; + } + + // One cache segment: \(cacheLinesPerSegment) chained lines written at cache[seg * \(cacheLinesPerSegment * 16)]. in_j = prev ^ (sigma || K || seg || j || tag), prev_0 = 0. + \(fn) void mh_cache_segment(\(wptr) cache, \(U) seg) { + \(U) prev[16]; \(U) x[16]; \(U) y[16]; + for (\(U) i = 0u; i < 16u; ++i) prev[i] = 0u; + for (\(U) j = 0u; j < MH_SEGMENT_LINES; ++j) { + x[0] = \(hex(chachaSigma[0])) ^ prev[0]; x[1] = \(hex(chachaSigma[1])) ^ prev[1]; x[2] = \(hex(chachaSigma[2])) ^ prev[2]; x[3] = \(hex(chachaSigma[3])) ^ prev[3]; + + """ + for i in 0..<8 { s += " x[\(4 + i)] = \(hex(K[i])) ^ prev[\(4 + i)];\n" } + s += """ + x[12] = seg ^ prev[12]; x[13] = j ^ prev[13]; x[14] = \(hex(cacheTag[0])) ^ prev[14]; x[15] = \(hex(cacheTag[1])) ^ prev[15]; + mh_chacha_block(x, y); + \(wptr) line = cache + ((seg * MH_SEGMENT_LINES + j) * 16u); + for (\(U) i = 0u; i < 16u; ++i) { line[i] = y[i]; prev[i] = y[i]; } + } + } + + // M_r: per word (s ^ (RC + rk)) * MUL, then a column round and a diagonal round with the seed-drawn rotations. + \(fn) void mh_mixer(\(lptr) s, \(U) rk) { + + """ + for i in 0..<16 { s += " s[\(i)] = (s[\(i)] ^ (\(hex(C[i])) + rk)) * \(hex(M[i]));\n" } + let c = (0..<4).map { "\(R[$0])u" }.joined(separator: ", "), d = (4..<8).map { "\(R[$0])u" }.joined(separator: ", ") + s += """ + MH_QR(s[0], s[4], s[8], s[12], \(c)) MH_QR(s[1], s[5], s[9], s[13], \(c)) + MH_QR(s[2], s[6], s[10], s[14], \(c)) MH_QR(s[3], s[7], s[11], s[15], \(c)) + MH_QR(s[0], s[5], s[10], s[15], \(d)) MH_QR(s[1], s[6], s[11], s[12], \(d)) + MH_QR(s[2], s[7], s[8], s[13], \(d)) MH_QR(s[3], s[4], s[9], s[14], \(d)) + } + + // Item t: 16 words. s = (K, t * MUL[i] + RC[i]); \(itemRounds) rounds of mixer + cache line s[0] & mask; final mixer. + \(fn) void mh_item(\(cptr) cache, \(U) t, \(lptr) s) { + + """ + for i in 0..<8 { s += " s[\(i)] = \(hex(K[i]));\n" } + for i in 0..<8 { s += " s[\(8 + i)] = t * \(hex(M[i])) + \(hex(C[i]));\n" } + s += """ + for (\(U) r = 0u; r < \(itemRounds)u; ++r) { + mh_mixer(s, 0x9E3779B9u * (r + 1u)); + \(cptr) line = cache + ((s[0] & MH_CACHE_LINE_MASK) * 16u); + for (\(U) i = 0u; i < 16u; ++i) s[i] ^= line[i]; + } + mh_mixer(s, 0x9E3779B9u * \(itemRounds + 1)u); + } + // dataset[w] without the dataset: derive item w >> 4 and take word w & 15. + \(fn) \(U) mh_word(\(cptr) cache, \(U) w) { \(U) s[16]; mh_item(cache, w >> 4u, s); return s[w & 15u]; } + + """ + return s +} + +// Metal library with the cache fill and dataset build kernels for one day key. +func memhardMSL(_ mp: MixParams) -> String { + return """ + #include + using namespace metal; + \(emitMemhardCore(mp, cuda: false)) + // One thread per segment (2^\(Int(log2(Double(cacheSegments)))) threads). + kernel void igneum_cache_fill(device uint* cache [[buffer(0)]], uint gid [[thread_position_in_grid]]) { + mh_cache_segment(cache, gid); + } + // One thread per 64-byte item (dataset words / 16 threads). + kernel void igneum_build(device const uint* cache [[buffer(0)]], device uint* dataset [[buffer(1)]], + uint gid [[thread_position_in_grid]]) { + uint s[16]; + mh_item(cache, gid, s); + device uint* d = dataset + gid * 16u; + for (uint i = 0u; i < 16u; ++i) d[i] = s[i]; + } + + """ +} + +// How the hash kernel gets dataset words. .stored reads the buffer (the honest kernel). The two inline +// variants are the shortcut measurements: every load recomputes the word instead of reading the dataset. +enum LoadSource { + case stored + case inlineClosed(UInt32, UInt32) // closed form: ds_elem(index, d0, d1), no memory read at all + case inlineMemhard(MixParams) // memory-hard: mh_word(cache, index), 8 dependent cache reads per word +} + +func generateMSL(_ p: Program, datasetLog2: Int, source: LoadSource = .stored) -> String { let mask = UInt32((1 << datasetLog2) - 1) var s = """ #include @@ -422,7 +558,15 @@ func generateMSL(_ p: Program, datasetLog2: Int, inlineDay: (UInt32, UInt32)? = return x; } - kernel void igneum_hash(device const uint* dataset [[buffer(0)]], + """ + if p.hasWide { s += " #define WMASK (MASK & ~31u)\n\n" } + var buffer0 = "device const uint* dataset [[buffer(0)]]" + if case .inlineMemhard(let mp) = source { + s += emitMemhardCore(mp, cuda: false) + "\n" + buffer0 = "device const uint* cache [[buffer(0)]]" + } + s += """ + kernel void igneum_hash(\(buffer0), device ulong* out [[buffer(1)]], constant uint& baseNonce [[buffer(2)]], uint gid [[thread_position_in_grid]]) { @@ -430,10 +574,20 @@ func generateMSL(_ p: Program, datasetLog2: Int, inlineDay: (UInt32, UInt32)? = uint r0, r1, r2, r3, r4, r5, r6, r7; """ + if p.hasWide { s += " uint lane = gid & 31u;\n" } for i in 0..<8 { s += " { uint x = nonce ^ SEEDW[\(i)]; x += 0x9e3779b9u * \(i + 1)u; x = splitmix32(x); r\(i) = x ^ SEEDW[\((i + 1) & 7)]; }\n" } s += "\n for (uint it = 0u; it < \(Program.iterations)u; ++it) {\n uint sel = r0;\n" + // The word index expression for a load: plain = a & MASK; wide = lane 0's a, aligned to 32 words, plus lane. + func wordIndex(_ a: String, wide: Bool) -> String { wide ? "(simd_broadcast(\(a), 0) & WMASK) + lane" : "\(a) & MASK" } + func fetch(_ idx: String) -> String { + switch source { + case .stored: return "dataset[\(idx)]" + case .inlineClosed(let d0, let d1): return "ds_elem(\(idx), \(hex(d0)), \(hex(d1)))" + case .inlineMemhard: return "mh_word(cache, \(idx))" + } + } for (k, ins) in p.instrs.enumerated() { let d = "r\(ins.dst)", a = "r\(ins.a)", b = "r\(ins.b)" var line: String @@ -448,9 +602,8 @@ func generateMSL(_ p: Program, datasetLog2: Int, inlineDay: (UInt32, UInt32)? = case .rotr: line = "\(d) = rotr_var(\(d), \(a));" case .mad: line = "\(d) = \(a) * \(b) + \(d);" case .shfl: line = "\(d) = \(d) ^ simd_shuffle_xor(\(a), (ushort)\(ins.mask));" - case .load: - if let dd = inlineDay { line = "\(d) = \(d) ^ ds_elem(\(a) & MASK, \(hex(dd.0)), \(hex(dd.1)));" } - else { line = "\(d) = \(d) ^ dataset[\(a) & MASK];" } + case .load: line = "\(d) = \(d) ^ \(fetch(wordIndex(a, wide: false)));" + case .wload: line = "\(d) = \(d) ^ \(fetch(wordIndex(a, wide: true)));" } s += " \(line) // \(k)\n" } @@ -485,16 +638,20 @@ kernel void igneum_fill(device uint* dataset [[buffer(0)]], // MARK: - CPU reference interpreter for one 32-lane warp -func cpuWarp(_ p: Program, baseNonce: UInt32, day: (UInt32, UInt32), mask: UInt32) -> [UInt64] { - cpuWarpTraced(p, baseNonce: baseNonce, day: day, mask: mask, trace: nil) +func cpuWarp(_ p: Program, baseNonce: UInt32, ds: DatasetSource) -> [UInt64] { + cpuWarpTraced(p, baseNonce: baseNonce, ds: ds, trace: nil) } // Same interpreter with an optional hook. When `trace` is set it is called before every instruction with // (iteration, instruction index, the 8 registers of lane 0). The edge-case tests use it to prove that the // operand values they were built to produce really occurred. The bench passes nil. -func cpuWarpTraced(_ p: Program, baseNonce: UInt32, day: (UInt32, UInt32), mask: UInt32, +// Loads are batched across the 32 lanes: the indices are gathered, DatasetSource.fetch answers all 32, then +// the xors are applied. For the closed form this is the old per-lane formula; for the memory-hard dataset it +// lets the 32 independent item derivations overlap their cache misses (see deriveItems). +func cpuWarpTraced(_ p: Program, baseNonce: UInt32, ds: DatasetSource, trace: ((Int, Int, [UInt32]) -> Void)?) -> [UInt64] { let lanes = 32 + let mask = ds.mask var r = [UInt32](repeating: 0, count: lanes * 8) // r[lane*8 + reg] for lane in 0...allocate(capacity: lanes) + let val = UnsafeMutablePointer.allocate(capacity: lanes) + defer { idx.deallocate(); val.deallocate() } for it in 0.. String { return o + "\"" } -func generateCUDA(_ p: Program) -> String { +// memhard: nil for a closed-form pack (the original fill kernel), MixParams for the memory-hard pack +// (cache fill and dataset build kernels; the shared core comes from memhard.h in the same pack). +func generateCUDA(_ p: Program, memhard: MixParams?) -> String { var s = """ // Generated by proto-metal/igneum-bench --export-pack for seed "\(p.seedString)". Do not edit by hand. // Bit-exact twin of the Metal kernel for the same seed (see proto-cuda/CHECKLIST.md and program.metal). @@ -584,6 +755,7 @@ func generateCUDA(_ p: Program) -> String { #include #include #include "program.h" + \(memhard != nil ? "#include \"memhard.h\"" : "") __device__ __forceinline__ uint32_t splitmix32(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; @@ -604,12 +776,38 @@ func generateCUDA(_ p: Program) -> String { return x; } - // dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel. - __global__ void igneum_fill(uint32_t* ds, uint32_t n, uint32_t d0, uint32_t d1) { - uint32_t i = blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) ds[i] = ds_elem(i, d0, d1); - } + """ + if memhard == nil { + s += """ + // dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel. + __global__ void igneum_fill(uint32_t* ds, uint32_t n, uint32_t d0, uint32_t d1) { + uint32_t i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) ds[i] = ds_elem(i, d0, d1); + } + + """ + } else { + s += """ + // Memory-hard dataset (MEMHARD.md). One thread per cache segment; one thread per 64-byte dataset item. + // The core functions (mh_cache_segment, mh_item) are in memhard.h and are also compiled for the host. + __global__ void igneum_cache_fill(uint32_t* cache, uint32_t nSegments) { + uint32_t seg = blockIdx.x * blockDim.x + threadIdx.x; + if (seg < nSegments) mh_cache_segment(cache, seg); + } + __global__ void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems) { + uint32_t t = blockIdx.x * blockDim.x + threadIdx.x; + if (t < nItems) { + uint32_t s[16]; + mh_item(cache, t, s); + uint32_t* d = ds + (size_t)t * 16u; + for (uint32_t i = 0u; i < 16u; ++i) d[i] = s[i]; + } + } + + """ + } + s += """ // One hash per thread. blockDim.x is a multiple of 32; lane = threadIdx.x & 31 and every // __shfl_xor_sync stays inside the lane's own warp, exactly like simd_shuffle_xor inside a // 32-wide Metal SIMD group. Control flow is uniform, so the full 0xffffffff member mask is valid. @@ -619,6 +817,7 @@ func generateCUDA(_ p: Program) -> String { uint32_t r0, r1, r2, r3, r4, r5, r6, r7; """ + if p.hasWide { s += " uint32_t lane = threadIdx.x & 31u;\n uint32_t wmask = mask & ~31u;\n" } for i in 0..<8 { let addc = 0x9e3779b9 &* UInt32(i + 1) s += " { uint32_t x = nonce ^ \(hex(p.seed[i])); x += \(hex(addc)); x = splitmix32(x); r\(i) = x ^ \(hex(p.seed[(i + 1) & 7])); } // SEEDW[\(i)], 0x9e3779b9u * \(i + 1)u, SEEDW[\((i + 1) & 7)]\n" @@ -640,6 +839,7 @@ func generateCUDA(_ p: Program) -> String { case .mad: line = "\(d) = \(a) * \(b) + \(d);" case .shfl: line = "\(d) = \(d) ^ __shfl_xor_sync(0xffffffffu, \(a), \(ins.mask));" case .load: line = "\(d) = \(d) ^ ds[\(a) & mask];" + case .wload: line = "\(d) = \(d) ^ ds[(__shfl_sync(0xffffffffu, \(a), 0) & wmask) + lane];" } s += " \(line) // \(k) \(ins.op.rawValue)\n" } @@ -651,14 +851,40 @@ func generateCUDA(_ p: Program) -> String { } // Host-side launch wrappers. Declared in program.h, called from host.cu. - cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1) { - if (nWords == 0u) return cudaErrorInvalidValue; - uint32_t block = 256u; - uint32_t grid = (nWords + block - 1u) / block; - igneum_fill<<>>(ds, nWords, d0, d1); - return cudaGetLastError(); - } + """ + if memhard == nil { + s += """ + cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1) { + if (nWords == 0u) return cudaErrorInvalidValue; + uint32_t block = 256u; + uint32_t grid = (nWords + block - 1u) / block; + igneum_fill<<>>(ds, nWords, d0, d1); + return cudaGetLastError(); + } + + """ + } else { + s += """ + cudaError_t igneum_launch_cache_fill(uint32_t* cache, uint32_t nSegments) { + if (nSegments == 0u) return cudaErrorInvalidValue; + uint32_t block = 256u; + uint32_t grid = (nSegments + block - 1u) / block; + igneum_cache_fill<<>>(cache, nSegments); + return cudaGetLastError(); + } + + cudaError_t igneum_launch_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems) { + if (nItems == 0u) return cudaErrorInvalidValue; + uint32_t block = 256u; + uint32_t grid = (nItems + block - 1u) / block; + igneum_build<<>>(ds, cache, nItems); + return cudaGetLastError(); + } + + """ + } + s += """ cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, uint32_t nonces, uint32_t blockWarps) { if (blockWarps == 0u || blockWarps > 32u) return cudaErrorInvalidValue;