bump version to v4.0.2 (#93)

Co-authored-by: n13 <n13@users.noreply.github.com>

.github/workflows/ci.yml

@@ -86,6 +86,11 @@ jobs:
       - name: build benchmark binary
         run: cargo build -p miner-cli --release
         timeout-minutes: 60
+      - name: upload linux release binary
+        uses: actions/upload-artifact@v4
+        with:
+          name: quantus-miner-linux-x86_64
+          path: target/release/quantus-miner
       - name: run cpu benchmark
         run: ./target/release/quantus-miner benchmark --cpu-workers 2 --duration 5
         timeout-minutes: 10

.github/workflows/dependency-cooldown-audit.yml

@@ -0,0 +1,13 @@
+name: Dependency cooldown audit
+
+on:
+  schedule:
+    - cron: '0 6 * * 1'
+  workflow_dispatch:
+
+permissions:
+  contents: read
+
+jobs:
+  audit:
+    uses: Quantus-Network/shared-workflows/.github/workflows/dependency-cooldown-audit.yml@v1

.github/workflows/dependency-cooldown.yml

@@ -0,0 +1,14 @@
+name: Dependency cooldown
+
+on:
+  pull_request:
+    # labeled/unlabeled/edited are required so the gate re-runs when an
+    # emergency bypass label or reason is added.
+    types: [opened, synchronize, reopened, labeled, unlabeled, edited]
+
+permissions:
+  contents: read
+
+jobs:
+  dependency-cooldown:
+    uses: Quantus-Network/shared-workflows/.github/workflows/dependency-cooldown.yml@v1

Cargo.lock

@@ -577,7 +577,7 @@ dependencies = [
 
 [[package]]
 name = "engine-cpu"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "criterion",
  "hex",
@@ -588,7 +588,7 @@ dependencies = [
 
 [[package]]
 name = "engine-gpu"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "bytemuck",
  "criterion",
@@ -1476,7 +1476,7 @@ dependencies = [
 
 [[package]]
 name = "metrics"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "log",
@@ -1504,7 +1504,7 @@ dependencies = [
 
 [[package]]
 name = "miner-cli"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "clap",
  "engine-cpu",
@@ -1522,7 +1522,7 @@ dependencies = [
 
 [[package]]
 name = "miner-service"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "crossbeam-channel",
@@ -1543,7 +1543,7 @@ dependencies = [
 
 [[package]]
 name = "miner-telemetry"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "futures",
@@ -2003,7 +2003,7 @@ dependencies = [
 
 [[package]]
 name = "pool-service"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "clap",
@@ -2049,7 +2049,7 @@ dependencies = [
 
 [[package]]
 name = "pow-core"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "hex",
  "primitive-types 0.13.1",
@@ -2291,7 +2291,7 @@ dependencies = [
 
 [[package]]
 name = "quic-transport"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "quantus-miner-api",
@@ -2871,7 +2871,7 @@ dependencies = [
 
 [[package]]
 name = "solver-wasm"
-version = "4.0.0"
+version = "4.0.2"
 dependencies = [
  "primitive-types 0.13.1",
  "qpow-math",

Cargo.toml

@@ -18,7 +18,7 @@ resolver = "2"
 edition = "2021"
 authors = ["Quantus Network"]
 description = "Quantus External Miner Workspace"
-version = "4.0.0"
+version = "4.0.2"
 
 [workspace.dependencies]
 anyhow = "1"

crates/engine-gpu/examples/trusted_hashrate.rs

@@ -46,9 +46,10 @@ impl Runner {
             .await
             .unwrap();
 
+        let kernel = engine_gpu::Kernel::for_adapter(&adapter);
         let desc = wgpu::ShaderModuleDescriptor {
-            label: None,
-            source: wgpu::ShaderSource::Wgsl(include_str!("../src/mining_u64.wgsl").into()),
+            label: Some(kernel.label()),
+            source: wgpu::ShaderSource::Wgsl(kernel.source().into()),
         };
         let shader = if trusted {
             unsafe {

crates/engine-gpu/src/kernels/mining_u64.wgsl

@@ -1,7 +1,5 @@
-// Quantus Mining Shader - native-u64 variant (requires wgpu Features::SHADER_INT64)
-// Implements Poseidon2 hash over the Goldilocks field with plonky2-style lazy reduction:
-// values live in [0, 2^64) and are only canonicalized when squeezed out.
-// Must produce byte-identical results to mining.wgsl / qp-poseidon-core.
+// Native-u64 Poseidon2 mining kernel. Default path for non-Apple GPUs.
+// Requires wgpu Features::SHADER_INT64. Bit-exact with mining.wgsl / pow_core.
 
 @group(0) @binding(0) var<storage, read_write> results: array<atomic<u32>>;
 // Sponge state after absorbing header + high nonce half (12 felts as LE u32 pairs),

crates/engine-gpu/src/kernels/mining_u64_apple.wgsl

@@ -0,0 +1,650 @@
+// Native-u64 Poseidon2 mining kernel. Apple Metal path only.
+// Requires wgpu Features::SHADER_INT64. Bit-exact with mining.wgsl / pow_core.
+
+@group(0) @binding(0) var<storage, read_write> results: array<atomic<u32>>;
+// Sponge state after absorbing header + high nonce half (12 felts as LE u32 pairs),
+// precomputed on the host per batch. See pow_core::mining_midstate.
+@group(0) @binding(1) var<storage, read> midstate: array<u32, 24>;
+@group(0) @binding(2) var<storage, read> start_nonce: array<u32, 16>;
+@group(0) @binding(3) var<storage, read> difficulty_target: array<u32, 16>;
+@group(0) @binding(4) var<storage, read> dispatch_config: array<u32, 3>;
+
+const P64: u64 = 0xFFFFFFFF00000001lu;
+// EPS64 = 2^32 - 1 = 2^64 mod P
+const EPS64: u64 = 0xFFFFFFFFlu;
+
+// Padded with a trailing zero so the internal layer can add the next round's
+// constant unconditionally.
+const RC_INTERNAL: array<u64, 23> = array<u64, 23>(
+    0x97f7798a784ad863lu, 0xd1d2bf082f60d4f0lu, 0x69a377a79f9ad206lu, 0xa9d06906a3858e24lu, 0x295275001eede5b5lu, 0x5874e441117bd746lu, 0x8a084bbba8ed86cclu, 0x3defd7645cde6425lu, 0x3998cfe6871cc137lu, 0x3e52ef8bca48314alu, 0x964a209f85dc9ecclu, 0x3fcc9ee82cc4577elu, 0x8e79b4a5d0096d6dlu, 0x8492362ad2392556lu, 0xee72f470262574d6lu, 0x1e0e18496da2444alu, 0x0f3a74bf215eaac6lu, 0x1b061b76a1c0ded3lu, 0x192c42d86803d7a6lu, 0xf6d49ff997ae0260lu, 0x3ec372e7a0fa3786lu, 0x5538cdf4f23445d3lu, 0lu
+);
+const RC_INITIAL: array<array<u64, 12>, 4> = array<array<u64, 12>, 4>(
+    array<u64, 12>(0xc002e770975b1607lu, 0xbca51a8dfe14593alu, 0x72938dfbe774f7f9lu, 0xe4f2fe29e03234aclu, 0xd5e0ba2f541b6449lu, 0xec33b868f3cc46c1lu, 0x486dcb55419d475alu, 0x6c1cb2a358cc24f1lu, 0xe3f30d509a1436bblu, 0xd9a64f068dca7c29lu, 0xe59b3f57aabba1aelu, 0x2a3dd4505b478fdclu),
+    array<u64, 12>(0xada1f8dc7676ed25lu, 0x2711aa8b5509d516lu, 0x4ae6acd0c9c92897lu, 0x56eb3d6b5256d67alu, 0x1f7a9d55923bf51elu, 0x3600427d397a7f68lu, 0xe5076df75b72c3d0lu, 0xfcd59aa12c6090adlu, 0xcd895e8c68b57a9elu, 0x41df7ef9d730ae3elu, 0xee3e2b889abe977dlu, 0xd29bb7edbeb9c405lu),
+    array<u64, 12>(0x7d5c08eef608e382lu, 0x89ae889caaf0802clu, 0xb35a8e976d2af617lu, 0xdb14234eafaf5173lu, 0x78f04462d48b1c98lu, 0x265293b0e47ce88alu, 0x999a649b69b9d32flu, 0x64b0a186698e01d3lu, 0xee0b22d0dfae8bb8lu, 0x4fd53e50ca04a7eelu, 0x5762bfe181f25047lu, 0xf51593e2beb5e3bdlu),
+    array<u64, 12>(0x1e5e2b5760e32477lu, 0x622462a1f9aaaeedlu, 0xaa284b3ecdb222aelu, 0x63c8e72f542bf3fclu, 0x3ba588cacb43b5e0lu, 0x23eda6f3c99150ddlu, 0xaad3bea4baac9a5alu, 0xe9da8d699b94184alu, 0xcdb13f4cd93e024clu, 0x902cbd0956f655e3lu, 0x5b4e40ffc759532flu, 0xde795c20a2357af7lu)
+);
+const RC_TERMINAL: array<array<u64, 12>, 4> = array<array<u64, 12>, 4>(
+    array<u64, 12>(0x7b72c539e0ea4c6elu, 0x144573dae2ce9976lu, 0x802028b68f35fc88lu, 0x6d36c5022c4fe7c2lu, 0xa205d0ffa9b9def3lu, 0xf6e7e38b1ea6ba2flu, 0x34f7909ae5258d64lu, 0xb0464d9d77b97fcalu, 0x64ddb9d5de7e00a6lu, 0x0ed0d75c27975d97lu, 0x1cbb36f11127338blu, 0x6673e505cfd0b6balu),
+    array<u64, 12>(0x605f902830872e01lu, 0x3fd5eb927e95fe4flu, 0xe81025b5a24c69cdlu, 0xf7d0ce75de23f74elu, 0xf39942b6a8585089lu, 0x6d808a08f7b71df6lu, 0xf8806b6588f49a8blu, 0x57df2d8c2a32107alu, 0x16e7c2074d654a2dlu, 0x213de241fcf33835lu, 0xb0f2b8905a0976f6lu, 0xd8e3cf2bbd355417lu),
+    array<u64, 12>(0xe498691679d9330flu, 0x763b45d2a3821b28lu, 0x0908bf65eb0a1f0dlu, 0x7691eb2d194b24f4lu, 0x0e43551233ae13b2lu, 0x93c393dbfc2fe76flu, 0x98f607485d48cdealu, 0xe3d95f30309819c0lu, 0x1ef581a93eaf6acflu, 0x0b24c1b7a030fca4lu, 0x624370be5670b327lu, 0x5f1e28615a11e486lu),
+    array<u64, 12>(0xfe04051f909e042blu, 0x7257e5b147fd3803lu, 0xe6ae134bb82f2e78lu, 0x5711fd5cf4784511lu, 0xf83a42660c08c0bclu, 0x2cd8c96d9a3ce855lu, 0x7d2ffb1bb0e17271lu, 0x85ae1528caea3811lu, 0x52a345d5c7adb0b8lu, 0x504c4c51f3faee94lu, 0xbce34a649cfccaf9lu, 0xe0a3389266fb6dc9lu)
+);
+const MDS_DIAG: array<u64, 12> = array<u64, 12>(
+    0xc3b6c08e23ba9300lu, 0xd84b5de94a324fb6lu, 0x0d0c371c5b35b84flu, 0x7964f570e7188037lu, 0x5daf18bbd996604blu, 0x6743bc47b9595257lu, 0x5528b9362c59bb70lu, 0xac45e25b7127b68blu, 0xa2077d7dfbb606b5lu, 0xf3faac6faee378aelu, 0x0c6388b51545e883lu, 0xd27dbb6944917b60lu
+);
+
+// a + b mod P in lazy form. Wrapping carries fold back via 2^64 ≡ EPS64 (mod P).
+fn gf64_add(a: u64, b: u64) -> u64 {
+    let s0 = a + b;
+    let c1 = s0 < a;
+    let s1 = s0 + select(0lu, EPS64, c1);
+    let c2 = c1 && (s1 < s0);
+    return s1 + select(0lu, EPS64, c2);
+}
+
+// Sum accumulator: value = lo + carries * 2^64, folded once via 2^64 ≡ EPS64 (mod P).
+struct Acc {
+    lo: u64,
+    carries: u32,
+}
+
+fn acc_add(a: Acc, b: u64) -> Acc {
+    let s = a.lo + b;
+    return Acc(s, a.carries + select(0u, 1u, s < a.lo));
+}
+
+fn acc_add2(a: Acc, b: Acc) -> Acc {
+    let s = a.lo + b.lo;
+    return Acc(s, a.carries + b.carries + select(0u, 1u, s < a.lo));
+}
+
+fn acc_fold(a: Acc) -> u64 {
+    let c = u64(a.carries);
+    let t = a.lo + ((c << 32u) - c);
+    return t + select(0lu, EPS64, t < a.lo);
+}
+
+struct U128 {
+    lo: u64,
+    hi: u64,
+}
+
+// Reduce a 128-bit value (lo + hi*2^64) mod P using
+// 2^64 ≡ EPS64 and 2^96 ≡ -1 (mod P).
+fn gf64_reduce(v: U128) -> u64 {
+    let hi_hi = v.hi >> 32u;
+    let hi_lo = v.hi & EPS64;
+    var t0 = v.lo - hi_hi;
+    t0 = t0 - select(0lu, EPS64, v.lo < hi_hi);
+    let t1 = hi_lo * EPS64;
+    let t2 = t0 + t1;
+    return t2 + select(0lu, EPS64, t2 < t0);
+}
+
+fn mul_wide(a: u64, b: u64) -> U128 {
+    let a_lo = a & EPS64;
+    let a_hi = a >> 32u;
+    let b_lo = b & EPS64;
+    let b_hi = b >> 32u;
+    let ll = a_lo * b_lo;
+    let lh = a_lo * b_hi;
+    let hl = a_hi * b_lo;
+    let hh = a_hi * b_hi;
+    let mid = (ll >> 32u) + (lh & EPS64) + (hl & EPS64);
+    return U128((mid << 32u) | (ll & EPS64), hh + (lh >> 32u) + (hl >> 32u) + (mid >> 32u));
+}
+
+// (a*b + addend) mod P for b <= 2^64 - 2^32 (all MDS_DIAG entries): the addend's
+// value and carries are folded into the 128-bit product before reduction.
+fn gf64_mul_add(a: u64, b: u64, addend: Acc) -> u64 {
+    let v = mul_wide(a, b);
+    let lo = v.lo + addend.lo;
+    return gf64_reduce(U128(lo, v.hi + u64(addend.carries) + select(0lu, 1lu, lo < v.lo)));
+}
+
+fn gf64_mul(a: u64, b: u64) -> u64 {
+    return gf64_reduce(mul_wide(a, b));
+}
+
+fn gf64_sqr(a: u64) -> u64 {
+    let a_lo = a & EPS64;
+    let a_hi = a >> 32u;
+    let ll = a_lo * a_lo;
+    let lh = a_lo * a_hi;
+    let hh = a_hi * a_hi;
+    let mid = (ll >> 32u) + ((lh & EPS64) << 1u);
+    return gf64_reduce(U128((mid << 32u) | (ll & EPS64), hh + ((lh >> 32u) << 1u) + (mid >> 32u)));
+}
+
+fn gf64_sbox(x: u64) -> u64 {
+    let x2 = gf64_sqr(x);
+    let x4 = gf64_sqr(x2);
+    let x6 = gf64_mul(x4, x2);
+    return gf64_mul(x6, x);
+}
+
+fn gf64_canon(a: u64) -> u64 {
+    return a - select(0lu, P64, a >= P64);
+}
+
+// 4x4 MDS circ(2, 3, 1, 1) on one chunk, results left unreduced.
+fn mds4(x0: u64, x1: u64, x2: u64, x3: u64) -> array<Acc, 4> {
+    let t01 = acc_add(Acc(x0, 0u), x1);
+    let t23 = acc_add(Acc(x2, 0u), x3);
+    let t0123 = acc_add2(t01, t23);
+    let t01123 = acc_add(t0123, x1);
+    let t01233 = acc_add(t0123, x3);
+    return array<Acc, 4>(
+        acc_add2(t01123, t01),
+        acc_add2(t01123, acc_add(Acc(x2, 0u), x2)),
+        acc_add2(t01233, t23),
+        acc_add2(t01233, acc_add(Acc(x0, 0u), x0))
+    );
+}
+
+// External linear layer: 4x4 MDS on each chunk, then circulant sums, plus the
+// next round's constants. Additions are accumulated unreduced (at most 27
+// carries) and folded once per output.
+fn ext_layer64(state: ptr<function, array<u64, 12>>, rc: array<u64, 12>) {
+    var y: array<Acc, 12>;
+    for (var chunk = 0u; chunk < 3u; chunk++) {
+        let o = chunk * 4u;
+        let m = mds4((*state)[o], (*state)[o + 1u], (*state)[o + 2u], (*state)[o + 3u]);
+        y[o] = m[0];
+        y[o + 1u] = m[1];
+        y[o + 2u] = m[2];
+        y[o + 3u] = m[3];
+    }
+    for (var k = 0u; k < 4u; k++) {
+        let s = acc_add2(acc_add2(y[k], y[k + 4u]), y[k + 8u]);
+        (*state)[k] = acc_fold(acc_add(acc_add2(y[k], s), rc[k]));
+        (*state)[k + 4u] = acc_fold(acc_add(acc_add2(y[k + 4u], s), rc[k + 4u]));
+        (*state)[k + 8u] = acc_fold(acc_add(acc_add2(y[k + 8u], s), rc[k + 8u]));
+    }
+}
+
+// Internal linear layer: diagonal matrix plus full sum, plus the next round's
+// constant on element 0 (the S-box output, which is summed last).
+fn int_layer64(state: ptr<function, array<u64, 12>>, rc0: u64) {
+    let s12 = acc_add(Acc((*state)[1], 0u), (*state)[2]);
+    let s34 = acc_add(Acc((*state)[3], 0u), (*state)[4]);
+    let s56 = acc_add(Acc((*state)[5], 0u), (*state)[6]);
+    let s78 = acc_add(Acc((*state)[7], 0u), (*state)[8]);
+    let s910 = acc_add(Acc((*state)[9], 0u), (*state)[10]);
+    let s1234 = acc_add2(s12, s34);
+    let s5678 = acc_add2(s56, s78);
+    let s91011 = acc_add(s910, (*state)[11]);
+    let sum = acc_add(acc_add2(acc_add2(s1234, s5678), s91011), (*state)[0]);
+    (*state)[0] = gf64_mul_add((*state)[0], MDS_DIAG[0], acc_add(sum, rc0));
+    (*state)[1] = gf64_mul_add((*state)[1], MDS_DIAG[1], sum);
+    (*state)[2] = gf64_mul_add((*state)[2], MDS_DIAG[2], sum);
+    (*state)[3] = gf64_mul_add((*state)[3], MDS_DIAG[3], sum);
+    (*state)[4] = gf64_mul_add((*state)[4], MDS_DIAG[4], sum);
+    (*state)[5] = gf64_mul_add((*state)[5], MDS_DIAG[5], sum);
+    (*state)[6] = gf64_mul_add((*state)[6], MDS_DIAG[6], sum);
+    (*state)[7] = gf64_mul_add((*state)[7], MDS_DIAG[7], sum);
+    (*state)[8] = gf64_mul_add((*state)[8], MDS_DIAG[8], sum);
+    (*state)[9] = gf64_mul_add((*state)[9], MDS_DIAG[9], sum);
+    (*state)[10] = gf64_mul_add((*state)[10], MDS_DIAG[10], sum);
+    (*state)[11] = gf64_mul_add((*state)[11], MDS_DIAG[11], sum);
+}
+
+fn sbox_lanes(state: ptr<function, array<u64, 12>>, lanes: u32) {
+    for (var i = 0u; i < lanes; i++) {
+        (*state)[i] = gf64_sbox((*state)[i]);
+    }
+}
+
+fn add_rc(state: ptr<function, array<u64, 12>>, rc: array<u64, 12>) {
+    for (var i = 0u; i < 12u; i++) {
+        (*state)[i] = gf64_add((*state)[i], rc[i]);
+    }
+}
+
+const RC_ZERO: array<u64, 12> = array<u64, 12>(0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu);
+// Constants added by each external layer, which feeds the S-box that follows it:
+// entries 0..3 precede the initial rounds, entry 4 carries the first internal
+// round's constant on element 0, entries 5..7 precede terminal rounds 1..3.
+const RC_EXT: array<array<u64, 12>, 9> = array<array<u64, 12>, 9>(
+    RC_INITIAL[0], RC_INITIAL[1], RC_INITIAL[2], RC_INITIAL[3],
+    array<u64, 12>(RC_INTERNAL[0], 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu),
+    RC_TERMINAL[1], RC_TERMINAL[2], RC_TERMINAL[3], RC_ZERO
+);
+
+// One loop drives all 30 rounds so each layer is emitted once: step 0 is the
+// initial external layer, steps 1..4 and 27..30 are external rounds, steps
+// 5..26 are internal rounds. Round constants are added by the linear layer
+// preceding each S-box (the last internal layer is followed by the first
+// terminal round's constants).
+fn permute64(state: ptr<function, array<u64, 12>>) {
+    for (var k = 0u; k < 31u; k++) {
+        let is_ext = k < 5u || k > 26u;
+        if (k > 0u) {
+            sbox_lanes(state, select(1u, 12u, is_ext));
+        }
+        if (is_ext) {
+            ext_layer64(state, RC_EXT[select(k, k - 22u, k > 26u)]);
+        } else {
+            int_layer64(state, RC_INTERNAL[k - 4u]);
+            if (k == 26u) {
+                add_rc(state, RC_TERMINAL[0]);
+            }
+        }
+    }
+}
+
+fn bswap32(v: u32) -> u32 {
+    return ((v & 0xFFu) << 24u) | ((v & 0xFF00u) << 8u) | ((v >> 8u) & 0xFF00u) | (v >> 24u);
+}
+
+@compute @workgroup_size(256)
+fn mining_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    if (atomicLoad(&results[0]) != 0u) {
+        return;
+    }
+    let thread_id = global_id.x;
+    let total_threads = dispatch_config[0];
+    let nonces_per_thread = dispatch_config[1];
+    let total_nonces = dispatch_config[2];
+    if (thread_id >= total_threads) {
+        return;
+    }
+    let base_index = thread_id * nonces_per_thread;
+
+    // Hoist uniform storage reads out of the nonce loop
+    var mid: array<u64, 12>;
+    for (var i = 0u; i < 12u; i++) {
+        mid[i] = (u64(midstate[2u * i + 1u]) << 32u) | u64(midstate[2u * i]);
+    }
+    var tgt: array<u32, 16>;
+    for (var i = 0u; i < 16u; i++) {
+        tgt[i] = difficulty_target[i];
+    }
+    var nonce_base: array<u32, 16>;
+    for (var i = 0u; i < 16u; i++) {
+        nonce_base[i] = start_nonce[i];
+    }
+
+    for (var j = 0u; j < nonces_per_thread; j = j + 1u) {
+        let logical_index = base_index + j;
+        if (logical_index >= total_nonces) {
+            break;
+        }
+        if (j > 0u && atomicLoad(&results[0]) != 0u) {
+            return;
+        }
+
+        // The host guarantees a batch never carries into the high nonce half
+        // (limbs 8..15), so only the low 256 bits are incremented here.
+        var current_nonce: array<u32, 16>;
+        let val0 = nonce_base[0];
+        let sum0 = val0 + logical_index;
+        current_nonce[0] = sum0;
+        var carry = select(0u, 1u, sum0 < val0);
+        for (var i = 1u; i < 8u; i++) {
+            let val = nonce_base[i];
+            let sum = val + carry;
+            current_nonce[i] = sum;
+            carry = select(0u, 1u, sum < val);
+        }
+        for (var i = 8u; i < 16u; i++) {
+            current_nonce[i] = nonce_base[i];
+        }
+
+        // Resume the sponge from the precomputed midstate: absorb the low
+        // nonce half, pad, squeeze twice (3 permutations instead of 5).
+        var st = mid;
+        for (var i = 0u; i < 8u; i++) {
+            st[i] = gf64_add(st[i], u64(bswap32(current_nonce[7u - i])));
+        }
+        // Squeeze-and-compare phases share one inlined permutation: phase 0
+        // pads after absorbing, phase 1 yields the most significant 256 bits of
+        // the hash, which decide hash-vs-target on their own unless they exactly
+        // equal the target's high half, and only candidates run phase 2 for the
+        // low half. Byte-swapped hash words are produced on demand.
+        var hash_le: array<u32, 16>;
+        var cmp = 0u;
+        var below = false;
+        for (var phase = 0u; phase < 3u; phase++) {
+            permute64(&st);
+            if (phase == 0u) {
+                st[0] = gf64_add(st[0], 1lu);
+                st[1] = gf64_add(st[1], 1lu);
+                continue;
+            }
+            var words: array<u32, 8>;
+            for (var i = 0u; i < 4u; i++) {
+                let c = gf64_canon(st[i]);
+                words[2u * i] = bswap32(u32(c & EPS64));
+                words[2u * i + 1u] = bswap32(u32(c >> 32u));
+            }
+            let base = select(15u, 7u, phase == 2u);
+            for (var i = 0u; i < 8u; i++) {
+                hash_le[base - i] = words[i];
+            }
+            if (phase == 1u) {
+                for (var i = 0u; i < 8u; i++) {
+                    let h = words[i];
+                    let t = tgt[15u - i];
+                    if (h != t) {
+                        cmp = select(2u, 1u, h > t);
+                        break;
+                    }
+                }
+                if (cmp == 1u) {
+                    break;
+                }
+            } else {
+                below = cmp == 2u;
+                if (!below) {
+                    for (var i = 0u; i < 8u; i++) {
+                        let h = words[i];
+                        let t = tgt[7u - i];
+                        if (h != t) {
+                            below = h < t;
+                            break;
+                        }
+                    }
+                }
+            }
+        }
+        if (cmp == 1u) {
+            continue;
+        }
+
+        if (below) {
+            if (atomicExchange(&results[0], 1u) == 0u) {
+                for (var i = 0u; i < 16u; i++) {
+                    atomicStore(&results[1u + i], current_nonce[i]);
+                    atomicStore(&results[17u + i], hash_le[i]);
+                }
+            }
+            return;
+        }
+    }
+}
+
+// ---------------------------------------------------------------------------
+// Compatibility layer: same API as mining.wgsl, backed by the u64 core above.
+// Only used by the component test harness; the mining kernel never calls it.
+// All outputs are canonical, matching the reference implementation.
+// ---------------------------------------------------------------------------
+
+struct GoldilocksField {
+    limb0: u32,
+    limb1: u32,
+}
+
+const INTERNAL_CONSTANTS: array<array<u32, 2>, 22> = array<array<u32, 2>, 22>(
+    array<u32, 2>(2018170979u, 2549578122u),
+    array<u32, 2>(794875120u, 3520249608u),
+    array<u32, 2>(2677723654u, 1772320679u),
+    array<u32, 2>(2743438884u, 2849007878u),
+    array<u32, 2>(518907317u, 693269760u),
+    array<u32, 2>(293328710u, 1484055617u),
+    array<u32, 2>(2834138828u, 2315799483u),
+    array<u32, 2>(1558078501u, 1039128420u),
+    array<u32, 2>(2266808631u, 966316006u),
+    array<u32, 2>(3393728842u, 1045622667u),
+    array<u32, 2>(2245828300u, 2521440415u),
+    array<u32, 2>(751064958u, 1070374632u),
+    array<u32, 2>(3490278765u, 2390340773u),
+    array<u32, 2>(3526960470u, 2224174634u),
+    array<u32, 2>(639988950u, 4000511088u),
+    array<u32, 2>(1839350858u, 504240201u),
+    array<u32, 2>(559852230u, 255489215u),
+    array<u32, 2>(2713771731u, 453385078u),
+    array<u32, 2>(1745082278u, 422331096u),
+    array<u32, 2>(2544763488u, 4141129721u),
+    array<u32, 2>(2700752774u, 1052996327u),
+    array<u32, 2>(4063512019u, 1429786100u)
+);
+
+const INITIAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(2539329031u, 3221415792u),
+        array<u32, 2>(4262746426u, 3164936845u),
+        array<u32, 2>(3883202553u, 1922272763u),
+        array<u32, 2>(3761386668u, 3841130025u),
+        array<u32, 2>(1411081289u, 3588274735u),
+        array<u32, 2>(4090250945u, 3962812520u),
+        array<u32, 2>(1100826458u, 1215155029u),
+        array<u32, 2>(1489773809u, 1813820067u),
+        array<u32, 2>(2585015995u, 3824356688u),
+        array<u32, 2>(2378857513u, 3651555078u),
+        array<u32, 2>(2864423342u, 3852156759u),
+        array<u32, 2>(1531416540u, 708695120u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(1987505445u, 2913073372u),
+        array<u32, 2>(1426707734u, 655469195u),
+        array<u32, 2>(3385403543u, 1256631504u),
+        array<u32, 2>(1381422714u, 1458257259u),
+        array<u32, 2>(2453402910u, 528129365u),
+        array<u32, 2>(964329320u, 905986685u),
+        array<u32, 2>(1534247888u, 3842469367u),
+        array<u32, 2>(744525997u, 4241857185u),
+        array<u32, 2>(1756723870u, 3448331916u),
+        array<u32, 2>(3610291774u, 1105166073u),
+        array<u32, 2>(2596181885u, 3997051784u),
+        array<u32, 2>(3199845381u, 3533420525u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(4127777666u, 2103183598u),
+        array<u32, 2>(2867888172u, 2309916828u),
+        array<u32, 2>(1831532055u, 3009056407u),
+        array<u32, 2>(2947502451u, 3675530062u),
+        array<u32, 2>(3565886616u, 2029012066u),
+        array<u32, 2>(3833391242u, 642945968u),
+        array<u32, 2>(1773785903u, 2577032347u),
+        array<u32, 2>(1770914259u, 1689297286u),
+        array<u32, 2>(3752758200u, 3993707216u),
+        array<u32, 2>(3389302766u, 1339375184u),
+        array<u32, 2>(2180141127u, 1466089441u),
+        array<u32, 2>(3199591357u, 4111832034u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(1625498743u, 509487959u),
+        array<u32, 2>(4188712685u, 1646551713u),
+        array<u32, 2>(3451003566u, 2854767422u),
+        array<u32, 2>(1412166652u, 1674110767u),
+        array<u32, 2>(3410212320u, 1000704202u),
+        array<u32, 2>(3381743837u, 602777331u),
+        array<u32, 2>(3131873882u, 2866003620u),
+        array<u32, 2>(2610174026u, 3923414377u),
+        array<u32, 2>(3644719692u, 3450945356u),
+        array<u32, 2>(1458984419u, 2418851081u),
+        array<u32, 2>(3344519983u, 1531855103u),
+        array<u32, 2>(2721413879u, 3732495392u)
+    )
+);
+
+const TERMINAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(3773451374u, 2071119161u),
+        array<u32, 2>(3805190518u, 340095962u),
+        array<u32, 2>(2402679944u, 2149591222u),
+        array<u32, 2>(743434178u, 1832305922u),
+        array<u32, 2>(2847530739u, 2718290175u),
+        array<u32, 2>(514243119u, 4142392203u),
+        array<u32, 2>(3844443492u, 888639642u),
+        array<u32, 2>(2008645578u, 2957397405u),
+        array<u32, 2>(3732799654u, 1692252629u),
+        array<u32, 2>(664231319u, 248567644u),
+        array<u32, 2>(287781771u, 482031345u),
+        array<u32, 2>(3486561978u, 1718871301u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(814165505u, 1616875560u),
+        array<u32, 2>(2123759183u, 1070984082u),
+        array<u32, 2>(2722916813u, 3893372341u),
+        array<u32, 2>(3726899022u, 4157656693u),
+        array<u32, 2>(2824360073u, 4086907574u),
+        array<u32, 2>(4155973110u, 1837140488u),
+        array<u32, 2>(2297731723u, 4169165669u),
+        array<u32, 2>(707924090u, 1474243980u),
+        array<u32, 2>(1298483757u, 384287239u),
+        array<u32, 2>(4243798069u, 557703745u),
+        array<u32, 2>(1510569718u, 2968696976u),
+        array<u32, 2>(3174388759u, 3638808363u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(2044277519u, 3835193622u),
+        array<u32, 2>(2743212840u, 1983595986u),
+        array<u32, 2>(3943309069u, 151568229u),
+        array<u32, 2>(424355060u, 1989274413u),
+        array<u32, 2>(867046322u, 239293714u),
+        array<u32, 2>(4230997871u, 2479068123u),
+        array<u32, 2>(1565052394u, 2566260552u),
+        array<u32, 2>(815274432u, 3822673712u),
+        array<u32, 2>(1051683535u, 519405993u),
+        array<u32, 2>(2687564964u, 186958263u),
+        array<u32, 2>(1450226471u, 1648586942u),
+        array<u32, 2>(1511122054u, 1595811937u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(2426274859u, 4261676319u),
+        array<u32, 2>(1207777283u, 1918363057u),
+        array<u32, 2>(3090099832u, 3870167883u),
+        array<u32, 2>(4101522705u, 1460796764u),
+        array<u32, 2>(201900220u, 4164567654u),
+        array<u32, 2>(2587682901u, 752404845u),
+        array<u32, 2>(2967564913u, 2100296475u),
+        array<u32, 2>(3404347409u, 2242778408u),
+        array<u32, 2>(3350048952u, 1386431957u),
+        array<u32, 2>(4093308564u, 1347177553u),
+        array<u32, 2>(2633812729u, 3169012324u),
+        array<u32, 2>(1727753673u, 3768793234u)
+    )
+);
+
+fn gf_pack(g: GoldilocksField) -> u64 {
+    return (u64(g.limb1) << 32u) | u64(g.limb0);
+}
+
+fn gf_unpack(v: u64) -> GoldilocksField {
+    let c = gf64_canon(v);
+    return GoldilocksField(u32(c & EPS64), u32(c >> 32u));
+}
+
+fn gf_from_limbs(l0: u32, l1: u32) -> GoldilocksField {
+    return GoldilocksField(l0, l1);
+}
+
+fn gf_zero() -> GoldilocksField {
+    return GoldilocksField(0u, 0u);
+}
+
+fn gf_one() -> GoldilocksField {
+    return GoldilocksField(1u, 0u);
+}
+
+fn gf_from_u32(val: u32) -> GoldilocksField {
+    return GoldilocksField(val, 0u);
+}
+
+fn gf_from_u64_parts(low: u32, high: u32) -> GoldilocksField {
+    return gf_unpack((u64(high) << 32u) | u64(low));
+}
+
+fn gf_from_const(val: array<u32, 2>) -> GoldilocksField {
+    return gf_from_u64_parts(val[0], val[1]);
+}
+
+fn gf_add(a: GoldilocksField, b: GoldilocksField) -> GoldilocksField {
+    return gf_unpack(gf64_add(gf_pack(a), gf_pack(b)));
+}
+
+fn gf_mul(a: GoldilocksField, b: GoldilocksField) -> GoldilocksField {
+    return gf_unpack(gf64_mul(gf_pack(a), gf_pack(b)));
+}
+
+fn sbox(x: GoldilocksField) -> GoldilocksField {
+    return gf_unpack(gf64_sbox(gf_pack(x)));
+}
+
+fn state_pack(state: ptr<function, array<GoldilocksField, 12>>, out: ptr<function, array<u64, 12>>) {
+    for (var i = 0u; i < 12u; i++) {
+        (*out)[i] = gf_pack((*state)[i]);
+    }
+}
+
+fn state_unpack(v: ptr<function, array<u64, 12>>, state: ptr<function, array<GoldilocksField, 12>>) {
+    for (var i = 0u; i < 12u; i++) {
+        (*state)[i] = gf_unpack((*v)[i]);
+    }
+}
+
+fn external_linear_layer(state: ptr<function, array<GoldilocksField, 12>>) {
+    var st: array<u64, 12>;
+    state_pack(state, &st);
+    ext_layer64(&st, RC_ZERO);
+    state_unpack(&st, state);
+}
+
+fn internal_linear_layer(state: ptr<function, array<GoldilocksField, 12>>) {
+    var st: array<u64, 12>;
+    state_pack(state, &st);
+    int_layer64(&st, 0lu);
+    state_unpack(&st, state);
+}
+
+fn poseidon2_permute(state: ptr<function, array<GoldilocksField, 12>>) {
+    var st: array<u64, 12>;
+    state_pack(state, &st);
+    permute64(&st);
+    state_unpack(&st, state);
+}
+
+fn bytes_to_field_elements(input: array<u32, 24>) -> array<GoldilocksField, 25> {
+    var felts: array<GoldilocksField, 25>;
+    for (var i = 0u; i < 24u; i++) {
+        felts[i] = gf_from_u32(input[i]);
+    }
+    felts[24] = gf_one();
+    return felts;
+}
+
+fn field_elements_to_bytes(felts: array<GoldilocksField, 4>) -> array<u32, 8> {
+    var result: array<u32, 8>;
+    for (var i = 0u; i < 4u; i++) {
+        result[i * 2u] = felts[i].limb0;
+        result[i * 2u + 1u] = felts[i].limb1;
+    }
+    return result;
+}
+
+fn poseidon2_hash_squeeze_twice(input: array<u32, 24>) -> array<u32, 16> {
+    var st: array<u64, 12>;
+    for (var i = 0u; i < 12u; i++) {
+        st[i] = 0lu;
+    }
+    for (var chunk = 0u; chunk < 3u; chunk++) {
+        for (var i = 0u; i < 8u; i++) {
+            st[i] = gf64_add(st[i], u64(input[chunk * 8u + i]));
+        }
+        permute64(&st);
+    }
+    st[0] = gf64_add(st[0], 1lu);
+    st[1] = gf64_add(st[1], 1lu);
+    permute64(&st);
+
+    var result: array<u32, 16>;
+    for (var i = 0u; i < 4u; i++) {
+        let c = gf64_canon(st[i]);
+        result[2u * i] = u32(c & EPS64);
+        result[2u * i + 1u] = u32(c >> 32u);
+    }
+    permute64(&st);
+    for (var i = 0u; i < 4u; i++) {
+        let c = gf64_canon(st[i]);
+        result[8u + 2u * i] = u32(c & EPS64);
+        result[8u + 2u * i + 1u] = u32(c >> 32u);
+    }
+    return result;
+}
+
+fn hash_squeeze_twice(input: array<u32, 24>) -> array<u32, 16> {
+    return poseidon2_hash_squeeze_twice(input);
+}

crates/engine-gpu/src/kernels/mod.rs

@@ -0,0 +1,116 @@
+//! Poseidon2 mining kernels.
+//!
+//! Same `mining_main` bindings; must stay bit-exact with `pow_core`.
+//!
+//! - Apple Metal + `SHADER_INT64` → Apple Metal u64 (`mining_u64_apple.wgsl`)
+//! - other GPUs + `SHADER_INT64` → native u64 (`mining_u64.wgsl`)
+//! - no `SHADER_INT64` → 32-bit fallback (`mining.wgsl`)
+
+#[derive(Clone, Copy, Debug, Eq, PartialEq)]
+pub enum Kernel {
+    U32,
+    Default,
+    Apple,
+}
+
+impl Kernel {
+    pub const fn needs_int64(self) -> bool {
+        !matches!(self, Self::U32)
+    }
+
+    pub const fn id(self) -> &'static str {
+        match self {
+            Self::U32 => "u32",
+            Self::Default => "u64",
+            Self::Apple => "u64-apple",
+        }
+    }
+
+    pub const fn label(self) -> &'static str {
+        match self {
+            Self::U32 => "32-bit",
+            Self::Default => "native-u64",
+            Self::Apple => "native-u64 Apple Metal",
+        }
+    }
+
+    pub const fn source(self) -> &'static str {
+        match self {
+            Self::U32 => include_str!("mining.wgsl"),
+            Self::Default => include_str!("mining_u64.wgsl"),
+            Self::Apple => include_str!("mining_u64_apple.wgsl"),
+        }
+    }
+
+    pub fn for_adapter(adapter: &wgpu::Adapter) -> Self {
+        Self::for_adapter_info(&adapter.get_info(), adapter.features())
+    }
+
+    pub fn for_adapter_info(info: &wgpu::AdapterInfo, features: wgpu::Features) -> Self {
+        if !features.contains(wgpu::Features::SHADER_INT64) {
+            return Self::U32;
+        }
+        if info.backend == wgpu::Backend::Metal {
+            Self::Apple
+        } else {
+            Self::Default
+        }
+    }
+
+    pub const fn all() -> &'static [Self] {
+        &[Self::U32, Self::Default, Self::Apple]
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn info(backend: wgpu::Backend) -> wgpu::AdapterInfo {
+        wgpu::AdapterInfo {
+            name: "test".into(),
+            vendor: 0,
+            device: 0,
+            device_type: wgpu::DeviceType::DiscreteGpu,
+            driver: String::new(),
+            driver_info: String::new(),
+            backend,
+        }
+    }
+
+    #[test]
+    fn metal_with_int64_selects_apple() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Metal), wgpu::Features::SHADER_INT64),
+            Kernel::Apple
+        );
+    }
+
+    #[test]
+    fn vulkan_with_int64_selects_default() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Vulkan), wgpu::Features::SHADER_INT64),
+            Kernel::Default
+        );
+    }
+
+    #[test]
+    fn dx12_with_int64_selects_default() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Dx12), wgpu::Features::SHADER_INT64),
+            Kernel::Default
+        );
+    }
+
+    #[test]
+    fn no_int64_falls_back_to_u32() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Metal), wgpu::Features::empty()),
+            Kernel::U32
+        );
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Vulkan), wgpu::Features::empty()),
+            Kernel::U32
+        );
+    }
+}

crates/engine-gpu/src/lib.rs

@@ -2,10 +2,13 @@
 #![deny(unsafe_code)]
 
 mod gpu_tiers;
+mod kernels;
 
 pub mod end_to_end_tests;
 pub mod tests;
 
+pub use kernels::Kernel;
+
 use engine_cpu::{CancelCheck, Candidate, EngineStatus, FoundOrigin, MinerEngine, Range};
 use pow_core::{format_hashrate, format_u512, JobContext};
 use primitive_types::U512;
@@ -153,10 +156,10 @@ impl GpuContext {
 /// Create the mining shader module without naga's runtime bounds checks and
 /// loop bounding (~9% faster kernels).
 ///
-/// SAFETY: the sources are the static mining shaders compiled into this binary;
+/// SAFETY: the sources are the static mining kernels compiled into this binary;
 /// every buffer access is a constant-bounded loop index into fixed-size
 /// bindings the engine itself allocates, and all loops have static bounds
-/// (verified by the component test suites against both shader variants).
+/// (verified by the component test suites against every kernel path).
 #[allow(unsafe_code)]
 fn create_trusted_shader(device: &wgpu::Device, shader_source: &str) -> wgpu::ShaderModule {
     unsafe {
@@ -357,15 +360,13 @@ impl GpuEngine {
             );
             log::debug!(target: "gpu_engine", "Adapter {i} raw info: {info:?}");
 
-            // Prefer the native-u64 shader where supported (Apple/NVIDIA/modern AMD):
-            // Goldilocks arithmetic on u64 is far cheaper than 32-bit limb emulation.
-            let use_u64 = adapter.features().contains(wgpu::Features::SHADER_INT64);
+            let kernel = Kernel::for_adapter(&adapter);
 
             // Try to initialize this adapter with a proper timeout.
             // If the driver hangs, we'll skip this adapter after the timeout.
             let device_future = adapter.request_device(&wgpu::DeviceDescriptor {
                 label: Some("Mining Device"),
-                required_features: if use_u64 {
+                required_features: if kernel.needs_int64() {
                     wgpu::Features::SHADER_INT64
                 } else {
                     wgpu::Features::empty()
@@ -412,18 +413,14 @@ impl GpuEngine {
             // Shader and pipeline creation are synchronous - can't timeout, but usually fast
             let pipeline_start = std::time::Instant::now();
 
-            let shader_source = if use_u64 {
-                include_str!("mining_u64.wgsl")
-            } else {
-                include_str!("mining.wgsl")
-            };
             log::info!(
                 target: "gpu_engine",
-                "GPU device {i} ({}) using {} shader",
+                "GPU device {i} ({}) using {} [{}]",
                 info.name,
-                if use_u64 { "native-u64" } else { "32-bit" }
+                kernel.label(),
+                kernel.id()
             );
-            let shader = create_trusted_shader(&device, shader_source);
+            let shader = create_trusted_shader(&device, kernel.source());
 
             let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                 label: Some("Mining Pipeline"),

crates/engine-gpu/src/main.rs

@@ -13,37 +13,28 @@ async fn main() -> Result<(), Box<dyn std::error::Error>> {
         .expect("no GPU adapter");
 
     let mut failures = 0usize;
+    let has_int64 = adapter.features().contains(wgpu::Features::SHADER_INT64);
 
-    {
-        let (device, queue) = adapter
-            .request_device(&wgpu::DeviceDescriptor::default())
-            .await?;
-        failures += run_suite(
-            &device,
-            &queue,
-            include_str!("mining.wgsl"),
-            "32-bit (mining.wgsl)",
-        )
-        .await;
-    }
-
-    if adapter.features().contains(wgpu::Features::SHADER_INT64) {
+    for kernel in engine_gpu::Kernel::all() {
+        if kernel.needs_int64() && !has_int64 {
+            println!(
+                "\nSHADER_INT64 not supported on this adapter; skipping {}",
+                kernel.label()
+            );
+            continue;
+        }
         let (device, queue) = adapter
             .request_device(&wgpu::DeviceDescriptor {
-                label: Some("u64 Test Device"),
-                required_features: wgpu::Features::SHADER_INT64,
+                label: Some(kernel.label()),
+                required_features: if kernel.needs_int64() {
+                    wgpu::Features::SHADER_INT64
+                } else {
+                    wgpu::Features::empty()
+                },
                 ..Default::default()
             })
             .await?;
-        failures += run_suite(
-            &device,
-            &queue,
-            include_str!("mining_u64.wgsl"),
-            "native-u64 (mining_u64.wgsl)",
-        )
-        .await;
-    } else {
-        println!("\nSHADER_INT64 not supported on this adapter; skipping mining_u64.wgsl suite");
+        failures += run_suite(&device, &queue, kernel.source(), kernel.label()).await;
     }
 
     if failures > 0 {

examples/systemd/overrides/10-shared-hardware.conf

@@ -14,8 +14,8 @@
 #
 # Notes:
 # - CPUAffinity must be a subset of the service’s allowed cpuset (e.g., cgroups).
-# - The miner logs the detected cpuset mask at startup and, when metrics are enabled,
-#   exports a miner_effective_cpus gauge for dashboards.
+# - The miner counts CPUs via the process affinity mask and exports a
+#   miner_effective_cpus gauge on the always-on metrics endpoint.
 # - Adjust the CPU list and worker count to match your hardware and cpuset.
 #
 # Example below assumes an 8‑way logical CPU system (IDs 0‑7).
@@ -40,9 +40,8 @@ IOSchedulingPriority=4
 # Default time-sharing policy is fine for CPU-bound workloads when pinned.
 CPUSchedulingPolicy=other
 
-# Optional: pin worker count and enable metrics (uncomment and adjust as needed).
-# Environment=MINER_ENGINE=cpu-fast
-# Environment=MINER_WORKERS=4
+# Optional: pin worker count and change the metrics port (uncomment and adjust as needed).
+# Environment=MINER_CPU_WORKERS=4
 # Environment=MINER_METRICS_PORT=9900
 # Environment=RUST_LOG=info,miner=debug
 

examples/systemd/overrides/20-dedicated-hardware.conf

@@ -15,11 +15,11 @@
 # - CPUAffinity MUST be a subset of the service’s allowed cpuset (cgroup v2:
 #     /sys/fs/cgroup/cpuset.cpus.effective). If your deployment already constrains
 #   the miner with a cpuset, you can omit CPUAffinity here to inherit that mask.
-# - The miner logs the detected cpuset mask at startup (debug) and, when metrics
-#   are enabled, exports miner_effective_cpus for dashboards.
-# - Set MINER_WORKERS to match the number of CPUs in your affinity mask if you want
-#   full utilization; otherwise the miner default (50% of effective CPUs) is used
-#   when --workers is omitted.
+# - The miner counts CPUs via the process affinity mask and exports
+#   miner_effective_cpus on the always-on metrics endpoint.
+# - Set MINER_CPU_WORKERS to match the number of CPUs in your affinity mask if you
+#   want full utilization; otherwise the miner auto-detects when --cpu-workers is
+#   omitted.
 
 [Service]
 # Bind miner to ALL logical CPUs intended for mining. EDIT this to match your host.
@@ -46,10 +46,9 @@ IOSchedulingPriority=0
 TasksMax=infinity
 MemoryMax=infinity
 
-# Optional: pin worker count and enable metrics (uncomment and adjust).
-# Ensure MINER_WORKERS matches the size of CPUAffinity (or the cpuset mask) to use all CPUs.
-# Environment=MINER_ENGINE=cpu-fast
-# Environment=MINER_WORKERS=16
+# Optional: pin worker count and change the metrics port (uncomment and adjust).
+# Ensure MINER_CPU_WORKERS matches the size of CPUAffinity (or the cpuset mask) to use all CPUs.
+# Environment=MINER_CPU_WORKERS=16
 # Environment=MINER_METRICS_PORT=9900
 # Environment=RUST_LOG=info,miner=debug
 
@@ -64,5 +63,5 @@ MemoryMax=infinity
 #   and compare with:
 #     cat /sys/fs/cgroup/cpuset.cpus.effective
 #
-# - The service logs the cpuset mask at startup and the metrics exporter (if enabled)
-#   exposes miner_effective_cpus for observability.
+# - The always-on metrics exporter exposes miner_effective_cpus for
+#   observability of the detected CPU count (MINER_METRICS_PORT changes the port).

examples/systemd/quantus-miner.service

@@ -15,14 +15,24 @@
 #
 # Notes:
 #   - The miner CLI supports environment variables for all flags (see README):
-#       MINER_PORT, MINER_METRICS_PORT, MINER_WORKERS, MINER_ENGINE,
-#       MINER_PROGRESS_CHUNK_MS,
-#       MINER_MANIP_SOLVED_BLOCKS, MINER_MANIP_BASE_DELAY_NS,
-#       MINER_MANIP_STEP_BATCH, MINER_MANIP_THROTTLE_CAP.
+#       MINER_NODE_ADDR, MINER_AUTH_TOKEN_FILE (or MINER_AUTH_TOKEN),
+#       MINER_TLS_CERT_SHA256_FILE (or MINER_TLS_CERT_SHA256),
+#       MINER_CPU_WORKERS, MINER_GPU_DEVICES, MINER_GPU_BATCH_SIZE,
+#       MINER_CPU_BATCH_SIZE, MINER_GPU_THROTTLE_MS, MINER_METRICS_PORT,
+#       MINER_ALLOW_INTEGRATED, MINER_VERBOSE.
 #     This unit relies on those env vars (set in /etc/default|/etc/sysconfig) so
 #     ExecStart can remain stable and simple.
-#   - The service logs the detected cpuset mask (if any) and emits a
-#     miner_effective_cpus metric when metrics are enabled.
+#   - The node requires miner auth: MINER_AUTH_TOKEN_FILE and
+#     MINER_TLS_CERT_SHA256_FILE (or their inline variants) are mandatory —
+#     the miner exits at startup without them. Copy the node's
+#     miner-auth-token and miner-tls-cert-sha256 files somewhere this service
+#     can read (e.g. /etc/quantus-miner/): ProtectHome=true below blocks
+#     paths under /home and /root, where the node's base path usually lives.
+#   - The miner counts CPUs via the process affinity mask (so CPUAffinity=
+#     is respected) and emits a miner_effective_cpus metric.
+#   - The Prometheus exporter is ALWAYS on and binds plaintext HTTP on
+#     0.0.0.0:9900 by default; MINER_METRICS_PORT only changes the port.
+#     Firewall the port if the host is reachable from untrusted networks.
 
 [Unit]
 Description=Quantus External Miner Service
@@ -53,7 +63,9 @@ Environment=RUST_LOG=info
 
 # Launch the miner. All configuration should be passed via environment vars
 # (preferred) or via EXTRA_MINER_FLAGS in the env file for additional arguments.
-ExecStart=/usr/local/bin/quantus-miner $EXTRA_MINER_FLAGS
+# MINER_AUTH_TOKEN_FILE and MINER_TLS_CERT_SHA256_FILE must be set in the env
+# file (see readme.md) or the miner exits immediately.
+ExecStart=/usr/local/bin/quantus-miner serve $EXTRA_MINER_FLAGS
 
 # Graceful shutdown and restart behavior
 Restart=always

examples/systemd/readme.md

@@ -18,6 +18,18 @@ Prerequisites
 - The quantus-miner binary installed at /usr/local/bin/quantus-miner (or adjust ExecStart).
 - A service account (recommended):
   - sudo useradd --system --no-create-home --shell /usr/sbin/nologin quantus
+- The node's miner auth token and TLS cert fingerprint (required since the node
+  authenticates miners). The node creates both on first start with
+  --miner-listen-port and logs their paths (default:
+  `<base-path>/chains/<chain>/miner-auth-token` and
+  `<base-path>/chains/<chain>/miner-tls-cert-sha256`). Copy them where the
+  service can read them — the unit sets ProtectHome=true, so paths under /home
+  or /root are not readable:
+  - sudo install -d -m 0755 /etc/quantus-miner
+  - sudo install -m 0640 -o root -g quantus \
+      "<base-path>/chains/<chain>/miner-auth-token" /etc/quantus-miner/miner-auth-token
+  - sudo install -m 0644 \
+      "<base-path>/chains/<chain>/miner-tls-cert-sha256" /etc/quantus-miner/miner-tls-cert-sha256
 
 Install (unit)
 1) Copy the service file
@@ -26,21 +38,22 @@ Install (unit)
 2) Create a writable working directory (managed by systemd via StateDirectory)
    sudo install -d -o quantus -g quantus /var/lib/quantus-miner
 
-3) (Optional) Provide environment variables
+3) Provide environment variables (auth vars are REQUIRED)
    - Debian/Ubuntu:   sudoedit /etc/default/quantus-miner
    - RHEL/CentOS/Fed: sudoedit /etc/sysconfig/quantus-miner
 
-   Common variables (examples):
-   MINER_ENGINE=cpu
-   MINER_PORT=9833
-   MINER_METRICS_PORT=9900         # enable Prometheus exporter
-   MINER_WORKERS=4                 # leave unset to use default (50% of effective CPUs)
-   MINER_PROGRESS_CHUNK_MS=2000
-   # Throttling engine (cpu-chain-manipulator) knobs:
-   # MINER_MANIP_SOLVED_BLOCKS=0
-   # MINER_MANIP_BASE_DELAY_NS=500000
-   # MINER_MANIP_STEP_BATCH=10000
-   # MINER_MANIP_THROTTLE_CAP=0
+   Common variables (examples). Note: systemd EnvironmentFile= does not strip
+   inline comments — keep each assignment on its own line with nothing after
+   the value.
+   MINER_NODE_ADDR=127.0.0.1:9833
+   # Required: node auth token + TLS cert pin (see Prerequisites)
+   MINER_AUTH_TOKEN_FILE=/etc/quantus-miner/miner-auth-token
+   MINER_TLS_CERT_SHA256_FILE=/etc/quantus-miner/miner-tls-cert-sha256
+   # CPU worker threads; leave unset to auto-detect (~50% of available CPUs)
+   MINER_CPU_WORKERS=4
+   MINER_GPU_DEVICES=0
+   # Prometheus exporter port (the exporter is always on; this only changes the port)
+   MINER_METRICS_PORT=9900
    # Extra CLI flags (kept stable ExecStart):
    # EXTRA_MINER_FLAGS="--some-future-flag value"
 
@@ -65,31 +78,44 @@ Dedicated hardware (miner only)
   sudo systemctl restart quantus-miner
 
 Configuration reference (environment variables)
-- MINER_ENGINE
-  - cpu (default), cpu-chain-manipulator
-  - gpu for high-performance GPU mining
-- MINER_PORT
-  - HTTP API port (default 9833)
+- MINER_NODE_ADDR
+  - Address of the node's miner QUIC endpoint (default 127.0.0.1:9833).
+- MINER_AUTH_TOKEN_FILE (required, or MINER_AUTH_TOKEN inline)
+  - Path to a copy of the node's miner-auth-token file. The file variant is
+    preferred so the secret stays off the command line and out of `systemctl show`.
+- MINER_TLS_CERT_SHA256_FILE (required, or MINER_TLS_CERT_SHA256 inline)
+  - Path to a copy of the node's miner-tls-cert-sha256 file (the fingerprint is
+    also printed in the node's startup logs).
+- MINER_CPU_WORKERS
+  - CPU worker threads. If unset, auto-detected.
+- MINER_GPU_DEVICES
+  - Number of GPU devices to use. If unset, auto-detected.
+- MINER_GPU_BATCH_SIZE / MINER_CPU_BATCH_SIZE
+  - Nonces/hashes per cancellation check (defaults 1000000 / 10000).
+- MINER_GPU_THROTTLE_MS
+  - Delay between GPU batches in milliseconds (default 0 = no throttle).
 - MINER_METRICS_PORT
-  - Enable Prometheus exporter when set (e.g., 9900). If unset, metrics exporter is disabled.
-- MINER_WORKERS
-  - Worker threads (logical CPUs). If unset, defaults to ~50% of effective CPUs (clamped to [1, effective-1]).
-- MINER_PROGRESS_CHUNK_MS
-  - Target milliseconds for per-thread progress updates (default 2000ms).
-- Throttling engine (cpu-chain-manipulator) knobs
-  - MINER_MANIP_SOLVED_BLOCKS, MINER_MANIP_BASE_DELAY_NS, MINER_MANIP_STEP_BATCH, MINER_MANIP_THROTTLE_CAP
+  - Prometheus exporter port (default 9900). The exporter is ALWAYS on and
+    binds plaintext HTTP on all interfaces (0.0.0.0); this variable only
+    changes the port — there is no disable or loopback-only option. Firewall
+    the port or restrict it to your monitoring network.
+- MINER_ALLOW_INTEGRATED
+  - Allow integrated GPUs even when discrete GPUs are present.
 - EXTRA_MINER_FLAGS
   - Optional extra CLI flags appended to ExecStart.
 
 CPU affinity, cpusets, and workers
-- The miner detects the effective CPU capacity (logical CPUs) visible to the process by preferring cgroup v2 cpuset (cpuset.cpus.effective), falling back to v1, else using all logical CPUs.
-- At startup (debug level), the miner logs the detected cpuset mask (if any).
-- A Prometheus gauge miner_effective_cpus is emitted (when metrics are enabled) with the effective count for dashboards/alerts.
-- If --workers (or MINER_WORKERS) exceeds effective CPUs, it is clamped and a warning is logged.
-- If omitted, the miner defaults to ~50% of effective CPUs (but always at least 1 and less than or equal to effective-1).
+- The miner counts available CPUs via the process CPU affinity mask (num_cpus),
+  so a systemd CPUAffinity= setting (or a cgroup cpuset) is reflected in the count.
+- If MINER_CPU_WORKERS is unset, the miner uses ~50% of the detected CPUs
+  (at least 1) and logs the choice at startup ("Auto-detected N CPU workers").
+- An explicit MINER_CPU_WORKERS value is used as-is — it is NOT clamped to the
+  affinity mask, so an oversized value oversubscribes the pinned CPUs.
+- A Prometheus gauge miner_effective_cpus is emitted with the detected count for dashboards/alerts.
 - When pinning CPUAffinity at the systemd level:
   - Ensure CPUAffinity is a subset of the cgroup cpuset mask.
-  - Consider setting MINER_WORKERS to match the number of CPUs in the affinity mask if you want full utilization, or rely on the default 50% policy.
+  - Set MINER_CPU_WORKERS to the number of CPUs in the affinity mask for full
+    utilization, or leave it unset for the ~50% default.
 
 Security hardening (in the unit)
 - NoNewPrivileges=true
@@ -101,6 +127,9 @@ Security hardening (in the unit)
 - RestrictSUIDSGID=true
 - SystemCallFilter=@system-service
 Adjust or relax as needed for your environment.
+Note: the Prometheus exporter always listens on 0.0.0.0:<MINER_METRICS_PORT>
+(default 9900) — the unit cannot disable it or bind it to loopback. Firewall
+the port if the host is reachable from untrusted networks.
 
 Validation and troubleshooting
 - Check service status and logs:
@@ -110,20 +139,28 @@ Validation and troubleshooting
   taskset -cp "$pid"
 - Verify cpuset mask (cgroup v2):
   cat /sys/fs/cgroup/cpuset.cpus.effective
-- Metrics:
-  - If MINER_METRICS_PORT is set, curl http://127.0.0.1:<port>/metrics
+- Metrics (always on, default port 9900):
+  - curl http://127.0.0.1:${MINER_METRICS_PORT:-9900}/metrics
   - Look for miner_effective_cpus and per-job/thread metrics.
 - Common pitfalls:
   - ExecStart path wrong (ensure /usr/local/bin/quantus-miner exists and is executable).
   - Service user/group missing (create quantus or adjust User/Group).
+  - Miner exits immediately with "miner auth token required" / "TLS cert fingerprint required":
+    set MINER_AUTH_TOKEN_FILE and MINER_TLS_CERT_SHA256_FILE in the env file (see Prerequisites).
+  - Auth/TLS files unreadable: ProtectHome=true blocks /home and /root; copy the
+    files to /etc/quantus-miner/ and make them readable by the quantus user.
+  - Node rejects the miner after the node's base path changed or its credential
+    files (miner-auth-token, miner-tls-cert.der/-key.der) were deleted and
+    regenerated: re-copy both files and restart. (A plain purge-chain does NOT
+    rotate them — it removes only the database, so the same token and cert are
+    reloaded on the next start.)
   - CPUAffinity not a subset of the cgroup cpuset (adjust cpuset or affinity).
-  - MINER_ENGINE set to gpu-* (currently unimplemented, exits with clear error).
   - Insufficient permissions to write WorkingDirectory (systemd StateDirectory creates /var/lib/quantus-miner with correct ownership).
 
 Operational tips
-- For shared machines: prefer 10-shared-hardware.conf and leave MINER_WORKERS unset (defaults to ~50%).
-- For dedicated machines: use 20-dedicated-hardware.conf and set MINER_WORKERS to the number of CPUs in CPUAffinity (or omit CPUAffinity to inherit cpuset).
-- Use RUST_LOG=info,miner=debug temporarily to verify startup detection (cpuset mask, effective CPUs) and to observe mining loop behavior; then turn back down to reduce log volume.
+- For shared machines: prefer 10-shared-hardware.conf and leave MINER_CPU_WORKERS unset (auto-detect).
+- For dedicated machines: use 20-dedicated-hardware.conf and set MINER_CPU_WORKERS to the number of CPUs in CPUAffinity (or omit CPUAffinity to inherit cpuset).
+- Use RUST_LOG=info,miner=debug temporarily to verify startup detection (worker auto-detection, GPU discovery) and to observe mining loop behavior; then turn back down to reduce log volume.
 
 Support
 - Repository: https://github.com/Quantus-Network/quantus-miner

rust-toolchain

@@ -1,4 +1,4 @@
 [toolchain]
-channel = "stable"
+channel = "1.93.0"
 components = ["clippy", "rustfmt"]
 profile = "minimal"