igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity

Domains attached and nameservers moved to Vercel; project notes

Browse source 
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs 2026-10-03 15:49:46 +00:00
parent 105354a2ce
commit c3da619571
1 changed files with 252 additions and 26 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

278
proto-metal/main.swift
View file

@ -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..<Program.count {
var roll = rng.below(100)
var op = Op.add
for (o, w) in opWeights { if roll < w { op = o; break }; roll -= w }
for (o, w) in weights { if roll < w { op = o; break }; roll -= w }
let dst = rng.below(8)
var a = rng.below(7); if a >= 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 <metal_stdlib>
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 <metal_stdlib>
@ -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..<lanes {
let nonce = baseNonce &+ UInt32(lane)
@ -506,6 +663,9 @@ func cpuWarpTraced(_ p: Program, baseNonce: UInt32, day: (UInt32, UInt32), mask:
}
}
var tmp = [UInt32](repeating: 0, count: lanes)
let idx = UnsafeMutablePointer<UInt32>.allocate(capacity: lanes)
let val = UnsafeMutablePointer<UInt32>.allocate(capacity: lanes)
defer { idx.deallocate(); val.deallocate() }
for it in 0..<Program.iterations {
for lane in 0..<lanes { tmp[lane] = r[lane * 8] } // sel = r0 at the top of the iteration
let sel = tmp
@ -515,6 +675,16 @@ func cpuWarpTraced(_ p: Program, baseNonce: UInt32, day: (UInt32, UInt32), mask:
case .shfl:
for lane in 0..<lanes { tmp[lane] = r[lane * 8 + ins.a] }
for lane in 0..<lanes { r[lane * 8 + ins.dst] ^= tmp[lane ^ ins.mask] }
case .load:
for lane in 0..<lanes { idx[lane] = r[lane * 8 + ins.a] & mask }
ds.fetch(UnsafePointer(idx), lanes, val)
for lane in 0..<lanes { r[lane * 8 + ins.dst] ^= val[lane] }
case .wload:
// Lane 0's register, masked, aligned down to 32 words; lane l reads word base + l.
let base = (r[ins.a] & mask) & ~31
for lane in 0..<lanes { idx[lane] = base + UInt32(lane) }
ds.fetch(UnsafePointer(idx), lanes, val)
for lane in 0..<lanes { r[lane * 8 + ins.dst] ^= val[lane] }
default:
for lane in 0..<lanes {
let base = lane * 8
@ -532,8 +702,7 @@ func cpuWarpTraced(_ p: Program, baseNonce: UInt32, day: (UInt32, UInt32), mask:
case .rotl: v = rotl32(d, ins.rot)
case .rotr: v = rotr32(d, a)
case .mad: v = (a &* r[base + ins.b]) &+ d
case .load: v = d ^ datasetElem(a & mask, day.0, day.1)
case .shfl: v = d // unreachable
case .load, .wload, .shfl: v = d // unreachable, handled above
}
r[base + ins.dst] = v
}
@ -576,7 +745,9 @@ func jstr(_ s: String) -> 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 <cuda_runtime.h>
#include <cstdint>
#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<<<grid, block>>>(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<<<grid, block>>>(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<<<grid, block>>>(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<<<grid, block>>>(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;