bump version to v4.0.0 (#85 )

Co-authored-by: n13 <n13@users.noreply.github.com>
Send miner auth token in Ready handshake. (#78 )
2026-08-13 19:34:16 +08:00 · 2026-08-13 09:50:03 +08:00 · 2026-08-12 15:41:38 +08:00 · 2026-08-12 14:24:05 +08:00 · 2026-08-12 12:10:58 +08:00 · 2026-08-12 12:04:44 +08:00
41 changed files with 4853 additions and 1223 deletions
@@ -1,42 +0,0 @@
-# Git
-.git
-.gitignore
-.github
-
-# Build artifacts
-target/
-*.rs.bk
-*.pdb
-
-# IDE
-.vscode
-.idea
-*.swp
-*.swo
-*~
-
-# Documentation
-docs/
-*.md
-!README.md
-
-# Examples and tests (not needed for runtime)
-examples/
-
-# Other
-.DS_Store
-Thumbs.db
-
-# CI/CD
-.github/
-
-# License
-LICENSE
-
-# Dockerfiles themselves
-Dockerfile*
-.dockerignore
-
-# Cargo cache
-.cargo/
-
@@ -1,145 +0,0 @@
-name: Docker Image
-
-on:
-  workflow_dispatch:
-    inputs:
-      specific_version:
-        description: "Optional: Specify a full version (e.g., v0.3.0 or v0.3.1-beta.1) to build. If empty, uses the latest release tag. MUST start with 'v'."
-        required: false
-        default: ""
-
-env:
-  CARGO_TERM_COLOR: always
-
-jobs:
-  determine_version:
-    runs-on: ubuntu-latest
-    outputs:
-      version_with_v: ${{ steps.version_info.outputs.version_with_v }}
-      is_valid_version: ${{ steps.version_info.outputs.is_valid_version }}
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-        with:
-          fetch-depth: 0
-
-      - name: Get version information
-        id: version_info
-        run: |
-          set -ex
-          TARGET_VERSION=""
-          SPECIFIC_VERSION="${{ github.event.inputs.specific_version }}"
-
-          if [[ -n "$SPECIFIC_VERSION" ]]; then
-            if [[ ! "$SPECIFIC_VERSION" =~ ^v[0-9]+\.[0-9]+\.[0-9]+(-[a-zA-Z0-9.-]+)?$ ]]; then
-              echo "::error::Specified version '$SPECIFIC_VERSION' is not a valid format. It must start with 'v' (e.g., v0.3.0 or v0.3.1-beta.1)."
-              echo "is_valid_version=false" >> $GITHUB_OUTPUT
-              exit 1
-            fi
-            TARGET_VERSION="$SPECIFIC_VERSION"
-            echo "Using specified version: $TARGET_VERSION"
-          else
-            echo "No specific version provided, determining latest release tag..."
-            TARGET_VERSION=$(git tag --list 'v*' --sort=-v:refname | head -n 1)
-            if [[ -z "$TARGET_VERSION" ]]; then
-              echo "::error::No version tags starting with 'v' (e.g., vX.Y.Z) found in the repository."
-              echo "is_valid_version=false" >> $GITHUB_OUTPUT
-              exit 1
-            fi
-            echo "Latest release version found: $TARGET_VERSION"
-          fi
-
-          echo "version_with_v=$TARGET_VERSION" >> $GITHUB_OUTPUT
-          echo "is_valid_version=true" >> $GITHUB_OUTPUT
-
-  build_and_publish_image:
-    needs: determine_version
-    if: needs.determine_version.outputs.is_valid_version == 'true'
-    runs-on: ubuntu-latest
-    permissions:
-      contents: read
-      packages: write
-    env:
-      GHCR_IMAGE_PATH: ghcr.io/quantus-network/quantus-miner
-      TARGET_VERSION_WITH_V: ${{ needs.determine_version.outputs.version_with_v }}
-
-    steps:
-      - name: Set up QEMU
-        uses: docker/setup-qemu-action@v3
-    
-      - name: Set up Docker Buildx
-        uses: docker/setup-buildx-action@v3
-
-      - name: Log in to GitHub Container Registry
-        uses: docker/login-action@v3
-        with:
-          registry: ghcr.io
-          username: ${{ github.actor }}
-          password: ${{ secrets.GITHUB_TOKEN }}
-
-      - name: Check if Docker image already exists and fail if so
-        id: check_image
-        run: |
-          set -ex
-          IMAGE_TO_CHECK="${{ env.GHCR_IMAGE_PATH }}:${{ env.TARGET_VERSION_WITH_V }}"
-          echo "Checking for image: $IMAGE_TO_CHECK"
-          
-          if docker manifest inspect "$IMAGE_TO_CHECK" > /dev/null 2>&1; then
-            echo "::error::Image $IMAGE_TO_CHECK already exists. Aborting."
-            exit 1
-          else
-            echo "Image $IMAGE_TO_CHECK does not exist. Proceeding with build."
-          fi
-
-      - name: Checkout code at tag version
-        uses: actions/checkout@v4
-        with:
-          ref: ${{ env.TARGET_VERSION_WITH_V }}
-
-      - name: Check if Dockerfile exists in tag
-        id: check_dockerfile
-        run: |
-          if [ ! -f Dockerfile ]; then
-            echo "::warning::Dockerfile not found in tag ${{ env.TARGET_VERSION_WITH_V }}, falling back to main branch Dockerfile"
-            echo "use_main=true" >> $GITHUB_OUTPUT
-          else
-            echo "Dockerfile found in tag"
-            echo "use_main=false" >> $GITHUB_OUTPUT
-          fi
-
-      - name: Checkout main branch for Dockerfile (if needed)
-        if: steps.check_dockerfile.outputs.use_main == 'true'
-        uses: actions/checkout@v4
-        with:
-          ref: main
-          sparse-checkout: |
-            Dockerfile
-            .dockerignore
-          sparse-checkout-cone-mode: false
-          path: dockerfile-source
-
-      - name: Copy Dockerfile from main (if needed)
-        if: steps.check_dockerfile.outputs.use_main == 'true'
-        run: |
-          cp dockerfile-source/Dockerfile ./
-          [ -f dockerfile-source/.dockerignore ] && cp dockerfile-source/.dockerignore ./ || true
-
-      - name: Build and push Docker image
-        uses: docker/build-push-action@v5
-        with:
-          context: .
-          file: ./Dockerfile
-          push: true
-          platforms: linux/amd64,linux/arm64
-          provenance: false
-          tags: |
-            ${{ env.GHCR_IMAGE_PATH }}:${{ env.TARGET_VERSION_WITH_V }}
-            ${{ env.GHCR_IMAGE_PATH }}:latest
-          cache-from: type=gha
-          cache-to: type=gha,mode=max
-
-      - name: Print completion message
-        run: |
-          echo "Successfully built and pushed ${{ env.GHCR_IMAGE_PATH }}:${{ env.TARGET_VERSION_WITH_V }}"
-          echo "Image is also tagged as latest: ${{ env.GHCR_IMAGE_PATH }}:latest"
-
@@ -377,6 +377,12 @@ dependencies = [
 "unicode-xid",
 ]

+[[package]]
+name = "constant_time_eq"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7c74b8349d32d297c9134b8c88677813a227df8f779daa29bfc29c183fe3dca6"
+
 [[package]]
 name = "core-foundation"
 version = "0.9.4"
@@ -571,7 +577,7 @@ dependencies = [

 [[package]]
 name = "engine-cpu"
-version = "3.2.0"
+version = "4.0.0"
 dependencies = [
 "criterion",
 "hex",
@@ -582,7 +588,7 @@ dependencies = [

 [[package]]
 name = "engine-gpu"
-version = "3.2.0"
+version = "4.0.0"
 dependencies = [
 "bytemuck",
 "criterion",
@@ -596,9 +602,11 @@ dependencies = [
 "qp-plonky2",
 "qp-plonky2-field",
 "qp-poseidon-constants",
- "qp-poseidon-core",
+ "qp-poseidon-core 3.1.0",
 "rand 0.9.2",
 "rand_chacha 0.9.0",
+ "regex",
+ "tokio",
 "wgpu",
 ]

@@ -628,7 +636,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "39cab71617ae0d63f51a36d69f866391735b51691dbda63cf6f96d042b63efeb"
 dependencies = [
 "libc",
- "windows-sys 0.61.2",
+ "windows-sys 0.52.0",
 ]

 [[package]]
@@ -1268,7 +1276,7 @@ checksum = "3640c1c38b8e4e43584d8df18be5fc6b0aa314ce6ebf51b53313d4306cca8e46"
 dependencies = [
 "hermit-abi",
 "libc",
- "windows-sys 0.61.2",
+ "windows-sys 0.52.0",
 ]

 [[package]]
@@ -1468,7 +1476,7 @@ dependencies = [

 [[package]]
 name = "metrics"
-version = "3.2.0"
+version = "4.0.0"
 dependencies = [
 "anyhow",
 "log",
@@ -1496,7 +1504,7 @@ dependencies = [

 [[package]]
 name = "miner-cli"
-version = "3.2.0"
+version = "4.0.0"
 dependencies = [
 "clap",
 "engine-cpu",
@@ -1507,13 +1515,14 @@ dependencies = [
 "miner-service",
 "num_cpus",
 "primitive-types 0.13.1",
+ "quic-transport",
 "rand 0.9.2",
 "tokio",
 ]

 [[package]]
 name = "miner-service"
-version = "3.2.0"
+version = "4.0.0"
 dependencies = [
 "anyhow",
 "crossbeam-channel",
@@ -1527,14 +1536,14 @@ dependencies = [
 "pow-core",
 "primitive-types 0.13.1",
 "quantus-miner-api",
+ "quic-transport",
 "quinn",
- "rustls 0.21.12",
 "tokio",
 ]

 [[package]]
 name = "miner-telemetry"
-version = "3.2.0"
+version = "4.0.0"
 dependencies = [
 "anyhow",
 "futures",
@@ -1992,6 +2001,28 @@ dependencies = [
 "plotters-backend",
 ]

+[[package]]
+name = "pool-service"
+version = "4.0.0"
+dependencies = [
+ "anyhow",
+ "clap",
+ "constant_time_eq",
+ "env_logger",
+ "hex",
+ "log",
+ "pow-core",
+ "primitive-types 0.13.1",
+ "quantus-miner-api",
+ "quic-transport",
+ "quinn",
+ "rand 0.9.2",
+ "serde",
+ "serde_json",
+ "tokio",
+ "warp",
+]
+
 [[package]]
 name = "portable-atomic"
 version = "1.13.1"
@@ -2018,10 +2049,11 @@ dependencies = [

 [[package]]
 name = "pow-core"
-version = "3.2.0"
+version = "4.0.0"
 dependencies = [
 "hex",
 "primitive-types 0.13.1",
+ "qp-poseidon-core 3.1.0",
 "qpow-math",
 ]

@@ -2229,6 +2261,12 @@ version = "2.1.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "78172860cd960773c72e97fba07ead9e5a1c13086c5746283744933e2e4a577b"

+[[package]]
+name = "qp-poseidon-core"
+version = "3.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5872607e25ea4ee5fb37e64bf1462168e1a36a4e719cdc8a105533c708253918"
+
 [[package]]
 name = "qpow-math"
 version = "0.1.0"
@@ -2237,19 +2275,33 @@ dependencies = [
 "hex",
 "log",
 "primitive-types 0.13.1",
- "qp-poseidon-core",
+ "qp-poseidon-core 2.1.0",
 ]

 [[package]]
 name = "quantus-miner-api"
-version = "0.2.0"
-source = "git+https://github.com/Quantus-Network/chain.git?tag=v0.7.1-q-day-2#5ea94cb1e5347b12d091b6fa9f9234ceb897ec42"
+version = "0.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "725bbf1de0bc7d2ace83f8e49ea754be5f45b84bc9471ddd3f417a9e2bb20553"
 dependencies = [
 "serde",
 "serde_json",
 "tokio",
 ]

+[[package]]
+name = "quic-transport"
+version = "4.0.0"
+dependencies = [
+ "anyhow",
+ "quantus-miner-api",
+ "quinn",
+ "rustls 0.21.12",
+ "serde_json",
+ "sha2",
+ "tokio",
+]
+
 [[package]]
 name = "quinn"
 version = "0.10.2"
@@ -2749,6 +2801,17 @@ dependencies = [
 "digest",
 ]

+[[package]]
+name = "sha2"
+version = "0.10.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a7507d819769d01a365ab707794a4084392c824f54a7a6a7862f8c3d0892b283"
+dependencies = [
+ "cfg-if",
+ "cpufeatures 0.2.17",
+ "digest",
+]
+
 [[package]]
 name = "shlex"
 version = "1.3.0"
@@ -2806,6 +2869,14 @@ dependencies = [
 "windows-sys 0.61.2",
 ]

+[[package]]
+name = "solver-wasm"
+version = "4.0.0"
+dependencies = [
+ "primitive-types 0.13.1",
+ "qpow-math",
+]
+
 [[package]]
 name = "spin"
 version = "0.5.2"
@@ -3691,7 +3762,7 @@ version = "0.1.11"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "c2a7b1c03c876122aa43f3020e6c3c3ee5c05081c9a00739faf7503aeba10d22"
 dependencies = [
- "windows-sys 0.61.2",
+ "windows-sys 0.48.0",
 ]

 [[package]]
@@ -6,7 +6,10 @@ members = [
 	"crates/miner-cli",
 	"crates/miner-service",
 	"crates/miner-telemetry",
+	"crates/pool-service",
 	"crates/pow-core",
+	"crates/quic-transport",
+	"crates/solver-wasm",
 ]
 default-members = ["crates/miner-cli"]
 resolver = "2"
@@ -15,7 +18,7 @@ resolver = "2"
 edition = "2021"
 authors = ["Quantus Network"]
 description = "Quantus External Miner Workspace"
-version = "3.2.0"
+version = "4.0.0"

 [workspace.dependencies]
 anyhow = "1"
@@ -29,9 +32,9 @@ num-bigint = { version = "0.4", features = ["rand"] }
 num-traits = "0.2"
 num_cpus = "1.16"
 primitive-types = { version = "0.13.1", default-features = false }
-qp-poseidon-core = { version = "2.1.0", default-features = false }
+qp-poseidon-core = { version = "3.1.0", default-features = false }
 qpow-math = { git = "https://github.com/Quantus-Network/chain.git", tag = "v0.7.1-q-day-2", package = "qpow-math", default-features = false }
-quantus-miner-api = { git = "https://github.com/Quantus-Network/chain.git", tag = "v0.7.1-q-day-2" }
+quantus-miner-api = "0.3.0"
 rand = { version = "0.9", default-features = false, features = ["std", "std_rng", "thread_rng"] }
 serde = { version = "1.0", features = ["derive"] }
 serde_json = { version = "1.0.132", default-features = false }
@@ -1,55 +0,0 @@
-# syntax=docker/dockerfile:1
-
-############################
-# Builder stage
-############################
-FROM rust:1.85-slim-bookworm AS builder
-
-# Install build dependencies
-RUN apt-get update \
- && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \
-    pkg-config \
-    libssl-dev \
-    ca-certificates \
- && rm -rf /var/lib/apt/lists/*
-
-WORKDIR /build
-
-# Copy workspace files
-COPY Cargo.toml Cargo.lock rust-toolchain taplo.toml ./
-COPY crates ./crates
-COPY tests ./tests
-
-# Build the miner-cli in release mode
-RUN cargo build --release -p miner-cli --locked
-
-# Strip debug symbols to reduce binary size
-RUN strip target/release/quantus-miner || true
-
-############################
-# Runtime-only stage
-############################
-FROM debian:bookworm-slim
-
-# Install runtime dependencies
-RUN apt-get update \
- && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \
-    ca-certificates \
- && rm -rf /var/lib/apt/lists/*
-
-# Copy binary from builder stage
-COPY --from=builder /build/target/release/quantus-miner /usr/local/bin/quantus-miner
-
-# Expose miner API port and metrics port
-EXPOSE 9833 9900
-
-# Run as unprivileged user
-RUN useradd --system --uid 10001 quantus
-USER 10001:10001
-
-# Default working directory
-WORKDIR /data
-
-# Start the miner
-ENTRYPOINT ["quantus-miner"]
-
@@ -1,235 +0,0 @@
-# External Miner Protocol Specification
-
-This document defines the QUIC-based protocol for communication between the Quantus Network node and external QPoW miner services.
-
-## Overview
-
-The node delegates the mining task (finding a valid nonce) to external miner services over persistent QUIC connections. The node provides the necessary parameters (header hash, difficulty) and each external miner independently searches for a valid nonce according to the QPoW rules defined in the `qpow-math` crate. Miners push results back when found.
-
-### Key Benefits of QUIC
-
- **Lower latency**: Results are pushed immediately when found (no polling)
- **Connection resilience**: Built-in connection migration and recovery
- **Multiplexed streams**: Multiple operations on single connection
- **Built-in TLS**: Encrypted by default
-
-## Architecture
-
-### Connection Model
-
-```
-                           ┌─────────────────────────────────┐
-                           │            Node                 │
-                           │   (QUIC Server on port 9833)    │
-                           │                                 │
-┌──────────┐               │  Broadcasts: NewJob             │
-│  Miner 1 │ ──connect───► │  Receives: JobResult            │
-└──────────┘               │                                 │
-                           │  Supports multiple miners       │
-┌──────────┐               │  First valid result wins        │
-│  Miner 2 │ ──connect───► │                                 │
-└──────────┘               └─────────────────────────────────┘
-                           
-┌──────────┐                         
-│  Miner 3 │ ──connect───►           
-└──────────┘                         
-```
-
- **Node** acts as the QUIC server, listening on port 9833 (default)
- **Miners** act as QUIC clients, connecting to the node
- Single bidirectional stream per miner connection
- Connection persists across multiple mining jobs
- Multiple miners can connect simultaneously
-
-### Multi-Miner Operation
-
-When multiple miners are connected:
-1. Node broadcasts the same `NewJob` to all connected miners
-2. Each miner independently selects a random starting nonce
-3. First miner to find a valid solution sends `JobResult`
-4. Node uses the first valid result, ignores subsequent results for same job
-5. New job broadcast implicitly cancels work on all miners
-
-### Message Types
-
-The protocol uses only **two message types**:
-
-| Direction | Message | Description |
-|-----------|---------|-------------|
-| Node → Miner | `NewJob` | Submit a mining job (implicitly cancels any previous job) |
-| Miner → Node | `JobResult` | Mining result (completed, failed, or cancelled) |
-
-### Wire Format
-
-Messages are length-prefixed JSON:
-
-```
-┌─────────────────┬─────────────────────────────────┐
-│ Length (4 bytes)│ JSON payload (MinerMessage)     │
-│ big-endian u32  │                                 │
-└─────────────────┴─────────────────────────────────┘
-```
-
-Maximum message size: 16 MB
-
-## Data Types
-
-See the `quantus-miner-api` crate for the canonical Rust definitions.
-
-### MinerMessage (Enum)
-
-```rust
-pub enum MinerMessage {
-    NewJob(MiningRequest),
-    JobResult(MiningResult),
-}
-```
-
-### MiningRequest
-
-| Field | Type | Description |
-|-------|------|-------------|
-| `job_id` | String | Unique identifier (UUID recommended) |
-| `mining_hash` | String | Header hash (64 hex chars, no 0x prefix) |
-| `distance_threshold` | String | Difficulty (U512 as decimal string) |
-
-Note: Nonce range is not specified - each miner independently selects a random starting point.
-
-### MiningResult
-
-| Field | Type | Description |
-|-------|------|-------------|
-| `status` | ApiResponseStatus | Result status (see below) |
-| `job_id` | String | Job identifier |
-| `nonce` | Option<String> | Winning nonce (U512 hex, no 0x prefix) |
-| `work` | Option<String> | Winning nonce as bytes (128 hex chars) |
-| `hash_count` | u64 | Number of nonces checked |
-| `elapsed_time` | f64 | Time spent mining (seconds) |
-
-### ApiResponseStatus (Enum)
-
-| Value | Description |
-|-------|-------------|
-| `completed` | Valid nonce found |
-| `failed` | Nonce range exhausted without finding solution |
-| `cancelled` | Job was cancelled (new job received) |
-| `running` | Job still in progress (not typically sent) |
-
-## Protocol Flow
-
-### Normal Mining Flow
-
-```
-Miner                                        Node
-  │                                            │
-  │──── QUIC Connect ─────────────────────────►│
-  │◄─── Connection Established ────────────────│
-  │                                            │
-  │◄─── NewJob { job_id: "abc", ... } ─────────│
-  │                                            │
-  │     (picks random nonce, starts mining)    │
-  │                                            │
-  │──── JobResult { job_id: "abc", ... } ─────►│ (found solution!)
-  │                                            │
-  │     (node submits block, gets new work)    │
-  │                                            │
-  │◄─── NewJob { job_id: "def", ... } ─────────│
-  │                                            │
-```
-
-### Job Cancellation (Implicit)
-
-When a new block arrives before the miner finds a solution, the node simply sends a new `NewJob`. The miner automatically cancels the previous job:
-
-```
-Miner                                        Node
-  │                                            │
-  │◄─── NewJob { job_id: "abc", ... } ─────────│
-  │                                            │
-  │     (mining "abc")                         │
-  │                                            │
-  │     (new block arrives at node!)           │
-  │                                            │
-  │◄─── NewJob { job_id: "def", ... } ─────────│
-  │                                            │
-  │     (cancels "abc", starts "def")          │
-  │                                            │
-  │──── JobResult { job_id: "def", ... } ─────►│
-```
-
-### Miner Connect During Active Job
-
-When a miner connects while a job is active, it immediately receives the current job:
-
-```
-Miner (new)                                  Node
-  │                                            │ (already mining job "abc")
-  │──── QUIC Connect ─────────────────────────►│
-  │◄─── Connection Established ────────────────│
-  │◄─── NewJob { job_id: "abc", ... } ─────────│ (current job sent immediately)
-  │                                            │
-  │     (joins mining effort)                  │
-```
-
-### Stale Result Handling
-
-If a result arrives for an old job, the node discards it:
-
-```
-Miner                                        Node
-  │                                            │
-  │◄─── NewJob { job_id: "abc", ... } ─────────│
-  │                                            │
-  │◄─── NewJob { job_id: "def", ... } ─────────│ (almost simultaneous)
-  │                                            │
-  │──── JobResult { job_id: "abc", ... } ─────►│ (stale, node ignores)
-  │                                            │
-  │──── JobResult { job_id: "def", ... } ─────►│ (current, node uses)
-```
-
-## Configuration
-
-### Node
-
-```bash
-# Listen for external miner connections on port 9833
-quantus-node --miner-listen-port 9833
-```
-
-### Miner
-
-```bash
-# Connect to node
-quantus-miner serve --node-addr 127.0.0.1:9833
-```
-
-## TLS Configuration
-
-The node generates a self-signed TLS certificate at startup. The miner skips certificate verification by default (insecure mode). For production deployments, consider:
-
-1. **Certificate pinning**: Configure the miner to accept only specific certificate fingerprints
-2. **Proper CA**: Use certificates signed by a trusted CA
-3. **Network isolation**: Run node and miner on a private network
-
-## Error Handling
-
-### Connection Loss
-
-The miner automatically reconnects with exponential backoff:
- Initial delay: 1 second
- Maximum delay: 30 seconds
-
-The node continues operating with remaining connected miners.
-
-### Validation Errors
-
-If the miner receives an invalid `MiningRequest`, it sends a `JobResult` with status `failed`.
-
-## Notes
-
- All hex values should be sent **without** the `0x` prefix
- The miner implements validation logic from `qpow_math::is_valid_nonce`
- The node uses the `work` field from `MiningResult` to construct `QPoWSeal`
- ALPN protocol identifier: `quantus-miner`
- Each miner independently generates a random nonce starting point using cryptographically secure randomness
- With a 512-bit nonce space, collision between miners is statistically impossible
@@ -16,18 +16,34 @@ The binary will be available at `target/release/quantus-miner`.

 ## Running

-```bash
-# CPU-only mining (default: auto-detected CPU cores)
-./target/release/quantus-miner serve --cpu-workers 4
+The node requires a shared auth token and TLS cert pin. Both live under the
+node's chain config dir (`<base-path>/chains/<chain>/`):

-# GPU-only mining 
-./target/release/quantus-miner serve --gpu-devices 1
+- `miner-auth-token` — shared secret (**not** logged by the node; read this file)
+- `miner-tls-cert-sha256` — SHA-256 of the miner TLS cert (also printed in node logs)
+
+```bash
+# Preferred: mount/read the node's chain config files
+./target/release/quantus-miner serve \
+  --node-addr 127.0.0.1:9833 \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
+  --cpu-workers 4
+
+# Or pass values directly (token from miner-auth-token; fingerprint also in node logs)
+./target/release/quantus-miner serve \
+  --node-addr 127.0.0.1:9833 \
+  --auth-token <TOKEN> \
+  --tls-cert-sha256 <FINGERPRINT> \
+  --gpu-devices 1

 # Hybrid CPU+GPU mining
-./target/release/quantus-miner serve --cpu-workers 4 --gpu-devices 1
-
-# Custom port and metrics
-./target/release/quantus-miner serve --cpu-workers 2 --port 8000 --metrics-port 9900
+./target/release/quantus-miner serve \
+  --node-addr 127.0.0.1:9833 \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
+  --cpu-workers 4 \
+  --gpu-devices 1
 ```

 ## Configuration
@@ -35,6 +51,10 @@ The binary will be available at `target/release/quantus-miner`.
 | Argument | Environment Variable | Description | Default |
 |----------|---------------------|-------------|---------|
 | `--node-addr <ADDR>` | `MINER_NODE_ADDR` | Node address to connect to | `127.0.0.1:9833` |
+| `--auth-token <TOKEN>` | `MINER_AUTH_TOKEN` | Shared secret from the node's `miner-auth-token` file (not logged) | required |
+| `--auth-token-file <PATH>` | `MINER_AUTH_TOKEN_FILE` | Read the shared secret from a file (preferred) | — |
+| `--tls-cert-sha256 <HEX>` | `MINER_TLS_CERT_SHA256` | SHA-256 of the node's miner TLS cert (`miner-tls-cert-sha256` / node logs) | required |
+| `--tls-cert-sha256-file <PATH>` | `MINER_TLS_CERT_SHA256_FILE` | Read the TLS cert fingerprint from a file | — |
 | `--cpu-workers <N>` | `MINER_CPU_WORKERS` | Number of CPU worker threads | Auto-detect |
 | `--gpu-devices <N>` | `MINER_GPU_DEVICES` | Number of GPU devices | Auto-detect |
 | `--gpu-batch-size <N>` | `MINER_GPU_BATCH_SIZE` | GPU batch size in nonces | 1000000 |
@@ -70,27 +90,46 @@ cargo build -p miner-cli --release

 ## Examples

+All `serve` examples need the auth token and TLS pin (files or inline values).
+
 ```bash
 # CPU mining with 8 workers
-./target/release/quantus-miner serve --cpu-workers 8
+./target/release/quantus-miner serve \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
+  --cpu-workers 8

 # Pure GPU mining
-./target/release/quantus-miner serve --gpu-devices 1
+./target/release/quantus-miner serve \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
+  --gpu-devices 1

 # GPU mining with throttle (reduce GPU utilization)
-./target/release/quantus-miner serve --gpu-devices 1 --gpu-throttle-ms 50
+./target/release/quantus-miner serve \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
+  --gpu-devices 1 --gpu-throttle-ms 50

 # Hybrid mining: 4 CPU + 1 GPU workers
-./target/release/quantus-miner serve --cpu-workers 4 --gpu-devices 1
+./target/release/quantus-miner serve \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
+  --cpu-workers 4 --gpu-devices 1

 # With verbose logging
-RUST_LOG=debug ./target/release/quantus-miner serve --cpu-workers 2 --gpu-devices 1
+RUST_LOG=debug ./target/release/quantus-miner serve \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
+  --cpu-workers 2 --gpu-devices 1

 # Production setup with metrics
 ./target/release/quantus-miner serve \
+  --node-addr 127.0.0.1:9833 \
+  --auth-token-file /path/to/miner-auth-token \
+  --tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
  --cpu-workers 6 \
  --gpu-devices 1 \
-  --port 9833 \
  --metrics-port 9900
 ```

@@ -98,25 +137,13 @@ RUST_LOG=debug ./target/release/quantus-miner serve --cpu-workers 2 --gpu-device

 The miner uses a QUIC-based protocol for communication with the node:

- **Transport**: QUIC with TLS 1.3 (self-signed certificates)
+- **Transport**: QUIC with TLS 1.3 (self-signed certificate, pinned by SHA-256)
+- **Auth**: `Ready { token }` must match the node's `miner-auth-token`
+- **ALPN**: `quantus-miner/2`
 - **Port**: 9833 (default)
- **Messages**: `NewJob` (from node) and `JobResult` (from miner)
+- **Messages**: `Ready` (miner→node), `NewJob` (node→miner), `JobResult` (miner→node)

-For full protocol specification, see `EXTERNAL_MINER_PROTOCOL.md`.
-
-## Docker
-
-```bash
-# Quick start
-docker pull ghcr.io/quantus-network/quantus-miner:latest
-docker run -d -p 9833:9833 -p 9900:9900 \
-  ghcr.io/quantus-network/quantus-miner:latest \
-  --cpu-workers 4 --metrics-port 9900
-
-# Build from source
-docker build -t quantus-miner .
-docker run -d -p 9833:9833 quantus-miner serve --cpu-workers 4
-```
+For full protocol specification, see the node's `MINING.md`.

 ## Benchmarking

@@ -0,0 +1,80 @@
+# Security Policy
+
+Quantus Network takes the security of our users seriously and welcomes reports
+from security researchers.
+
+## Reporting a vulnerability
+
+**Please do not open a public issue, pull request, or social-media post for a
+security vulnerability.** Public disclosure before a fix is available puts
+users' funds at risk. Use one of the private channels below and we will
+coordinate a fix and disclosure with you.
+
+1. **GitHub private security advisory (preferred).**
+   [Report a vulnerability](https://github.com/quantus-network/quantus-miner/security/advisories/new).
+   This gives you a private, structured thread with the maintainers and is the
+   fastest way to reach us.
+2. **Email.** Send details to **security@quantus.com**.
+
+### What to include
+
+- A clear description of the issue and its security impact.
+- A working proof of concept is required. It must run directly against this
+  project, using a release build or locally built version, and demonstrate the
+  reported behavior and security impact. Standalone code, mathematical
+  examples, or simulations that only reproduce the theory without exercising
+  the project do not satisfy this requirement.
+- Step-by-step reproduction instructions.
+- Affected platforms and versions.
+- Any relevant logs, addresses, or transaction IDs (for on-chain issues).
+
+If you used AI tooling to find or write up the report, please say so.
+
+## Our commitment (safe harbor)
+
+We consider security research conducted in good faith under this policy to be
+authorized. We will not pursue or support legal action against researchers who:
+
+- make a good-faith effort to avoid privacy violations, data destruction, and
+  interruption or degradation of our services;
+- only interact with accounts they own or have explicit permission to access;
+  and
+- give us a reasonable opportunity to fix an issue before disclosing it
+  publicly.
+
+If in doubt about whether an action is authorized, ask us first at
+security@quantus.com.
+
+## What to expect
+
+- **Acknowledgement:** within **2 business days**.
+- **Triage and initial assessment:** within **7 business days**.
+- **Coordinated disclosure:** we aim to ship a fix and coordinate public
+  disclosure within **90 days** of the report. We will keep you updated on
+  progress and agree on a disclosure date with you.
+- **Credit:** with your permission, we are happy to publicly credit you for the
+  report once a fix is released.
+
+## Scope
+
+**In scope:** the external mining service (CPU and GPU) in this repository — anything that could lead to theft or misattribution of mining rewards, invalid blocks being produced, or remote code execution or crash via its service API.
+
+**Out of scope:** issues in third-party services or nodes we do not operate;
+reports generated solely by automated scanners without a demonstrated impact;
+low-severity or informational issues on our marketing and landing websites;
+and social-engineering or physical attacks.
+
+## Rewards
+
+At our **sole discretion**, we may offer a reward for a valid report. To be
+eligible, a report must be submitted **privately** through one of the channels
+above and identify a genuine vulnerability with real impact on users. Trivial
+or low-impact findings, automated-scanner output without a working proof of
+concept, and already-known issues are **not** eligible. There is no fixed
+bounty and no guaranteed payout; whether a report qualifies, and any amount,
+are determined solely by Quantus Network.
+
+## Supported versions
+
+Only the latest release is supported; security fixes are delivered in new
+versions. Please keep your installation up to date.
@@ -0,0 +1,7 @@
+#!/usr/bin/env bash
+# Run the same clippy command as CI (see .github/workflows/ci.yml)
+set -euo pipefail
+
+cargo +nightly fmt
+taplo format
+cargo clippy --workspace --all-targets --all-features --locked -- -D warnings
@@ -51,6 +51,10 @@ pub enum EngineStatus {
    Cancelled {
        hash_count: u64,
    },
+    /// GPU device lost or unresponsive. Worker should exit permanently.
+    DeviceLost {
+        hash_count: u64,
+    },
 }

 /// Cancellation checker passed to search_range.
@@ -27,17 +27,20 @@ qp-poseidon-constants = "1.1"
 qp-plonky2-field = { version = "1.1.1" }
 plonky2 = { package = "qp-plonky2", version = "1.1.3" }

-wgpu = { version = "27.0.1" }                              # GPU compute library
-futures = "0.3"                                            # For async executor
-bytemuck = "1.16"                                          # For buffer mapping
+wgpu = { version = "27.0.1" }                                              # GPU compute library
+tokio = { workspace = true, features = ["rt", "rt-multi-thread", "time"] } # Async runtime with timeout support
+futures = "0.3"                                                            # For async oneshot channels in test modules
+bytemuck = "1.16"                                                          # For buffer mapping
 rand = { workspace = true, features = ["std", "std_rng"] }
-rand_chacha = "0.9"                                        # For deterministic random test generation
+rand_chacha = "0.9"                                                        # For deterministic random test generation
+regex = "1"                                                                # For GPU name pattern matching

 [dev-dependencies]
 hex = { workspace = true }
 criterion = "0.5"
 rand = { workspace = true, features = ["std", "std_rng"] }
 env_logger.workspace = true
+tokio = { workspace = true, features = ["macros", "rt", "rt-multi-thread", "time"] } # Add macros for #[tokio::main]

 [[bench]]
 name = "gpu_engine_bench"
@@ -6,9 +6,17 @@ use primitive_types::U512;
 use rand::RngCore;
 use std::sync::atomic::AtomicBool;

+/// Drop thread-local wgpu buffers before `GpuEngine` is dropped. Criterion
+/// creates a fresh engine per group; without this, TLS buffers outlive the
+/// device and the next group panics (`Buffer[…] does not exist`).
+fn teardown_gpu(engine: GpuEngine) {
+    GpuEngine::clear_worker_resources();
+    drop(engine);
+}
+
 fn bench_cpu_vs_gpu_small(c: &mut Criterion) {
    let cpu_engine = FastCpuEngine::new(10_000);
-    let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
+    let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
    let cancel_flag = AtomicBool::new(false);
    let cancel_check = AtomicBoolCancelCheck(&cancel_flag);

@@ -55,11 +63,12 @@ fn bench_cpu_vs_gpu_small(c: &mut Criterion) {
    });

    group.finish();
+    teardown_gpu(gpu_engine);
 }

 fn bench_cpu_vs_gpu_medium(c: &mut Criterion) {
    let cpu_engine = FastCpuEngine::new(10_000);
-    let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
+    let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
    let cancel_flag = AtomicBool::new(false);
    let cancel_check = AtomicBoolCancelCheck(&cancel_flag);

@@ -106,11 +115,12 @@ fn bench_cpu_vs_gpu_medium(c: &mut Criterion) {
    });

    group.finish();
+    teardown_gpu(gpu_engine);
 }

 fn bench_cpu_vs_gpu_large(c: &mut Criterion) {
    let cpu_engine = FastCpuEngine::new(10_000);
-    let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
+    let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
    let cancel_flag = AtomicBool::new(false);
    let cancel_check = AtomicBoolCancelCheck(&cancel_flag);

@@ -157,11 +167,12 @@ fn bench_cpu_vs_gpu_large(c: &mut Criterion) {
    });

    group.finish();
+    teardown_gpu(gpu_engine);
 }

 fn bench_solution_finding(c: &mut Criterion) {
    let cpu_engine = FastCpuEngine::new(10_000);
-    let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
+    let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
    let cancel_flag = AtomicBool::new(false);
    let cancel_check = AtomicBoolCancelCheck(&cancel_flag);

@@ -208,11 +219,12 @@ fn bench_solution_finding(c: &mut Criterion) {
    });

    group.finish();
+    teardown_gpu(gpu_engine);
 }

 fn bench_throughput_per_second(c: &mut Criterion) {
    let cpu_engine = FastCpuEngine::new(10_000);
-    let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
+    let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
    let cancel_flag = AtomicBool::new(false);
    let cancel_check = AtomicBoolCancelCheck(&cancel_flag);

@@ -259,10 +271,11 @@ fn bench_throughput_per_second(c: &mut Criterion) {
    });

    group.finish();
+    teardown_gpu(gpu_engine);
 }

 fn bench_gpu_batch_efficiency(c: &mut Criterion) {
-    let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
+    let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
    let cancel_flag = AtomicBool::new(false);
    let cancel_check = AtomicBoolCancelCheck(&cancel_flag);

@@ -335,6 +348,7 @@ fn bench_gpu_batch_efficiency(c: &mut Criterion) {
    });

    group.finish();
+    teardown_gpu(gpu_engine);
 }

 criterion_group!(
@@ -0,0 +1,63 @@
+use engine_cpu::{AtomicBoolCancelCheck, EngineStatus, MinerEngine, Range};
+use engine_gpu::GpuEngine;
+use primitive_types::U512;
+use rand::RngCore;
+use std::sync::atomic::AtomicBool;
+
+fn main() {
+    env_logger::init();
+    let jobs: usize = std::env::args()
+        .nth(1)
+        .map(|s| s.parse().expect("job count"))
+        .unwrap_or(25);
+
+    let engine = GpuEngine::try_new(1_000_000, 0, false).expect("GPU init failed");
+    let cancel_flag = AtomicBool::new(false);
+    let cancel = AtomicBoolCancelCheck(&cancel_flag);
+    let mut rng = rand::rng();
+
+    let mut found = 0usize;
+    for job in 0..jobs {
+        let mut header = [0u8; 32];
+        rng.fill_bytes(&mut header);
+        let difficulty = U512::from(100_000u64);
+        let ctx = engine.prepare_context(header, difficulty);
+
+        let start = if job == 0 {
+            // Cross a 2^256 boundary: the high nonce half changes mid-range,
+            // exercising the midstate batch clamp.
+            (U512::from(3u64) << 256) - U512::from(1_000u64)
+        } else {
+            let mut start_bytes = [0u8; 64];
+            rng.fill_bytes(&mut start_bytes);
+            // Keep clear of the very top so range arithmetic cannot wrap
+            start_bytes[0] = 0;
+            U512::from_big_endian(&start_bytes)
+        };
+        let range = Range {
+            start,
+            end: start + U512::from(10_000_000u64),
+        };
+
+        match engine.search_range(&ctx, range.clone(), &cancel) {
+            EngineStatus::Found { candidate, .. } => {
+                let cpu_hash = pow_core::hash_from_nonce(&ctx, candidate.nonce);
+                assert_eq!(
+                    cpu_hash, candidate.hash,
+                    "job {job}: GPU hash != CPU hash for nonce {}",
+                    candidate.nonce
+                );
+                assert!(cpu_hash < ctx.target, "job {job}: hash not below target");
+                assert!(
+                    candidate.nonce >= range.start && candidate.nonce <= range.end,
+                    "job {job}: nonce outside range"
+                );
+                found += 1;
+            }
+            EngineStatus::Exhausted { .. } => {}
+            other => panic!("job {job}: unexpected status {other:?}"),
+        }
+    }
+    assert!(found > 0, "no solutions found across {jobs} jobs");
+    println!("PARITY OK: {found}/{jobs} jobs found solutions, all verified against CPU");
+}
@@ -0,0 +1,30 @@
+fn main() {
+    let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
+        backends: wgpu::Backends::PRIMARY,
+        ..Default::default()
+    });
+    for adapter in instance.enumerate_adapters(wgpu::Backends::PRIMARY) {
+        let info = adapter.get_info();
+        let features = adapter.features();
+        println!("{} ({:?}, {:?})", info.name, info.device_type, info.backend);
+        println!(
+            "  SHADER_INT64: {}",
+            features.contains(wgpu::Features::SHADER_INT64)
+        );
+        println!(
+            "  SUBGROUP: {}",
+            features.contains(wgpu::Features::SUBGROUP)
+        );
+        println!(
+            "  TIMESTAMP_QUERY: {}",
+            features.contains(wgpu::Features::TIMESTAMP_QUERY)
+        );
+        let limits = adapter.limits();
+        println!(
+            "  max_workgroup_size_x: {}, max_invocations: {}, max_workgroups_per_dim: {}",
+            limits.max_compute_workgroup_size_x,
+            limits.max_compute_invocations_per_workgroup,
+            limits.max_compute_workgroups_per_dimension
+        );
+    }
+}
@@ -0,0 +1,52 @@
+use engine_cpu::{AtomicBoolCancelCheck, MinerEngine, Range};
+use engine_gpu::GpuEngine;
+use primitive_types::U512;
+use std::sync::atomic::AtomicBool;
+
+fn main() {
+    env_logger::init();
+    let args: Vec<String> = std::env::args().collect();
+    let total: u64 = args
+        .get(1)
+        .map(|s| s.parse().expect("total nonces"))
+        .unwrap_or(4_000_000);
+    let batch: u32 = args
+        .get(2)
+        .map(|s| s.parse().expect("batch size"))
+        .unwrap_or(1_000_000);
+
+    let engine = GpuEngine::try_new(batch, 0, false).expect("GPU init failed");
+    let cancel_flag = AtomicBool::new(false);
+    let cancel = AtomicBoolCancelCheck(&cancel_flag);
+
+    let header = [42u8; 32];
+    let ctx = engine.prepare_context(header, U512::from(u64::MAX));
+
+    let range = Range {
+        start: U512::from(1u64) << 200,
+        end: (U512::from(1u64) << 200) + U512::from(total - 1),
+    };
+
+    // Warmup
+    let warm = Range {
+        start: range.start,
+        end: range.start + U512::from(200_000u64),
+    };
+    engine.search_range(&ctx, warm, &cancel);
+
+    let start = std::time::Instant::now();
+    let status = engine.search_range(&ctx, range, &cancel);
+    let elapsed = start.elapsed().as_secs_f64();
+
+    let hashes = match status {
+        engine_cpu::EngineStatus::Exhausted { hash_count } => hash_count,
+        engine_cpu::EngineStatus::Found { hash_count, .. } => hash_count,
+        other => panic!("unexpected status: {other:?}"),
+    };
+    println!(
+        "{} hashes in {:.3}s = {:.3} MH/s",
+        hashes,
+        elapsed,
+        hashes as f64 / elapsed / 1e6
+    );
+}
@@ -0,0 +1,228 @@
+use pow_core::{hash_from_nonce, mining_midstate, JobContext};
+use primitive_types::U512;
+use std::time::Instant;
+
+fn u512_to_u32s_le(v: U512) -> [u32; 16] {
+    let bytes = v.to_little_endian();
+    let mut out = [0u32; 16];
+    for i in 0..16 {
+        out[i] = u32::from_le_bytes(bytes[i * 4..(i + 1) * 4].try_into().unwrap());
+    }
+    out
+}
+
+struct Runner {
+    device: wgpu::Device,
+    queue: wgpu::Queue,
+    pipeline: wgpu::ComputePipeline,
+    results: wgpu::Buffer,
+    midstate: wgpu::Buffer,
+    start_nonce: wgpu::Buffer,
+    target: wgpu::Buffer,
+    cfg: wgpu::Buffer,
+    staging: wgpu::Buffer,
+}
+
+impl Runner {
+    async fn new(trusted: bool) -> Self {
+        let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
+            backends: wgpu::Backends::PRIMARY,
+            ..Default::default()
+        });
+        let adapter = instance
+            .request_adapter(&wgpu::RequestAdapterOptions::default())
+            .await
+            .expect("no adapter");
+        assert!(
+            adapter.features().contains(wgpu::Features::SHADER_INT64),
+            "SHADER_INT64 required"
+        );
+        let (device, queue) = adapter
+            .request_device(&wgpu::DeviceDescriptor {
+                label: None,
+                required_features: wgpu::Features::SHADER_INT64,
+                ..Default::default()
+            })
+            .await
+            .unwrap();
+
+        let desc = wgpu::ShaderModuleDescriptor {
+            label: None,
+            source: wgpu::ShaderSource::Wgsl(include_str!("../src/mining_u64.wgsl").into()),
+        };
+        let shader = if trusted {
+            unsafe {
+                device.create_shader_module_trusted(desc, wgpu::ShaderRuntimeChecks::unchecked())
+            }
+        } else {
+            device.create_shader_module(desc)
+        };
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: None,
+            layout: None,
+            module: &shader,
+            entry_point: Some("mining_main"),
+            compilation_options: Default::default(),
+            cache: None,
+        });
+        let mk = |size: u64, usage| {
+            device.create_buffer(&wgpu::BufferDescriptor {
+                label: None,
+                size,
+                usage,
+                mapped_at_creation: false,
+            })
+        };
+        use wgpu::BufferUsages as U;
+        Runner {
+            pipeline,
+            results: mk(132, U::STORAGE | U::COPY_SRC | U::COPY_DST),
+            midstate: mk(96, U::STORAGE | U::COPY_DST),
+            start_nonce: mk(64, U::STORAGE | U::COPY_DST),
+            target: mk(64, U::STORAGE | U::COPY_DST),
+            cfg: mk(12, U::STORAGE | U::COPY_DST),
+            staging: mk(132, U::MAP_READ | U::COPY_DST),
+            device,
+            queue,
+        }
+    }
+
+    fn run_batch(&self, header: [u8; 32], start: U512, batch: u32, target: U512) -> f64 {
+        let nonce_be = start.to_big_endian();
+        let mid = mining_midstate(header, nonce_be[..32].try_into().unwrap());
+        let mut mid_u32 = [0u32; 24];
+        for (i, f) in mid.iter().enumerate() {
+            mid_u32[2 * i] = *f as u32;
+            mid_u32[2 * i + 1] = (*f >> 32) as u32;
+        }
+        self.queue
+            .write_buffer(&self.midstate, 0, bytemuck::cast_slice(&mid_u32));
+        self.queue.write_buffer(
+            &self.start_nonce,
+            0,
+            bytemuck::cast_slice(&u512_to_u32s_le(start)),
+        );
+        self.queue.write_buffer(
+            &self.target,
+            0,
+            bytemuck::cast_slice(&u512_to_u32s_le(target)),
+        );
+        self.queue
+            .write_buffer(&self.cfg, 0, bytemuck::cast_slice(&[batch, 1u32, batch]));
+        self.queue.write_buffer(&self.results, 0, &[0u8; 132]);
+
+        let layout = self.pipeline.get_bind_group_layout(0);
+        let entries: Vec<wgpu::BindGroupEntry<'_>> = [
+            (0, &self.results),
+            (1, &self.midstate),
+            (2, &self.start_nonce),
+            (3, &self.target),
+            (4, &self.cfg),
+        ]
+        .iter()
+        .map(|(i, b)| wgpu::BindGroupEntry {
+            binding: *i,
+            resource: b.as_entire_binding(),
+        })
+        .collect();
+        let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &layout,
+            entries: &entries,
+        });
+
+        let t0 = Instant::now();
+        let mut encoder = self
+            .device
+            .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+        {
+            let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: None,
+                timestamp_writes: None,
+            });
+            cpass.set_pipeline(&self.pipeline);
+            cpass.set_bind_group(0, &bind_group, &[]);
+            cpass.dispatch_workgroups(batch.div_ceil(256), 1, 1);
+        }
+        encoder.copy_buffer_to_buffer(&self.results, 0, &self.staging, 0, 132);
+        self.queue.submit(Some(encoder.finish()));
+
+        let slice = self.staging.slice(..);
+        let (tx, rx) = std::sync::mpsc::channel();
+        slice.map_async(wgpu::MapMode::Read, move |r| tx.send(r).unwrap());
+        loop {
+            let _ = self.device.poll(wgpu::PollType::Wait {
+                submission_index: None,
+                timeout: None,
+            });
+            if rx.try_recv().is_ok() {
+                break;
+            }
+        }
+        let elapsed = t0.elapsed().as_secs_f64();
+        self.staging.unmap();
+        elapsed
+    }
+
+    fn read_results(&self) -> [u32; 33] {
+        let slice = self.staging.slice(..);
+        let (tx, rx) = std::sync::mpsc::channel();
+        slice.map_async(wgpu::MapMode::Read, move |r| tx.send(r).unwrap());
+        let _ = self.device.poll(wgpu::PollType::Wait {
+            submission_index: None,
+            timeout: None,
+        });
+        rx.recv().unwrap().unwrap();
+        let data = slice.get_mapped_range();
+        let mut out = [0u32; 33];
+        out.copy_from_slice(bytemuck::cast_slice(&data));
+        drop(data);
+        self.staging.unmap();
+        out
+    }
+}
+
+fn main() {
+    let batch: u32 = std::env::args()
+        .nth(1)
+        .map(|s| s.parse().unwrap())
+        .unwrap_or(8_000_000);
+    let iters = 4u32;
+    let rt = tokio::runtime::Runtime::new().unwrap();
+
+    for trusted in [false, true] {
+        let label = if trusted {
+            "trusted (no checks)"
+        } else {
+            "checked (default)"
+        };
+        let runner = rt.block_on(Runner::new(trusted));
+
+        let header = [9u8; 32];
+        let ctx = JobContext::new(header, U512::one());
+        let start = (U512::from(7u64) << 300) | U512::from(123456789u64);
+        runner.run_batch(header, start, 256, ctx.target);
+        let r = runner.read_results();
+        assert_eq!(r[0], 1, "{label}: no solution with target=MAX");
+        let nonce = U512::from_little_endian(bytemuck::cast_slice(&r[1..17]));
+        let hash = U512::from_little_endian(bytemuck::cast_slice(&r[17..33]));
+        assert_eq!(hash, hash_from_nonce(&ctx, nonce), "{label}: hash mismatch");
+        println!("{label}: correctness OK");
+
+        let header = [42u8; 32];
+        let start = U512::from(1u64) << 200;
+        let target = U512::one();
+        runner.run_batch(header, start, batch, target);
+        let mut total = 0.0;
+        for i in 0..iters {
+            total += runner.run_batch(
+                header,
+                start + U512::from((i as u64 + 1) * batch as u64),
+                batch,
+                target,
+            );
+        }
+        let mhs = (batch as f64 * iters as f64) / total / 1e6;
+        println!("{label}: {batch} nonces x{iters} in {total:.3}s = {mhs:.3} MH/s");
+    }
+}
@@ -26,7 +26,7 @@ fn main() {

    // 3. Verify with GPU engine
    log::info!("Initializing GPU engine...");
-    let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
+    let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");

    // Search a small range around the valid nonce
    let gpu_range = Range {
@@ -56,6 +56,9 @@ fn main() {
        EngineStatus::Cancelled { .. } => {
            log::error!("FAILURE: GPU search cancelled!");
        }
+        EngineStatus::DeviceLost { .. } => {
+            log::error!("FAILURE: GPU device lost!");
+        }
        EngineStatus::Running { .. } => {
            log::error!("FAILURE: GPU returned Running status!");
        }
@@ -7,6 +7,7 @@ use wgpu::util::DeviceExt;
 pub async fn test_end_to_end_mining(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    println!("Running End-to-End Mining Test...");

@@ -28,15 +29,17 @@ pub async fn test_end_to_end_mining(

    // 4. Run GPU Mining for this specific nonce

-    // Header Buffer
-    let mut header_u32s = [0u32; 8];
-    for (i, item) in header_u32s.iter_mut().enumerate() {
-        let chunk = &ctx.header[i * 4..(i + 1) * 4];
-        *item = u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]);
+    // Midstate Buffer (header + high nonce half absorbed on CPU)
+    let nonce_be = nonce_val.to_big_endian();
+    let midstate = pow_core::mining_midstate(ctx.header, nonce_be[..32].try_into().unwrap());
+    let mut midstate_u32s = [0u32; 24];
+    for (i, felt) in midstate.iter().enumerate() {
+        midstate_u32s[2 * i] = *felt as u32;
+        midstate_u32s[2 * i + 1] = (*felt >> 32) as u32;
    }
-    let header_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
-        label: Some("Header Buffer"),
-        contents: bytemuck::cast_slice(&header_u32s),
+    let midstate_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+        label: Some("Midstate Buffer"),
+        contents: bytemuck::cast_slice(&midstate_u32s),
        usage: wgpu::BufferUsages::STORAGE,
    });

@@ -82,11 +85,11 @@ pub async fn test_end_to_end_mining(
    let zeros = vec![0u8; results_size];
    queue.write_buffer(&results_buffer, 0, &zeros);

-    // Dispatch config buffer: [total_threads, nonces_per_thread, total_nonces, cancel_check_interval]
-    let dispatch_config_data: [u32; 4] = [256, 1, 256, 10000];
+    // Dispatch config buffer: [total_threads, nonces_per_thread, total_nonces]
+    let dispatch_config_data: [u32; 3] = [256, 1, 256];
    let dispatch_config_buffer = device.create_buffer(&wgpu::BufferDescriptor {
        label: Some("Dispatch Config Buffer"),
-        size: 16,
+        size: 12,
        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
        mapped_at_creation: false,
    });
@@ -96,25 +99,9 @@ pub async fn test_end_to_end_mining(
        bytemuck::cast_slice(&dispatch_config_data),
    );

-    // Cancel flag buffer: 0 = running, 1 = cancel requested
-    let cancel_flag_data: [u32; 1] = [0];
-    let cancel_flag_buffer = device.create_buffer(&wgpu::BufferDescriptor {
-        label: Some("Cancel Flag Buffer"),
-        size: 4,
-        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
-        mapped_at_creation: false,
-    });
-    queue.write_buffer(
-        &cancel_flag_buffer,
-        0,
-        bytemuck::cast_slice(&cancel_flag_data),
-    );
-
-    // Load Shader
-    let shader_source = include_str!("mining.wgsl");
    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
        label: Some("Mining Shader"),
-        source: wgpu::ShaderSource::Wgsl(shader_source.into()),
+        source: wgpu::ShaderSource::Wgsl(shader_src.into()),
    });

    let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
@@ -137,7 +124,7 @@ pub async fn test_end_to_end_mining(
            },
            wgpu::BindGroupEntry {
                binding: 1,
-                resource: header_buffer.as_entire_binding(),
+                resource: midstate_buffer.as_entire_binding(),
            },
            wgpu::BindGroupEntry {
                binding: 2,
@@ -151,10 +138,6 @@ pub async fn test_end_to_end_mining(
                binding: 4,
                resource: dispatch_config_buffer.as_entire_binding(),
            },
-            wgpu::BindGroupEntry {
-                binding: 5,
-                resource: cancel_flag_buffer.as_entire_binding(),
-            },
        ],
    });

@@ -0,0 +1,748 @@
+//! GPU tier detection using regex patterns
+//!
+//! This module provides a table-driven approach to GPU detection,
+//! using regex patterns with word boundaries to avoid substring matching issues.
+
+use regex::Regex;
+use std::sync::LazyLock;
+
+/// GPU tier configuration
+struct GpuTier {
+    /// Regex pattern to match GPU name (case-insensitive)
+    pattern: &'static str,
+    /// Human-readable tier name
+    name: &'static str,
+    /// Divisor for max_workgroups (higher = more conservative)
+    workgroup_divisor: u32,
+    /// Minimum workgroups to use
+    min_workgroups: u32,
+}
+
+/// Compiled GPU tier with regex
+struct CompiledGpuTier {
+    regex: Regex,
+    name: &'static str,
+    workgroup_divisor: u32,
+    min_workgroups: u32,
+}
+
+impl CompiledGpuTier {
+    fn from_tier(tier: &GpuTier) -> Self {
+        Self {
+            regex: Regex::new(&format!("(?i){}", tier.pattern)).expect("Invalid GPU tier regex"),
+            name: tier.name,
+            workgroup_divisor: tier.workgroup_divisor,
+            min_workgroups: tier.min_workgroups,
+        }
+    }
+}
+
+// GPU tiers are checked in order - first match wins
+// Use word boundaries (\b) to avoid substring issues like "550" matching "5500"
+
+const NVIDIA_TIERS: &[GpuTier] = &[
+    // Blackwell (RTX 50 series)
+    GpuTier {
+        pattern: r"\b50[89]0\b",
+        name: "NVIDIA RTX 50 Flagship (Blackwell)",
+        workgroup_divisor: 6,
+        min_workgroups: 5120,
+    },
+    GpuTier {
+        pattern: r"\b50[67]0\b|rtx 50",
+        name: "NVIDIA RTX 50 (Blackwell)",
+        workgroup_divisor: 7,
+        min_workgroups: 4608,
+    },
+    // Ada Lovelace (RTX 40 series)
+    GpuTier {
+        pattern: r"\b40[89]0\b",
+        name: "NVIDIA RTX 40 Flagship (Ada)",
+        workgroup_divisor: 8,
+        min_workgroups: 4096,
+    },
+    GpuTier {
+        pattern: r"\b40[67]0\b|rtx 40",
+        name: "NVIDIA RTX 40 (Ada)",
+        workgroup_divisor: 10,
+        min_workgroups: 3072,
+    },
+    // Ampere/Turing (RTX 30/20 series)
+    GpuTier {
+        pattern: r"\b30[5-9]0\b|\b20[6-8]0\b|rtx 30|rtx 20",
+        name: "NVIDIA RTX 30/20 (Ampere/Turing)",
+        workgroup_divisor: 12,
+        min_workgroups: 2048,
+    },
+    // Turing/Pascal (GTX 16/10 series)
+    GpuTier {
+        pattern: r"\b16[56]0\b|\b10[3-8]0\b|gtx 16|gtx 10",
+        name: "NVIDIA GTX 16/10 (Turing/Pascal)",
+        workgroup_divisor: 16,
+        min_workgroups: 1024,
+    },
+    // Maxwell (GTX 900 series)
+    GpuTier {
+        pattern: r"\b9[5-8]0\b|gtx 9",
+        name: "NVIDIA GTX 900 (Maxwell)",
+        workgroup_divisor: 18,
+        min_workgroups: 768,
+    },
+    // Kepler/Maxwell (GTX 700 series)
+    GpuTier {
+        pattern: r"\b7[5-8]0\b|gtx 7",
+        name: "NVIDIA GTX 700 (Kepler/Maxwell)",
+        workgroup_divisor: 20,
+        min_workgroups: 512,
+    },
+    // Legacy GTX
+    GpuTier {
+        pattern: r"gtx [456]",
+        name: "NVIDIA GTX Legacy (Fermi/Kepler)",
+        workgroup_divisor: 24,
+        min_workgroups: 384,
+    },
+    GpuTier {
+        pattern: r"\bgtx\b",
+        name: "NVIDIA GTX (Unknown)",
+        workgroup_divisor: 20,
+        min_workgroups: 512,
+    },
+    // Mobile
+    GpuTier {
+        pattern: r"\bmx[1-5]\d0|geforce mx",
+        name: "NVIDIA MX (Mobile)",
+        workgroup_divisor: 24,
+        min_workgroups: 384,
+    },
+    GpuTier {
+        pattern: r"geforce gt\b|\bgt \d{3}\b",
+        name: "NVIDIA GT (Entry-Level)",
+        workgroup_divisor: 28,
+        min_workgroups: 256,
+    },
+    // Professional
+    GpuTier {
+        pattern: r"quadro|rtx a\d|tesla|\ba100\b|\bh100\b|\bl4\b",
+        name: "NVIDIA Quadro/Professional",
+        workgroup_divisor: 10,
+        min_workgroups: 2560,
+    },
+];
+
+const AMD_TIERS: &[GpuTier] = &[
+    // RDNA 4 (RX 9000 series)
+    GpuTier {
+        pattern: r"rx 9|\b90[78]0\b",
+        name: "AMD RX 9000 (RDNA 4)",
+        workgroup_divisor: 8,
+        min_workgroups: 4096,
+    },
+    // RDNA 3 Discrete
+    GpuTier {
+        pattern: r"\b7900\b",
+        name: "AMD RX 7900 (RDNA 3 Flagship)",
+        workgroup_divisor: 9,
+        min_workgroups: 3584,
+    },
+    GpuTier {
+        pattern: r"rx 7|\b7[6-8]00\b",
+        name: "AMD RX 7000 (RDNA 3)",
+        workgroup_divisor: 10,
+        min_workgroups: 3072,
+    },
+    // RDNA 3 APUs - check before discrete to avoid substring match
+    GpuTier {
+        pattern: r"\b780m\b|radeon 780m",
+        name: "AMD Radeon 780M (RDNA 3 APU)",
+        workgroup_divisor: 12,
+        min_workgroups: 2048,
+    },
+    GpuTier {
+        pattern: r"\b7[46]0m\b|radeon 7[46]0m",
+        name: "AMD Radeon 7x0M (RDNA 3 APU)",
+        workgroup_divisor: 16,
+        min_workgroups: 1024,
+    },
+    // RDNA 2 Discrete
+    GpuTier {
+        pattern: r"\b6[89][05]0\b",
+        name: "AMD RX 6900/6800 (RDNA 2 Flagship)",
+        workgroup_divisor: 12,
+        min_workgroups: 2560,
+    },
+    GpuTier {
+        pattern: r"\b6[67][05]0\b",
+        name: "AMD RX 6700/6600 (RDNA 2)",
+        workgroup_divisor: 14,
+        min_workgroups: 2048,
+    },
+    GpuTier {
+        pattern: r"\b6[45]00\b",
+        name: "AMD RX 6500/6400 (RDNA 2 Entry)",
+        workgroup_divisor: 22,
+        min_workgroups: 512,
+    },
+    GpuTier {
+        pattern: r"rx 6\d{3}",
+        name: "AMD RX 6000 (RDNA 2)",
+        workgroup_divisor: 14,
+        min_workgroups: 2048,
+    },
+    // RDNA 2 APUs - check before discrete
+    GpuTier {
+        pattern: r"\b680m\b|radeon 680m",
+        name: "AMD Radeon 680M (RDNA 2 APU)",
+        workgroup_divisor: 16,
+        min_workgroups: 1536,
+    },
+    GpuTier {
+        pattern: r"\b6[16]0m\b|radeon 6[16]0m",
+        name: "AMD Radeon 6x0M (RDNA 2 APU)",
+        workgroup_divisor: 22,
+        min_workgroups: 768,
+    },
+    // RDNA 1 (RX 5000 series) - 4-digit patterns to avoid matching Polaris 3-digit
+    GpuTier {
+        pattern: r"\b5700\b",
+        name: "AMD RX 5700 (RDNA 1)",
+        workgroup_divisor: 16,
+        min_workgroups: 1536,
+    },
+    GpuTier {
+        pattern: r"\b5[56]00\b|rx 5\d{3}",
+        name: "AMD RX 5000 (RDNA 1)",
+        workgroup_divisor: 18,
+        min_workgroups: 1024,
+    },
+    // Polaris (RX 400/500 series) - 3-digit models with boundaries
+    GpuTier {
+        pattern: r"rx [45]\d0\b|\b[45][6-9]0\b|\b590\b|rx 5.0x",
+        name: "AMD RX 500/400 (Polaris)",
+        workgroup_divisor: 20,
+        min_workgroups: 768,
+    },
+    // Vega
+    GpuTier {
+        pattern: r"radeon vii\b",
+        name: "AMD Radeon VII (Vega 20)",
+        workgroup_divisor: 12,
+        min_workgroups: 2048,
+    },
+    GpuTier {
+        pattern: r"vega\s*64",
+        name: "AMD Vega 64 (Discrete)",
+        workgroup_divisor: 14,
+        min_workgroups: 1536,
+    },
+    GpuTier {
+        pattern: r"vega\s*56",
+        name: "AMD Vega 56 (Discrete)",
+        workgroup_divisor: 16,
+        min_workgroups: 1280,
+    },
+    GpuTier {
+        pattern: r"vega",
+        name: "AMD Vega (APU)",
+        workgroup_divisor: 28,
+        min_workgroups: 384,
+    },
+    // GCN
+    GpuTier {
+        pattern: r"fury|nano",
+        name: "AMD R9 Fury/Nano (Fiji)",
+        workgroup_divisor: 16,
+        min_workgroups: 1280,
+    },
+    GpuTier {
+        pattern: r"r9.*(3[89]0|2[89]0)|\b[23][89]0x?\b",
+        name: "AMD R9 (GCN)",
+        workgroup_divisor: 20,
+        min_workgroups: 768,
+    },
+    GpuTier {
+        pattern: r"r7.*(3[67]0|2[67]0)|\b[23][67]0x?\b",
+        name: "AMD R7 (GCN)",
+        workgroup_divisor: 22,
+        min_workgroups: 512,
+    },
+    // Professional
+    GpuTier {
+        pattern: r"radeon pro|instinct mi|mi[123]\d0|firepro|w[567]\d{3}",
+        name: "AMD Radeon Pro/Instinct",
+        workgroup_divisor: 10,
+        min_workgroups: 2560,
+    },
+    // OEM/APU fallbacks
+    GpuTier {
+        pattern: r"radeon\s*\(tm\)\s*[67]\d{2}\b|radeon [67]\d{2}\b",
+        name: "AMD Radeon OEM (Polaris Rebrand)",
+        workgroup_divisor: 24,
+        min_workgroups: 512,
+    },
+    GpuTier {
+        pattern: r"radeon\s*(\(tm\)\s*)?graphics",
+        name: "AMD Radeon Graphics (APU)",
+        workgroup_divisor: 26,
+        min_workgroups: 384,
+    },
+];
+
+const INTEL_TIERS: &[GpuTier] = &[
+    // Battlemage (Arc B-Series)
+    GpuTier {
+        pattern: r"arc b|\bb5[78]0\b",
+        name: "Intel Arc B-Series (Battlemage)",
+        workgroup_divisor: 10,
+        min_workgroups: 2560,
+    },
+    // Alchemist Mobile - check BEFORE desktop (a770m contains a770)
+    GpuTier {
+        pattern: r"\ba7[37]0m\b",
+        name: "Intel Arc A7 Mobile (Alchemist)",
+        workgroup_divisor: 14,
+        min_workgroups: 1536,
+    },
+    GpuTier {
+        pattern: r"\ba5[57]0m\b",
+        name: "Intel Arc A5 Mobile (Alchemist)",
+        workgroup_divisor: 16,
+        min_workgroups: 1024,
+    },
+    GpuTier {
+        pattern: r"\ba3[57]0m\b",
+        name: "Intel Arc A3 Mobile (Alchemist)",
+        workgroup_divisor: 20,
+        min_workgroups: 512,
+    },
+    // Alchemist Desktop
+    GpuTier {
+        pattern: r"\ba7[57]0\b",
+        name: "Intel Arc A7 (Alchemist)",
+        workgroup_divisor: 12,
+        min_workgroups: 2048,
+    },
+    GpuTier {
+        pattern: r"\ba580\b",
+        name: "Intel Arc A5 (Alchemist)",
+        workgroup_divisor: 14,
+        min_workgroups: 1536,
+    },
+    GpuTier {
+        pattern: r"\ba3[18]0\b",
+        name: "Intel Arc A3 (Alchemist)",
+        workgroup_divisor: 18,
+        min_workgroups: 768,
+    },
+    GpuTier {
+        pattern: r"arc a|\barc\b",
+        name: "Intel Arc (Unknown)",
+        workgroup_divisor: 16,
+        min_workgroups: 1024,
+    },
+    // Integrated
+    GpuTier {
+        pattern: r"iris xe max",
+        name: "Intel Iris Xe Max (Discrete)",
+        workgroup_divisor: 20,
+        min_workgroups: 512,
+    },
+    GpuTier {
+        pattern: r"iris xe",
+        name: "Intel Iris Xe (Integrated)",
+        workgroup_divisor: 24,
+        min_workgroups: 384,
+    },
+    GpuTier {
+        pattern: r"iris pro",
+        name: "Intel Iris Pro (Integrated)",
+        workgroup_divisor: 26,
+        min_workgroups: 256,
+    },
+    GpuTier {
+        pattern: r"iris plus|iris\b",
+        name: "Intel Iris Plus (Integrated)",
+        workgroup_divisor: 26,
+        min_workgroups: 320,
+    },
+    GpuTier {
+        pattern: r"uhd.*(7\d{2}|graphics 7)",
+        name: "Intel UHD 700 (Integrated)",
+        workgroup_divisor: 26,
+        min_workgroups: 320,
+    },
+    GpuTier {
+        pattern: r"uhd.*(6\d{2}|graphics 6)",
+        name: "Intel UHD 600 (Integrated)",
+        workgroup_divisor: 28,
+        min_workgroups: 256,
+    },
+    GpuTier {
+        pattern: r"uhd",
+        name: "Intel UHD Graphics (Integrated)",
+        workgroup_divisor: 28,
+        min_workgroups: 256,
+    },
+    GpuTier {
+        pattern: r"hd graphics|hd [456]\d{2}",
+        name: "Intel HD Graphics (Integrated)",
+        workgroup_divisor: 30,
+        min_workgroups: 192,
+    },
+];
+
+const QUALCOMM_TIERS: &[GpuTier] = &[
+    // Snapdragon X (Adreno X1)
+    GpuTier {
+        pattern: r"x elite|x plus|adreno x|x1-[89]",
+        name: "Qualcomm Adreno X1 (Snapdragon X)",
+        workgroup_divisor: 14,
+        min_workgroups: 1536,
+    },
+    // Adreno 700 series (730, 740, 750, etc.)
+    GpuTier {
+        pattern: r"adreno\s*7[0-9]{2}\b|\b7[345]0\b",
+        name: "Qualcomm Adreno 700 Series",
+        workgroup_divisor: 16,
+        min_workgroups: 1024,
+    },
+    // Adreno 600 series (610, 612, 615, 618, 619, 620, 630, 640, 650, 660, etc.)
+    GpuTier {
+        pattern: r"adreno\s*6[0-9]{2}\b|\b6[1-6][0-9]\b",
+        name: "Qualcomm Adreno 600 Series",
+        workgroup_divisor: 20,
+        min_workgroups: 512,
+    },
+    // Adreno 500 series (505, 506, 508, 509, 510, 512, 530, 540)
+    GpuTier {
+        pattern: r"adreno\s*5[0-4][0-9]\b|\b5[0-4][0-9]\b",
+        name: "Qualcomm Adreno 500 Series",
+        workgroup_divisor: 24,
+        min_workgroups: 384,
+    },
+];
+
+const APPLE_TIERS: &[GpuTier] = &[
+    // M5 series
+    GpuTier {
+        pattern: r"m5 ultra",
+        name: "Apple M5 Ultra",
+        workgroup_divisor: 4,
+        min_workgroups: 1600,
+    },
+    GpuTier {
+        pattern: r"m5 max",
+        name: "Apple M5 Max",
+        workgroup_divisor: 4,
+        min_workgroups: 800,
+    },
+    GpuTier {
+        pattern: r"m5 pro",
+        name: "Apple M5 Pro",
+        workgroup_divisor: 4,
+        min_workgroups: 400,
+    },
+    GpuTier {
+        pattern: r"\bm5\b",
+        name: "Apple M5",
+        workgroup_divisor: 4,
+        min_workgroups: 200,
+    },
+    // M4 series
+    GpuTier {
+        pattern: r"m4 ultra",
+        name: "Apple M4 Ultra",
+        workgroup_divisor: 4,
+        min_workgroups: 1600,
+    },
+    GpuTier {
+        pattern: r"m4 max",
+        name: "Apple M4 Max",
+        workgroup_divisor: 4,
+        min_workgroups: 800,
+    },
+    GpuTier {
+        pattern: r"m4 pro",
+        name: "Apple M4 Pro",
+        workgroup_divisor: 4,
+        min_workgroups: 400,
+    },
+    GpuTier {
+        pattern: r"\bm4\b",
+        name: "Apple M4",
+        workgroup_divisor: 4,
+        min_workgroups: 200,
+    },
+    // M3 series
+    GpuTier {
+        pattern: r"m3 ultra",
+        name: "Apple M3 Ultra",
+        workgroup_divisor: 4,
+        min_workgroups: 1520,
+    },
+    GpuTier {
+        pattern: r"m3 max",
+        name: "Apple M3 Max",
+        workgroup_divisor: 4,
+        min_workgroups: 800,
+    },
+    GpuTier {
+        pattern: r"m3 pro",
+        name: "Apple M3 Pro",
+        workgroup_divisor: 4,
+        min_workgroups: 360,
+    },
+    GpuTier {
+        pattern: r"\bm3\b",
+        name: "Apple M3",
+        workgroup_divisor: 4,
+        min_workgroups: 200,
+    },
+    // M2 series
+    GpuTier {
+        pattern: r"m2 ultra",
+        name: "Apple M2 Ultra",
+        workgroup_divisor: 4,
+        min_workgroups: 1520,
+    },
+    GpuTier {
+        pattern: r"m2 max",
+        name: "Apple M2 Max",
+        workgroup_divisor: 4,
+        min_workgroups: 760,
+    },
+    GpuTier {
+        pattern: r"m2 pro",
+        name: "Apple M2 Pro",
+        workgroup_divisor: 4,
+        min_workgroups: 380,
+    },
+    GpuTier {
+        pattern: r"\bm2\b",
+        name: "Apple M2",
+        workgroup_divisor: 4,
+        min_workgroups: 200,
+    },
+    // M1 series
+    GpuTier {
+        pattern: r"m1 ultra",
+        name: "Apple M1 Ultra",
+        workgroup_divisor: 4,
+        min_workgroups: 1280,
+    },
+    GpuTier {
+        pattern: r"m1 max",
+        name: "Apple M1 Max",
+        workgroup_divisor: 4,
+        min_workgroups: 640,
+    },
+    GpuTier {
+        pattern: r"m1 pro",
+        name: "Apple M1 Pro",
+        workgroup_divisor: 4,
+        min_workgroups: 320,
+    },
+    GpuTier {
+        pattern: r"\bm1\b",
+        name: "Apple M1",
+        workgroup_divisor: 4,
+        min_workgroups: 160,
+    },
+];
+
+/// Compiled GPU tier tables (lazily initialized)
+struct GpuTierTables {
+    nvidia: Vec<CompiledGpuTier>,
+    amd: Vec<CompiledGpuTier>,
+    intel: Vec<CompiledGpuTier>,
+    qualcomm: Vec<CompiledGpuTier>,
+    apple: Vec<CompiledGpuTier>,
+}
+
+static GPU_TIERS: LazyLock<GpuTierTables> = LazyLock::new(|| GpuTierTables {
+    nvidia: NVIDIA_TIERS
+        .iter()
+        .map(CompiledGpuTier::from_tier)
+        .collect(),
+    amd: AMD_TIERS.iter().map(CompiledGpuTier::from_tier).collect(),
+    intel: INTEL_TIERS.iter().map(CompiledGpuTier::from_tier).collect(),
+    qualcomm: QUALCOMM_TIERS
+        .iter()
+        .map(CompiledGpuTier::from_tier)
+        .collect(),
+    apple: APPLE_TIERS.iter().map(CompiledGpuTier::from_tier).collect(),
+});
+
+/// Result of GPU tier matching
+pub struct GpuTierMatch {
+    pub name: &'static str,
+    pub workgroup_divisor: u32,
+    pub min_workgroups: u32,
+    pub is_fallback: bool,
+}
+
+/// Find matching tier from a list of compiled tiers
+fn find_matching_tier(name: &str, tiers: &[CompiledGpuTier]) -> Option<GpuTierMatch> {
+    tiers
+        .iter()
+        .find(|tier| tier.regex.is_match(name))
+        .map(|tier| GpuTierMatch {
+            name: tier.name,
+            workgroup_divisor: tier.workgroup_divisor,
+            min_workgroups: tier.min_workgroups,
+            is_fallback: false,
+        })
+}
+
+/// Detect GPU tier based on adapter info
+///
+/// Returns (tier_name, workgroup_divisor, min_workgroups, is_fallback)
+pub fn detect_gpu_tier(vendor_name: &str, vendor_id: u32, is_metal_backend: bool) -> GpuTierMatch {
+    let name_lower = vendor_name.to_lowercase();
+
+    // Determine vendor and find matching tier
+    if name_lower.contains("nvidia") || vendor_id == 0x10DE {
+        find_matching_tier(&name_lower, &GPU_TIERS.nvidia).unwrap_or(GpuTierMatch {
+            name: "NVIDIA Unknown",
+            workgroup_divisor: 20,
+            min_workgroups: 512,
+            is_fallback: true,
+        })
+    } else if name_lower.contains("amd") || name_lower.contains("radeon") || vendor_id == 0x1002 {
+        find_matching_tier(&name_lower, &GPU_TIERS.amd).unwrap_or(GpuTierMatch {
+            name: "AMD Unknown",
+            workgroup_divisor: 24,
+            min_workgroups: 512,
+            is_fallback: true,
+        })
+    } else if name_lower.contains("intel") || vendor_id == 0x8086 {
+        find_matching_tier(&name_lower, &GPU_TIERS.intel).unwrap_or(GpuTierMatch {
+            name: "Intel Unknown",
+            workgroup_divisor: 24,
+            min_workgroups: 256,
+            is_fallback: true,
+        })
+    } else if name_lower.contains("qualcomm")
+        || name_lower.contains("adreno")
+        || vendor_id == 0x5143
+    {
+        find_matching_tier(&name_lower, &GPU_TIERS.qualcomm).unwrap_or(GpuTierMatch {
+            name: "Qualcomm Adreno (Unknown)",
+            workgroup_divisor: 24,
+            min_workgroups: 384,
+            is_fallback: true,
+        })
+    } else if is_metal_backend {
+        // Apple Silicon - detected by Metal backend
+        find_matching_tier(&name_lower, &GPU_TIERS.apple).unwrap_or(GpuTierMatch {
+            name: "Apple Silicon Unknown",
+            workgroup_divisor: 4,
+            min_workgroups: 160,
+            is_fallback: true,
+        })
+    } else {
+        GpuTierMatch {
+            name: "Unknown GPU",
+            workgroup_divisor: 16,
+            min_workgroups: 512,
+            is_fallback: true,
+        }
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_nvidia_detection() {
+        // RTX 40 series
+        let tier = detect_gpu_tier("NVIDIA GeForce RTX 4090", 0x10DE, false);
+        assert_eq!(tier.name, "NVIDIA RTX 40 Flagship (Ada)");
+        assert!(!tier.is_fallback);
+
+        let tier = detect_gpu_tier("NVIDIA GeForce RTX 4070 Ti", 0x10DE, false);
+        assert_eq!(tier.name, "NVIDIA RTX 40 (Ada)");
+
+        // RTX 30 series
+        let tier = detect_gpu_tier("NVIDIA GeForce RTX 3080", 0x10DE, false);
+        assert_eq!(tier.name, "NVIDIA RTX 30/20 (Ampere/Turing)");
+
+        // GTX 10 series
+        let tier = detect_gpu_tier("NVIDIA GeForce GTX 1080 Ti", 0x10DE, false);
+        assert_eq!(tier.name, "NVIDIA GTX 16/10 (Turing/Pascal)");
+    }
+
+    #[test]
+    fn test_amd_rdna_vs_polaris() {
+        // RDNA 1 - should NOT match Polaris
+        let tier = detect_gpu_tier("AMD Radeon RX 5500 XT", 0x1002, false);
+        assert_eq!(tier.name, "AMD RX 5000 (RDNA 1)");
+        assert!(!tier.is_fallback);
+
+        let tier = detect_gpu_tier("AMD Radeon RX 5600 XT", 0x1002, false);
+        assert_eq!(tier.name, "AMD RX 5000 (RDNA 1)");
+
+        let tier = detect_gpu_tier("AMD Radeon RX 5700 XT", 0x1002, false);
+        assert_eq!(tier.name, "AMD RX 5700 (RDNA 1)");
+
+        // Polaris - should match Polaris
+        let tier = detect_gpu_tier("AMD Radeon RX 580", 0x1002, false);
+        assert_eq!(tier.name, "AMD RX 500/400 (Polaris)");
+
+        let tier = detect_gpu_tier("AMD Radeon RX 560X", 0x1002, false);
+        assert_eq!(tier.name, "AMD RX 500/400 (Polaris)");
+
+        let tier = detect_gpu_tier("AMD Radeon RX 550", 0x1002, false);
+        assert_eq!(tier.name, "AMD RX 500/400 (Polaris)");
+    }
+
+    #[test]
+    fn test_amd_vega_apu() {
+        let tier = detect_gpu_tier("AMD Radeon(TM) Vega 8 Graphics", 0x1002, false);
+        assert_eq!(tier.name, "AMD Vega (APU)");
+        assert!(!tier.is_fallback);
+    }
+
+    #[test]
+    fn test_intel_arc_mobile_vs_desktop() {
+        // Mobile should match mobile tier
+        let tier = detect_gpu_tier("Intel Arc A770M", 0x8086, false);
+        assert_eq!(tier.name, "Intel Arc A7 Mobile (Alchemist)");
+        assert!(!tier.is_fallback);
+
+        // Desktop should match desktop tier
+        let tier = detect_gpu_tier("Intel Arc A770", 0x8086, false);
+        assert_eq!(tier.name, "Intel Arc A7 (Alchemist)");
+        assert!(!tier.is_fallback);
+    }
+
+    #[test]
+    fn test_apple_silicon() {
+        let tier = detect_gpu_tier("Apple M1 Pro", 0, true);
+        assert_eq!(tier.name, "Apple M1 Pro");
+        assert!(!tier.is_fallback);
+
+        let tier = detect_gpu_tier("Apple M3 Max", 0, true);
+        assert_eq!(tier.name, "Apple M3 Max");
+        assert!(!tier.is_fallback);
+    }
+
+    #[test]
+    fn test_qualcomm_adreno_series() {
+        // Adreno 627 should match 600 series, NOT 700 series
+        let tier = detect_gpu_tier("Qualcomm Adreno 627", 0x5143, false);
+        assert_eq!(tier.name, "Qualcomm Adreno 600 Series");
+        assert!(!tier.is_fallback);
+
+        // Adreno 730 should match 700 series
+        let tier = detect_gpu_tier("Qualcomm Adreno 730", 0x5143, false);
+        assert_eq!(tier.name, "Qualcomm Adreno 700 Series");
+        assert!(!tier.is_fallback);
+
+        // Adreno 540 should match 500 series
+        let tier = detect_gpu_tier("Qualcomm Adreno 540", 0x5143, false);
+        assert_eq!(tier.name, "Qualcomm Adreno 500 Series");
+        assert!(!tier.is_fallback);
+    }
+}
@@ -1,177 +1,141 @@
-use futures::executor::block_on;
-
 mod end_to_end_tests;
 mod tests;

-fn main() {
-    block_on(run()).unwrap();
-}
-
-async fn run() -> Result<(), Box<dyn std::error::Error>> {
-    // Setup GPU device and queue
+#[tokio::main]
+async fn main() -> Result<(), Box<dyn std::error::Error>> {
    let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
-        backends: wgpu::Backends::METAL, // Force Metal on Apple
+        backends: wgpu::Backends::PRIMARY,
        ..Default::default()
    });
    let adapter = instance
        .request_adapter(&wgpu::RequestAdapterOptions::default())
        .await
-        .unwrap();
-    let (device, queue) = adapter
-        .request_device(&wgpu::DeviceDescriptor::default())
-        .await?;
+        .expect("no GPU adapter");

-    println!("Running Poseidon2 GPU Component Tests...\n");
+    let mut failures = 0usize;

-    // Run all component tests
-    if let Err(e) = tests::test_gf_from_const(&device, &queue).await {
-        eprintln!("❌ gf_from_const tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_gf_mul(&device, &queue).await {
-        eprintln!("❌ gf_mul tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_sbox(&device, &queue).await {
-        eprintln!("❌ S-box tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_mds_matrix(&device, &queue).await {
-        eprintln!("❌ MDS matrix tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_internal_linear_layer(&device, &queue).await {
-        eprintln!("❌ Internal linear layer tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_external_linear_layer(&device, &queue).await {
-        eprintln!("❌ External linear layer tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_poseidon2_initial_external_rounds(&device, &queue).await {
-        eprintln!("❌ Initial external rounds tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_poseidon2_terminal_external_rounds(&device, &queue).await {
-        eprintln!("❌ Terminal external rounds tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_poseidon2_constants_verification(&device, &queue).await {
-        eprintln!("❌ Constants verification test failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_poseidon2_internal_constants_verification(&device, &queue).await {
-        eprintln!("❌ Internal constants verification test failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_poseidon2_internal_rounds_only(&device, &queue).await {
-        eprintln!("❌ Internal rounds only test failed: {}", e);
-    }
-
-    if let Err(e) =
-        tests::test_poseidon2_terminal_external_constants_verification(&device, &queue).await
    {
-        eprintln!(
-            "❌ Terminal external constants verification test failed: {}",
-            e
-        );
+        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 let Err(e) = tests::test_poseidon2_permutation(&device, &queue).await {
-        eprintln!("❌ Poseidon2 permutation tests failed: {}", e);
+    if adapter.features().contains(wgpu::Features::SHADER_INT64) {
+        let (device, queue) = adapter
+            .request_device(&wgpu::DeviceDescriptor {
+                label: Some("u64 Test Device"),
+                required_features: wgpu::Features::SHADER_INT64,
+                ..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");
    }

-    if let Err(e) = tests::test_bytes_to_field_elements(&device, &queue).await {
-        eprintln!("❌ Bytes to field elements tests failed: {}", e);
+    if failures > 0 {
+        return Err(format!("{failures} test group(s) failed").into());
    }
-
-    if let Err(e) = tests::test_field_elements_to_bytes(&device, &queue).await {
-        eprintln!("❌ Field elements to bytes tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_poseidon2_squeeze_twice(&device, &queue).await {
-        eprintln!("❌ Poseidon2 squeeze-twice tests failed: {}", e);
-    }
-
-    if let Err(e) = tests::test_hash_squeeze_twice(&device, &queue).await {
-        eprintln!("❌ Hash squeeze twice tests failed: {}", e);
-    }
-
-    if let Err(e) = end_to_end_tests::test_end_to_end_mining(&device, &queue).await {
-        eprintln!("❌ End-to-end mining test failed: {}", e);
-    }
-
    println!("\nAll tests completed!");
-    if let Err(e) = end_to_end_tests::test_end_to_end_mining(&device, &queue).await {
-        eprintln!("❌ End-to-end mining test failed: {}", e);
-    }
-
-    println!("\nAll tests completed!");
-
-    // generate_correct_wgsl_constants();
    Ok(())
 }

-#[allow(dead_code)]
-fn generate_correct_wgsl_constants() {
-    use qp_poseidon_constants::*;
-
-    println!("🔧 Generating correct WGSL constants...");
-
-    println!("// Initial external round constants (4 rounds x 12 elements)");
-    println!("const INITIAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(");
-
-    for (round_idx, round) in POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW.iter().enumerate() {
-        println!("    array<array<u32, 2>, 12>(");
-        for (elem_idx, &value) in round.iter().enumerate() {
-            let low = value as u32;
-            let high = (value >> 32) as u32;
-            if elem_idx == 11 {
-                println!("        array<u32, 2>({}u, {}u)", low, high);
-            } else {
-                println!("        array<u32, 2>({}u, {}u),", low, high);
+async fn run_suite(
+    device: &wgpu::Device,
+    queue: &wgpu::Queue,
+    shader_src: &str,
+    label: &str,
+) -> usize {
+    println!("\n==== Running Poseidon2 GPU tests against {label} ====\n");
+    let mut failures = 0usize;
+    macro_rules! run {
+        ($name:literal, $fut:expr) => {
+            if let Err(e) = $fut.await {
+                eprintln!("❌ {} failed: {}", $name, e);
+                failures += 1;
            }
-        }
-        if round_idx == 3 {
-            println!("    )");
-        } else {
-            println!("    ),");
-        }
+        };
    }
-    println!(");");

-    println!("\n// Terminal external round constants (4 rounds x 12 elements)");
-    println!("const TERMINAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(");
+    run!(
+        "gf_from_const",
+        tests::test_gf_from_const(device, queue, shader_src)
+    );
+    run!("gf_mul", tests::test_gf_mul(device, queue, shader_src));
+    run!("sbox", tests::test_sbox(device, queue, shader_src));
+    run!(
+        "mds matrix",
+        tests::test_mds_matrix(device, queue, shader_src)
+    );
+    run!(
+        "internal linear layer",
+        tests::test_internal_linear_layer(device, queue, shader_src)
+    );
+    run!(
+        "external linear layer",
+        tests::test_external_linear_layer(device, queue, shader_src)
+    );
+    run!(
+        "initial external rounds",
+        tests::test_poseidon2_initial_external_rounds(device, queue, shader_src)
+    );
+    run!(
+        "terminal external rounds",
+        tests::test_poseidon2_terminal_external_rounds(device, queue, shader_src)
+    );
+    run!(
+        "constants verification",
+        tests::test_poseidon2_constants_verification(device, queue, shader_src)
+    );
+    run!(
+        "internal constants verification",
+        tests::test_poseidon2_internal_constants_verification(device, queue, shader_src)
+    );
+    run!(
+        "internal rounds only",
+        tests::test_poseidon2_internal_rounds_only(device, queue, shader_src)
+    );
+    run!(
+        "terminal external constants verification",
+        tests::test_poseidon2_terminal_external_constants_verification(device, queue, shader_src)
+    );
+    run!(
+        "poseidon2 permutation",
+        tests::test_poseidon2_permutation(device, queue, shader_src)
+    );
+    run!(
+        "bytes to field elements",
+        tests::test_bytes_to_field_elements(device, queue, shader_src)
+    );
+    run!(
+        "field elements to bytes",
+        tests::test_field_elements_to_bytes(device, queue, shader_src)
+    );
+    run!(
+        "poseidon2 squeeze-twice",
+        tests::test_poseidon2_squeeze_twice(device, queue, shader_src)
+    );
+    run!(
+        "hash squeeze twice",
+        tests::test_hash_squeeze_twice(device, queue, shader_src)
+    );
+    run!(
+        "end-to-end mining",
+        end_to_end_tests::test_end_to_end_mining(device, queue, shader_src)
+    );

-    for (round_idx, round) in POSEIDON2_TERMINAL_EXTERNAL_CONSTANTS_RAW.iter().enumerate() {
-        println!("    array<array<u32, 2>, 12>(");
-        for (elem_idx, &value) in round.iter().enumerate() {
-            let low = value as u32;
-            let high = (value >> 32) as u32;
-            if elem_idx == 11 {
-                println!("        array<u32, 2>({}u, {}u)", low, high);
-            } else {
-                println!("        array<u32, 2>({}u, {}u),", low, high);
-            }
-        }
-        if round_idx == 3 {
-            println!("    )");
-        } else {
-            println!("    ),");
-        }
-    }
-    println!(");");
-
-    println!("\n// Internal round constants (22 values)");
-    println!("const INTERNAL_CONSTANTS: array<array<u32, 2>, 22> = array<array<u32, 2>, 22>(");
-    for (idx, &value) in POSEIDON2_INTERNAL_CONSTANTS_RAW.iter().enumerate() {
-        let low = value as u32;
-        let high = (value >> 32) as u32;
-        if idx == 21 {
-            println!("    array<u32, 2>({}u, {}u)", low, high);
-        } else {
-            println!("    array<u32, 2>({}u, {}u),", low, high);
-        }
-    }
-    println!(");");
+    failures
 }
@@ -191,7 +191,9 @@ const MDS_MATRIX_DIAG_12: array<array<u32, 2>, 12> = array<array<u32, 2>, 12>(

 // Storage buffers
@group(0) @binding(0) var<storage, read_write> results: array<atomic<u32>>;
-@group(0) @binding(1) var<storage, read> header: array<u32, 8>;     // 32 bytes
+// 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>; // 64 bytes
@group(0) @binding(3) var<storage, read> difficulty_target: array<u32, 16>;    // 64 bytes (U512 target)
@group(0) @binding(4) var<storage, read> dispatch_config: array<u32, 3>; // [total_threads, nonces_per_thread, total_nonces]
@@ -787,17 +789,8 @@ fn hash_squeeze_twice(input: array<u32, 24>) -> array<u32, 16> {
    return poseidon2_hash_squeeze_twice(input);
 }

-// Check if hash < target (U512 comparison)
-fn is_below_target(hash: array<u32, 16>, difficulty_tgt: array<u32, 16>) -> bool {
-    // Compare from most significant to least significant
-    for (var i = 0u; i < 16u; i++) {
-        if (hash[15u - i] < difficulty_tgt[15u - i]) {
-            return true;
-        } else if (hash[15u - i] > difficulty_tgt[15u - i]) {
-            return false;
-        }
-    }
-    return false; // Equal, not below
+fn bswap32(v: u32) -> u32 {
+    return ((v & 0xFFu) << 24u) | ((v & 0xFF00u) << 8u) | ((v >> 8u) & 0xFF00u) | (v >> 24u);
 }

 // Main mining kernel
@@ -834,61 +827,90 @@ fn mining_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
            return;
        }

-        // current_nonce = start_nonce + logical_index
+        // current_nonce = start_nonce + logical_index. The host guarantees a batch
+        // never carries into the high nonce half (limbs 8..15).
        var current_nonce: array<u32, 16>;
        var carry: u32 = 0u;

-        // Add logical_index into the low limb and propagate carry through the U512 nonce
        let val0 = start_nonce[0];
        let sum0 = val0 + logical_index;
        current_nonce[0] = sum0;
        carry = select(0u, 1u, sum0 < val0);

-        // Propagate carry through remaining limbs
-        for (var i = 1u; i < 16u; i++) {
+        for (var i = 1u; i < 8u; i++) {
            let val = start_nonce[i];
            let sum = val + carry;
            current_nonce[i] = sum;
            carry = select(0u, 1u, sum < val);
        }
-
-        // Construct input (96 bytes = 24 u32s)
-        // Header (32 bytes = 8 u32s) followed by Nonce (64 bytes = 16 u32s)
-        var input: array<u32, 24>;
-        for (var i = 0u; i < 8u; i++) {
-            input[i] = header[i];
+        for (var i = 8u; i < 16u; i++) {
+            current_nonce[i] = start_nonce[i];
        }
-        // Nonce needs to be Big Endian in the byte stream for hashing.
-        // current_nonce is Little Endian words.
-        for (var i = 0u; i < 16u; i++) {
-            let val = current_nonce[15u - i];
-            // Reverse bytes
+
+        // Resume the sponge from the precomputed midstate: absorb the low nonce
+        // half (as Big Endian words), pad, squeeze twice (3 permutations instead of 5).
+        var state: array<GoldilocksField, 12>;
+        for (var i = 0u; i < 12u; i++) {
+            state[i] = gf_from_limbs(midstate[2u * i], midstate[2u * i + 1u]);
+        }
+        for (var i = 0u; i < 8u; i++) {
+            let val = current_nonce[7u - i];
            let rev = ((val & 0xFFu) << 24u) |
                      ((val & 0xFF00u) << 8u) |
                      ((val & 0xFF0000u) >> 8u) |
                      ((val & 0xFF000000u) >> 24u);
-            input[8u + i] = rev;
+            state[i] = gf_add(state[i], gf_from_u32(rev));
        }
+        poseidon2_permute(&state);
+        state[0] = gf_add(state[0], gf_one());
+        state[1] = gf_add(state[1], gf_one());
+        poseidon2_permute(&state);

-        // Hash (Big Endian)
-        let hash_be = hash_squeeze_twice(input);
-
-        // Convert to Little Endian for difficulty check and storage
+        // First squeeze 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. Only candidates pay for the second squeeze.
+        let first_output = field_elements_to_bytes(array<GoldilocksField, 4>(
+            state[0], state[1], state[2], state[3]
+        ));
        var hash_le: array<u32, 16>;
-        for (var i = 0u; i < 16u; i++) {
-            let val = hash_be[15u - i];
-            // Reverse bytes in u32
-            hash_le[i] = ((val & 0xFFu) << 24u) |
-                         ((val & 0xFF00u) << 8u) |
-                         ((val & 0xFF0000u) >> 8u) |
-                         ((val & 0xFF000000u) >> 24u);
+        for (var i = 0u; i < 8u; i++) {
+            hash_le[15u - i] = bswap32(first_output[i]);
+        }
+        var cmp = 0u;
+        for (var i = 0u; i < 8u; i++) {
+            let h = hash_le[15u - i];
+            let t = difficulty_target[15u - i];
+            if (h != t) {
+                cmp = select(2u, 1u, h > t);
+                break;
+            }
+        }
+        if (cmp == 1u) {
+            continue;
        }

-        // Check target
-        if (is_below_target(hash_le, difficulty_target)) {
+        poseidon2_permute(&state);
+        let second_output = field_elements_to_bytes(array<GoldilocksField, 4>(
+            state[0], state[1], state[2], state[3]
+        ));
+        for (var i = 0u; i < 8u; i++) {
+            hash_le[i] = bswap32(second_output[7u - i]);
+        }
+        var below = cmp == 2u;
+        if (!below) {
+            for (var i = 0u; i < 8u; i++) {
+                let h = hash_le[7u - i];
+                let t = difficulty_target[7u - i];
+                if (h != t) {
+                    below = h < t;
+                    break;
+                }
+            }
+        }
+
+        if (below) {
            // Try to claim the solution
            if (atomicExchange(&results[0], 1u) == 0u) {
-                // We won! Write nonce and hash
                // results layout: [0]=found, [1..16]=nonce, [17..32]=hash
                for (var i = 0u; i < 16u; i++) {
                    atomicStore(&results[1u + i], current_nonce[i]);
@@ -0,0 +1,575 @@
+// 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.
+
+@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;
+
+const RC_INTERNAL: array<u64, 22> = array<u64, 22>(
+    0x97f7798a784ad863lu, 0xd1d2bf082f60d4f0lu, 0x69a377a79f9ad206lu, 0xa9d06906a3858e24lu, 0x295275001eede5b5lu, 0x5874e441117bd746lu, 0x8a084bbba8ed86cclu, 0x3defd7645cde6425lu, 0x3998cfe6871cc137lu, 0x3e52ef8bca48314alu, 0x964a209f85dc9ecclu, 0x3fcc9ee82cc4577elu, 0x8e79b4a5d0096d6dlu, 0x8492362ad2392556lu, 0xee72f470262574d6lu, 0x1e0e18496da2444alu, 0x0f3a74bf215eaac6lu, 0x1b061b76a1c0ded3lu, 0x192c42d86803d7a6lu, 0xf6d49ff997ae0260lu, 0x3ec372e7a0fa3786lu, 0x5538cdf4f23445d3lu
+);
+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);
+}
+
+// Reduce a 128-bit value (lo + hi*2^64) mod P using
+// 2^64 ≡ EPS64 and 2^96 ≡ -1 (mod P).
+fn gf64_reduce(lo: u64, hi: u64) -> u64 {
+    let hi_hi = hi >> 32u;
+    let hi_lo = hi & EPS64;
+    var t0 = lo - hi_hi;
+    t0 = t0 - select(0lu, EPS64, lo < hi_hi);
+    let t1 = hi_lo * EPS64;
+    let t2 = t0 + t1;
+    return t2 + select(0lu, EPS64, t2 < t0);
+}
+
+fn gf64_mul(a: u64, b: u64) -> u64 {
+    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 = lh + hl;
+    let mid_c = select(0lu, 1lu, mid < lh);
+    let lo = ll + (mid << 32u);
+    let lo_c = select(0lu, 1lu, lo < ll);
+    let hi = hh + (mid >> 32u) + (mid_c << 32u) + lo_c;
+    return gf64_reduce(lo, hi);
+}
+
+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 = lh << 1u;
+    let mid_c = lh >> 63u;
+    let lo = ll + (mid << 32u);
+    let lo_c = select(0lu, 1lu, lo < ll);
+    let hi = hh + (mid >> 32u) + (mid_c << 32u) + lo_c;
+    return gf64_reduce(lo, hi);
+}
+
+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);
+}
+
+// External linear layer: 4x4 MDS on each chunk, then circulant sums.
+fn ext_layer64(state: ptr<function, array<u64, 12>>) {
+    for (var chunk = 0u; chunk < 3u; chunk++) {
+        let o = chunk * 4u;
+        let x0 = (*state)[o];
+        let x1 = (*state)[o + 1u];
+        let x2 = (*state)[o + 2u];
+        let x3 = (*state)[o + 3u];
+        let t01 = gf64_add(x0, x1);
+        let t23 = gf64_add(x2, x3);
+        let t0123 = gf64_add(t01, t23);
+        let t01123 = gf64_add(t0123, x1);
+        let t01233 = gf64_add(t0123, x3);
+        (*state)[o + 3u] = gf64_add(t01233, gf64_add(x0, x0));
+        (*state)[o + 1u] = gf64_add(t01123, gf64_add(x2, x2));
+        (*state)[o] = gf64_add(t01123, t01);
+        (*state)[o + 2u] = gf64_add(t01233, t23);
+    }
+    var sums: array<u64, 4>;
+    for (var k = 0u; k < 4u; k++) {
+        sums[k] = gf64_add(gf64_add((*state)[k], (*state)[k + 4u]), (*state)[k + 8u]);
+    }
+    for (var i = 0u; i < 12u; i++) {
+        (*state)[i] = gf64_add((*state)[i], sums[i % 4u]);
+    }
+}
+
+// Internal linear layer: diagonal matrix plus full sum.
+fn int_layer64(state: ptr<function, array<u64, 12>>) {
+    var sum = (*state)[0];
+    for (var i = 1u; i < 12u; i++) {
+        sum = gf64_add(sum, (*state)[i]);
+    }
+    for (var i = 0u; i < 12u; i++) {
+        (*state)[i] = gf64_add(gf64_mul((*state)[i], MDS_DIAG[i]), sum);
+    }
+}
+
+fn permute64(state: ptr<function, array<u64, 12>>) {
+    ext_layer64(state);
+    for (var r = 0u; r < 4u; r++) {
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_add((*state)[i], RC_INITIAL[r][i]);
+        }
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_sbox((*state)[i]);
+        }
+        ext_layer64(state);
+    }
+    for (var r = 0u; r < 22u; r++) {
+        (*state)[0] = gf64_sbox(gf64_add((*state)[0], RC_INTERNAL[r]));
+        int_layer64(state);
+    }
+    for (var r = 0u; r < 4u; r++) {
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_add((*state)[i], RC_TERMINAL[r][i]);
+        }
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_sbox((*state)[i]);
+        }
+        ext_layer64(state);
+    }
+}
+
+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: array<u64, 12>;
+        for (var i = 0u; i < 12u; i++) {
+            st[i] = mid[i];
+        }
+        for (var i = 0u; i < 8u; i++) {
+            st[i] = gf64_add(st[i], u64(bswap32(current_nonce[7u - i])));
+        }
+        permute64(&st);
+        st[0] = gf64_add(st[0], 1lu);
+        st[1] = gf64_add(st[1], 1lu);
+        permute64(&st);
+
+        // First squeeze 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. Only candidates pay for the second squeeze, and
+        // byte-swapped hash words are produced on demand during the compare.
+        var first: array<u32, 8>;
+        for (var i = 0u; i < 4u; i++) {
+            let c = gf64_canon(st[i]);
+            first[2u * i] = u32(c & EPS64);
+            first[2u * i + 1u] = u32(c >> 32u);
+        }
+        var cmp = 0u;
+        for (var i = 0u; i < 8u; i++) {
+            let h = bswap32(first[i]);
+            let t = tgt[15u - i];
+            if (h != t) {
+                cmp = select(2u, 1u, h > t);
+                break;
+            }
+        }
+        if (cmp == 1u) {
+            continue;
+        }
+
+        var hash_le: array<u32, 16>;
+        for (var i = 0u; i < 8u; i++) {
+            hash_le[15u - i] = bswap32(first[i]);
+        }
+        permute64(&st);
+        for (var i = 0u; i < 4u; i++) {
+            let c = gf64_canon(st[i]);
+            hash_le[7u - 2u * i] = bswap32(u32(c & EPS64));
+            hash_le[6u - 2u * i] = bswap32(u32(c >> 32u));
+        }
+        var below = cmp == 2u;
+        if (!below) {
+            for (var i = 0u; i < 8u; i++) {
+                let h = hash_le[7u - i];
+                let t = tgt[7u - i];
+                if (h != t) {
+                    below = h < t;
+                    break;
+                }
+            }
+        }
+
+        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);
+    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);
+    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);
+}
@@ -1166,6 +1166,7 @@ fn generate_true_internal_only_test_vectors() -> Vec<InternalRoundsTestCase> {
 pub async fn test_poseidon2_permutation(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_poseidon2_test_vectors();
    let total_tests = test_vectors.len();
@@ -1202,7 +1203,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -1370,6 +1371,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_initial_external_rounds(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    // Test just the initial 4 external rounds
    let test_vectors = generate_gf_from_const_test_vectors();
@@ -1415,7 +1417,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -1630,6 +1632,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_terminal_external_rounds(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    // Test just the terminal 4 external rounds
    let _total_tests = 2;
@@ -1674,7 +1677,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -1888,6 +1891,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_gf_from_const(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    // Test the gf_from_const function that converts constant arrays to GoldilocksField
    let test_vectors = generate_gf_from_const_test_vectors();
@@ -1908,7 +1912,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    output_fields[index] = gf_from_const(input_constants[index]);
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2023,6 +2027,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_constants_verification(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let mut passed_tests = 0;
    let mut failed_tests = Vec::new();
@@ -2049,7 +2054,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2160,6 +2165,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_internal_rounds_only(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    // Test ONLY the internal rounds (22 rounds) without initial/terminal external rounds
    let test_vectors = generate_true_internal_only_test_vectors();
@@ -2209,7 +2215,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-            include_str!("mining.wgsl")
+            shader_src
        );

        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2343,6 +2349,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_internal_constants_verification(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let mut passed_tests = 0;
    let mut failed_tests = Vec::new();
@@ -2364,7 +2371,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2462,6 +2469,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_terminal_external_rounds_issue(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    println!("🔍 Testing terminal external rounds issue for sequential input...");

@@ -2502,7 +2510,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2630,6 +2638,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_terminal_external_constants_verification(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let mut passed_tests = 0;
    let mut failed_tests = Vec::new();
@@ -2656,7 +2665,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2764,6 +2773,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_gf_mul(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_gf_mul_test_vectors();
    let total_tests = test_vectors.len();
@@ -2786,7 +2796,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    output[i] = gf_mul(input_a[i], input_b[i]);
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2947,6 +2957,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_mds_matrix(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_mds_test_vectors();
    let total_tests = test_vectors.len();
@@ -2955,7 +2966,7 @@ pub async fn test_mds_matrix(
    let mut failure_details = Vec::new();

    // Read the mining shader source and create a test wrapper
-    let mining_shader_source = include_str!("mining.wgsl");
+    let mining_shader_source = shader_src;
    let shader_source = format!(
        "
 {}
@@ -3215,6 +3226,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_internal_linear_layer(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_internal_linear_layer_test_vectors();
    let total_tests = test_vectors.len();
@@ -3223,7 +3235,7 @@ pub async fn test_internal_linear_layer(
    let mut failure_details = Vec::new();

    // Read the mining shader source and create a test wrapper
-    let mining_shader_source = include_str!("mining.wgsl");
+    let mining_shader_source = shader_src;
    let shader_source = format!(
        "
 {}
@@ -3448,6 +3460,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_external_linear_layer(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_external_linear_layer_test_vectors();
    let total_tests = test_vectors.len();
@@ -3485,7 +3498,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -3625,6 +3638,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_sbox(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_sbox_test_vectors();
    let total_tests = test_vectors.len();
@@ -3662,7 +3676,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    output_data[index] = result;
 }}
 ",
-            include_str!("mining.wgsl"),
+            shader_src,
            batch.len()
        );

@@ -3807,6 +3821,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_poseidon2_squeeze_twice(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_poseidon2_squeeze_twice_test_vectors();
    let total_tests = test_vectors.len();
@@ -3842,7 +3857,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -3985,6 +4000,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_bytes_to_field_elements(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_bytes_to_field_test_vectors();
    let total_tests = test_vectors.len();
@@ -4020,7 +4036,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -4168,6 +4184,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_field_elements_to_bytes(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_field_to_bytes_test_vectors();
    let total_tests = test_vectors.len();
@@ -4203,7 +4220,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -4336,6 +4353,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
 pub async fn test_hash_squeeze_twice(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
+    shader_src: &str,
 ) -> Result<(), Box<dyn std::error::Error>> {
    let test_vectors = generate_hash_squeeze_twice_test_vectors();
    let total_tests = test_vectors.len();
@@ -4371,7 +4389,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
    }}
 }}
 ",
-        include_str!("mining.wgsl")
+        shader_src
    );

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -8,6 +8,7 @@ description = "CLI binary to run the Quantus External Miner service"
 [dependencies]
 miner-service = { path = "../miner-service" }
 metrics = { path = "../metrics", features = ["http-exporter"] }
+quic-transport = { path = "../quic-transport" }
 clap = { workspace = true, features = ["derive", "env"] }
 tokio = { workspace = true, features = ["full"] }
 env_logger = { workspace = true }
@@ -3,6 +3,7 @@ use engine_cpu::{AtomicBoolCancelCheck, EngineRange, MinerEngine};
 use miner_service::{run, ServiceConfig};
 use primitive_types::U512;
 use rand::RngCore;
+use std::path::PathBuf;
 use std::sync::atomic::AtomicBool;
 use std::sync::Arc;
 use std::thread;
@@ -20,6 +21,34 @@ enum Command {
        #[arg(long, env = "MINER_NODE_ADDR", default_value = "127.0.0.1:9833")]
        node_addr: std::net::SocketAddr,

+        /// Shared auth token from the node's `miner-auth-token` file
+        /// (`<base-path>/chains/<chain>/miner-auth-token`). Prefer `--auth-token-file`.
+        #[arg(long, env = "MINER_AUTH_TOKEN", conflicts_with = "auth_token_file")]
+        auth_token: Option<String>,
+
+        /// Path to the node's `miner-auth-token` file (trimmed). Preferred over
+        /// `--auth-token` so the secret is not placed on the command line.
+        #[arg(long, env = "MINER_AUTH_TOKEN_FILE", conflicts_with = "auth_token")]
+        auth_token_file: Option<PathBuf>,
+
+        /// SHA-256 fingerprint of the node's miner TLS certificate (64 hex chars).
+        /// Prefer `--tls-cert-sha256-file` pointing at `miner-tls-cert-sha256`
+        /// (the node also logs this fingerprint).
+        #[arg(
+            long,
+            env = "MINER_TLS_CERT_SHA256",
+            conflicts_with = "tls_cert_sha256_file"
+        )]
+        tls_cert_sha256: Option<String>,
+
+        /// Path to the node's `miner-tls-cert-sha256` file.
+        #[arg(
+            long,
+            env = "MINER_TLS_CERT_SHA256_FILE",
+            conflicts_with = "tls_cert_sha256"
+        )]
+        tls_cert_sha256_file: Option<PathBuf>,
+
        /// Number of CPU worker threads to use for mining (default: auto-detect)
        #[arg(long = "cpu-workers", env = "MINER_CPU_WORKERS")]
        cpu_workers: Option<usize>,
@@ -29,11 +58,11 @@ enum Command {
        gpu_devices: Option<usize>,

        /// GPU batch size in nonces - controls how often GPU checks for cancellation
-        #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE)]
+        #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE, value_parser = clap::value_parser!(u32).range(1..))]
        gpu_batch_size: u32,

        /// CPU batch size in hashes - controls how often CPU checks for cancellation
-        #[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE)]
+        #[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE, value_parser = clap::value_parser!(u64).range(1..))]
        cpu_batch_size: u64,

        /// Port for Prometheus metrics HTTP endpoint (default: 9900)
@@ -52,6 +81,12 @@ enum Command {
        )]
        gpu_throttle_ms: u64,

+        /// Allow integrated GPUs (APUs) even when discrete GPUs are available.
+        /// By default, integrated GPUs are skipped when a discrete GPU is present
+        /// to avoid resource contention and driver instability.
+        #[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
+        allow_integrated: bool,
+
        /// Enable verbose logging
        #[arg(short, long, env = "MINER_VERBOSE")]
        verbose: bool,
@@ -68,17 +103,21 @@ enum Command {
        gpu_devices: Option<usize>,

        /// GPU batch size in nonces - controls how often GPU checks for cancellation
-        #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE)]
+        #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE, value_parser = clap::value_parser!(u32).range(1..))]
        gpu_batch_size: u32,

        /// CPU batch size in hashes - controls how often CPU checks for cancellation
-        #[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE)]
+        #[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE, value_parser = clap::value_parser!(u64).range(1..))]
        cpu_batch_size: u64,

        /// Benchmark duration in seconds (default: 10)
        #[arg(short, long, default_value_t = 10)]
        duration: u64,

+        /// Allow integrated GPUs (APUs) even when discrete GPUs are available
+        #[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
+        allow_integrated: bool,
+
        /// Enable verbose logging
        #[arg(short, long, env = "MINER_VERBOSE")]
        verbose: bool,
@@ -99,23 +138,53 @@ async fn main() {

    let Some(command) = args.command else {
        eprintln!("Error: No command provided. Use 'serve' to start mining (defaults to local node at 127.0.0.1:9833).");
-        eprintln!("Example: quantus-miner serve --node-addr 127.0.0.1:9833");
+        eprintln!(
+            "Example: quantus-miner serve --node-addr 127.0.0.1:9833 \
+             --auth-token-file /path/to/miner-auth-token \
+             --tls-cert-sha256-file /path/to/miner-tls-cert-sha256"
+        );
        std::process::exit(1);
    };

    match command {
        Command::Serve {
            node_addr,
+            auth_token,
+            auth_token_file,
+            tls_cert_sha256,
+            tls_cert_sha256_file,
            cpu_workers,
            gpu_devices,
            gpu_batch_size,
            cpu_batch_size,
            gpu_throttle_ms,
            metrics_port,
+            allow_integrated,
            verbose,
        } => {
            init_logger(verbose);

+            let auth_token = match resolve_auth_token(auth_token, auth_token_file) {
+                Ok(token) => token,
+                Err(e) => {
+                    eprintln!("Error: {e}");
+                    std::process::exit(1);
+                }
+            };
+            let tls_cert_sha256 =
+                match resolve_tls_cert_sha256(tls_cert_sha256, tls_cert_sha256_file) {
+                    Ok(fp) => fp,
+                    Err(e) => {
+                        eprintln!("Error: {e}");
+                        std::process::exit(1);
+                    }
+                };
+            // Fail closed on permanent misconfig before metrics/workers start.
+            if let Err(e) = quic_transport::validate_auth_config(&auth_token, &tls_cert_sha256) {
+                eprintln!("Error: {e}");
+                std::process::exit(1);
+            }
+
            log::info!("Starting external miner service...");

            // Start metrics HTTP server
@@ -130,11 +199,14 @@ async fn main() {

            let config = ServiceConfig {
                node_addr,
+                auth_token,
+                tls_cert_sha256,
                cpu_workers,
                gpu_devices,
                gpu_batch_size,
                cpu_batch_size,
                gpu_throttle_ms,
+                allow_integrated,
            };

            if let Err(e) = run(config).await {
@@ -149,6 +221,7 @@ async fn main() {
            gpu_batch_size,
            cpu_batch_size,
            duration,
+            allow_integrated,
            verbose,
        } => {
            init_logger(verbose);
@@ -158,18 +231,71 @@ async fn main() {
                gpu_batch_size,
                cpu_batch_size,
                duration,
+                allow_integrated,
            )
            .await;
        }
    }
 }

+fn resolve_auth_token(
+    auth_token: Option<String>,
+    auth_token_file: Option<PathBuf>,
+) -> Result<String, String> {
+    resolve_required_secret(
+        auth_token,
+        auth_token_file,
+        "--auth-token",
+        "--auth-token-file",
+        "miner auth token required: pass --auth-token-file <PATH> to the node's \
+         miner-auth-token file (or --auth-token <TOKEN>)",
+    )
+}
+
+fn resolve_tls_cert_sha256(value: Option<String>, file: Option<PathBuf>) -> Result<String, String> {
+    resolve_required_secret(
+        value,
+        file,
+        "--tls-cert-sha256",
+        "--tls-cert-sha256-file",
+        "TLS cert fingerprint required: pass --tls-cert-sha256-file <PATH> to the node's \
+         miner-tls-cert-sha256 file (or --tls-cert-sha256 <HEX>; also printed in node logs)",
+    )
+}
+
+fn resolve_required_secret(
+    value: Option<String>,
+    file: Option<PathBuf>,
+    value_flag: &str,
+    file_flag: &str,
+    missing_msg: &str,
+) -> Result<String, String> {
+    if let Some(value) = value {
+        let value = value.trim().to_string();
+        if value.is_empty() {
+            return Err(format!("{value_flag} is empty"));
+        }
+        return Ok(value);
+    }
+    if let Some(path) = file {
+        let contents = std::fs::read_to_string(&path)
+            .map_err(|e| format!("failed to read {file_flag} {}: {}", path.display(), e))?;
+        let value = contents.trim().to_string();
+        if value.is_empty() {
+            return Err(format!("{file_flag} {} is empty", path.display()));
+        }
+        return Ok(value);
+    }
+    Err(missing_msg.into())
+}
+
 fn init_logger(verbose: bool) {
    if std::env::var("RUST_LOG").is_err() {
+        // Filter out noisy wgpu/naga shader compilation logs
        let log_level = if verbose {
-            "debug,miner=debug,gpu_engine=debug,engine_cpu=debug"
+            "debug,miner=debug,gpu_engine=debug,engine_cpu=debug,wgpu=warn,wgpu_core=warn,wgpu_hal=warn,naga=warn"
        } else {
-            "info,miner=info,gpu_engine=info"
+            "info,miner=info,gpu_engine=info,wgpu=error,wgpu_core=error,wgpu_hal=error,naga=error"
        };
        std::env::set_var("RUST_LOG", log_level);
    }
@@ -182,18 +308,23 @@ async fn run_benchmark(
    gpu_batch_size: u32,
    cpu_batch_size: u64,
    duration: u64,
+    allow_integrated: bool,
 ) {
    let effective_cpu_workers = cpu_workers.unwrap_or_else(num_cpus::get);

    // Initialize GPU engine (no throttle for benchmark)
-    let (gpu_engine, effective_gpu_devices) =
-        match miner_service::resolve_gpu_configuration(gpu_devices, gpu_batch_size, 0) {
-            Ok((engine, count)) => (engine, count),
-            Err(e) => {
-                eprintln!("❌ ERROR: {}", e);
-                std::process::exit(1);
-            }
-        };
+    let (gpu_engine, effective_gpu_devices) = match miner_service::resolve_gpu_configuration(
+        gpu_devices,
+        gpu_batch_size,
+        0,
+        allow_integrated,
+    ) {
+        Ok((engine, count)) => (engine, count),
+        Err(e) => {
+            eprintln!("❌ ERROR: {}", e);
+            std::process::exit(1);
+        }
+    };

    let total_workers = effective_cpu_workers + effective_gpu_devices;

@@ -205,7 +336,13 @@ async fn run_benchmark(
        num_cpus::get()
    );
    println!("GPU Devices: {}", effective_gpu_devices);
-    println!("Duration: {} seconds", duration);
+    if effective_cpu_workers > 0 {
+        println!("CPU batch size: {cpu_batch_size} hashes");
+    }
+    if effective_gpu_devices > 0 {
+        println!("GPU batch size: {gpu_batch_size} nonces");
+    }
+    println!("Duration: {duration} seconds");
    println!();

    if total_workers == 0 {
@@ -237,8 +374,11 @@ async fn run_benchmark(
    let mut handles = Vec::new();
    let total_hashes = Arc::new(std::sync::Mutex::new(0u64));

-    let cpu_chunk = 10_000u64;
-    let gpu_chunk = 1_000_000u64;
+    // Floor range widths at the old constants: engines still batch at the flag's
+    // size internally, but tiny flags don't turn per-call harness overhead into
+    // the measured quantity.
+    let cpu_chunk = cpu_batch_size.max(10_000);
+    let gpu_chunk = (gpu_batch_size as u64).max(1_000_000);

    for worker_id in 0..total_workers {
        let (engine, nonces_per_batch) = if worker_id < effective_cpu_workers {
@@ -254,31 +394,40 @@ async fn run_benchmark(

        let handle = thread::spawn(move || {
            let stride = U512::from(1_000_000_000_000u64);
-            let worker_start = U512::from(worker_id as u64).saturating_mul(stride);
-            let worker_range = EngineRange {
-                start: worker_start,
-                end: worker_start
-                    .saturating_add(U512::from(nonces_per_batch))
-                    .saturating_sub(U512::from(1u64)),
-            };
+            let mut current = U512::from(worker_id as u64).saturating_mul(stride);
+            let step = U512::from(nonces_per_batch);

            loop {
                if cancel.load(std::sync::atomic::Ordering::Relaxed) {
                    break;
                }

+                let worker_range = EngineRange {
+                    start: current,
+                    end: current
+                        .saturating_add(step)
+                        .saturating_sub(U512::from(1u64)),
+                };
+
                let cancel_check = AtomicBoolCancelCheck(&cancel);
-                let result = engine.search_range(&ctx, worker_range.clone(), &cancel_check);
+                let result = engine.search_range(&ctx, worker_range, &cancel_check);

                match result {
                    engine_cpu::EngineStatus::Found { hash_count, .. }
                    | engine_cpu::EngineStatus::Exhausted { hash_count }
-                    | engine_cpu::EngineStatus::Cancelled { hash_count } => {
+                    | engine_cpu::EngineStatus::Cancelled { hash_count }
+                    | engine_cpu::EngineStatus::DeviceLost { hash_count } => {
                        *hashes.lock().unwrap() += hash_count;
                    }
                    engine_cpu::EngineStatus::Running { .. } => {}
                }

+                if matches!(result, engine_cpu::EngineStatus::DeviceLost { .. }) {
+                    break;
+                }
+
+                current = current.saturating_add(step);
+
                if start.elapsed() >= Duration::from_secs(duration) {
                    break;
                }
@@ -21,7 +21,7 @@ anyhow = { workspace = true }

 # QUIC transport
 quinn = "0.10"
-rustls = { version = "0.21", default-features = false, features = ["dangerous_configuration", "quic"] }
+quic-transport = { path = "../quic-transport" }
 getrandom = "0.2"

 # Protocol types from the node repo
@@ -23,6 +23,10 @@ use std::thread;
 pub struct ServiceConfig {
    /// Address of the node to connect to (e.g., "127.0.0.1:9833").
    pub node_addr: std::net::SocketAddr,
+    /// Shared secret that must match the node's miner auth token.
+    pub auth_token: String,
+    /// SHA-256 fingerprint (hex) of the node's miner TLS certificate DER.
+    pub tls_cert_sha256: String,
    /// Number of CPU worker threads to use for mining (None = auto-detect)
    pub cpu_workers: Option<usize>,
    /// Number of GPU devices to use for mining (None = auto-detect)
@@ -33,6 +37,8 @@ pub struct ServiceConfig {
    pub cpu_batch_size: u64,
    /// GPU throttle delay in milliseconds between batches (0 = no throttle)
    pub gpu_throttle_ms: u64,
+    /// Allow integrated GPUs even when discrete GPUs are available
+    pub allow_integrated: bool,
 }

 /// Engine type for tracking metrics per compute type.
@@ -206,12 +212,34 @@ impl WorkerPool {

        // Dispatch job to all workers using bounded channels (capacity 16).
        // Workers drain to get the latest job, so we just need room to queue.
+        let mut disconnected_count = 0;
        for (i, tx) in self.job_senders.iter().enumerate() {
            if let Err(e) = tx.try_send(job.clone()) {
+                match e {
+                    crossbeam_channel::TrySendError::Disconnected(_) => {
+                        // Worker thread has exited (e.g., device lost)
+                        disconnected_count += 1;
+                        log::debug!("Worker {i} channel disconnected (worker exited)");
+                    }
+                    crossbeam_channel::TrySendError::Full(_) => {
+                        log::warn!(
+                            "Failed to send job {new_job_id} to worker {i}: channel full - \
+                             worker may be stuck or jobs arriving too fast"
+                        );
+                    }
+                }
+            }
+        }
+
+        if disconnected_count > 0 {
+            let active = self.job_senders.len() - disconnected_count;
+            if active == 0 {
                log::error!(
-                    "Failed to send job {new_job_id} to worker {i}: {e}. \
-                     Channel full - worker may be stuck or jobs arriving too fast."
+                    "All workers have exited! No workers available to process jobs. \
+                     Consider restarting the miner."
                );
+            } else {
+                log::warn!("{disconnected_count} worker(s) have exited, {active} still active");
            }
        }

@@ -336,6 +364,7 @@ fn worker_loop(
            engine_cpu::EngineStatus::Found { .. } => "FOUND",
            engine_cpu::EngineStatus::Exhausted { .. } => "EXHAUSTED",
            engine_cpu::EngineStatus::Cancelled { .. } => "CANCELLED",
+            engine_cpu::EngineStatus::DeviceLost { .. } => "DEVICE_LOST",
            engine_cpu::EngineStatus::Running { .. } => "RUNNING",
        };
        log::debug!(
@@ -352,6 +381,7 @@ fn worker_loop(
                engine_cpu::EngineStatus::Found { hash_count, .. } => hash_count,
                engine_cpu::EngineStatus::Exhausted { hash_count } => hash_count,
                engine_cpu::EngineStatus::Cancelled { hash_count } => hash_count,
+                engine_cpu::EngineStatus::DeviceLost { hash_count } => hash_count,
                engine_cpu::EngineStatus::Running { .. } => 0,
            };
            log_worker_completion(
@@ -380,7 +410,7 @@ fn worker_loop(
                ..
            } => {
                log::info!(
-                    "{type_str} worker {thread_id} found solution! Nonce: {}, Hash: {} (job {job_id})",
+                    "🎯 {type_str} worker {thread_id} found solution! Nonce: {}, Hash: {} (job {job_id})",
                    format_u512(nonce),
                    format_u512(hash),
                );
@@ -400,6 +430,21 @@ fn worker_loop(
                log_worker_completion(type_str, thread_id, "new block", hash_count, search_elapsed);
                (None, hash_count)
            }
+            engine_cpu::EngineStatus::DeviceLost { hash_count } => {
+                log::error!(
+                    "{type_str} worker {thread_id} GPU device lost - worker exiting permanently"
+                );
+                // Send final result before exiting
+                let _ = result_tx.try_send(WorkerResult {
+                    thread_id,
+                    engine_type,
+                    job_id,
+                    candidate: None,
+                    hash_count,
+                    completed: true,
+                });
+                break; // Exit the worker loop
+            }
            engine_cpu::EngineStatus::Running { .. } => {
                // Should not happen for synchronous search
                (None, 0)
@@ -430,6 +475,7 @@ pub fn resolve_gpu_configuration(
    requested_devices: Option<usize>,
    batch_size: u32,
    throttle_ms: u64,
+    allow_integrated: bool,
 ) -> anyhow::Result<(Option<Arc<dyn MinerEngine>>, usize)> {
    // Explicit 0 means no GPU
    if requested_devices == Some(0) {
@@ -437,7 +483,7 @@ pub fn resolve_gpu_configuration(
    }

    // Try to initialize GPU engine
-    let engine = engine_gpu::GpuEngine::try_new(batch_size, throttle_ms);
+    let engine = engine_gpu::GpuEngine::try_new(batch_size, throttle_ms, allow_integrated);
    let engine = match engine {
        Ok(e) => e,
        Err(e) => {
@@ -482,6 +528,7 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
        config.gpu_devices,
        config.gpu_batch_size,
        config.gpu_throttle_ms,
+        config.allow_integrated,
    )?;

    // Resolve CPU workers
@@ -539,6 +586,8 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
    log::info!("🌐 Connecting to node at {node_addr}");
    quic::connect_and_mine(
        config.node_addr,
+        &config.auth_token,
+        &config.tls_cert_sha256,
        cpu_engine,
        gpu_engine,
        cpu_workers,
@@ -10,8 +10,6 @@ use std::time::{Duration, Instant};

 use engine_cpu::MinerEngine;
 use primitive_types::U512;
-use quinn::{ClientConfig, Endpoint};
-use rustls::client::ServerCertVerified;

 use quantus_miner_api::{
    read_message, write_message, ApiResponseStatus, MinerMessage, MiningResult,
@@ -28,6 +26,8 @@ use pow_core::format_hashrate;
 /// Uses a persistent worker pool to avoid thread creation overhead between jobs.
 pub async fn connect_and_mine(
    node_addr: SocketAddr,
+    auth_token: &str,
+    tls_cert_sha256: &str,
    cpu_engine: Option<Arc<dyn MinerEngine>>,
    gpu_engine: Option<Arc<dyn MinerEngine>>,
    cpu_workers: usize,
@@ -42,18 +42,39 @@ pub async fn connect_and_mine(
    loop {
        log::info!("⛏️ Connecting to node at {}...", node_addr);

-        match establish_connection(node_addr).await {
+        match establish_connection(node_addr, auth_token, tls_cert_sha256).await {
            Ok((connection, send, recv)) => {
                log::info!("⛏️ Connected to node at {}", node_addr);
-                reconnect_delay = Duration::from_secs(1);

-                if let Err(e) = handle_connection(connection, send, recv, &worker_pool).await {
-                    log::info!("⛏️ Connection lost: {}", e);
+                let mut authenticated = false;
+                if let Err(e) =
+                    handle_connection(connection, send, recv, &worker_pool, &mut authenticated)
+                        .await
+                {
                    // Cancel any running job when connection drops
                    worker_pool.cancel();
+                    if e.downcast_ref::<quic_transport::PermanentConnectError>()
+                        .is_some()
+                    {
+                        log::error!("⛏️ Permanent connection error (not retrying): {e}");
+                        return Err(e);
+                    }
+                    log::info!("⛏️ Connection lost: {}", e);
+                }
+                // Only clear backoff after the node accepted auth (first NewJob).
+                // connect() already treats explicit "auth failed" as permanent;
+                // this covers any other post-Ready close that looked like success.
+                if authenticated {
+                    reconnect_delay = Duration::from_secs(1);
                }
            }
            Err(e) => {
+                if e.downcast_ref::<quic_transport::PermanentConnectError>()
+                    .is_some()
+                {
+                    log::error!("⛏️ Permanent connection error (not retrying): {e}");
+                    return Err(e);
+                }
                log::warn!("⛏️ Failed to connect to node: {}", e);
            }
        }
@@ -64,38 +85,18 @@ pub async fn connect_and_mine(
    }
 }

-/// Establish a QUIC connection to the node.
+/// Establish a QUIC connection to the node (shared transport crate).
 async fn establish_connection(
    addr: SocketAddr,
+    auth_token: &str,
+    tls_cert_sha256: &str,
 ) -> anyhow::Result<(quinn::Connection, quinn::SendStream, quinn::RecvStream)> {
-    let mut crypto = rustls::ClientConfig::builder()
-        .with_safe_defaults()
-        .with_custom_certificate_verifier(Arc::new(InsecureCertVerifier))
-        .with_no_client_auth();
-
-    crypto.alpn_protocols = vec![b"quantus-miner".to_vec()];
-
-    let mut client_config = ClientConfig::new(Arc::new(crypto));
-
-    let mut transport_config = quinn::TransportConfig::default();
-    transport_config.keep_alive_interval(Some(Duration::from_secs(5)));
-    transport_config.max_idle_timeout(Some(Duration::from_secs(15).try_into().unwrap()));
-    client_config.transport_config(Arc::new(transport_config));
-
-    let mut endpoint = Endpoint::client("0.0.0.0:0".parse().unwrap())?;
-    endpoint.set_default_client_config(client_config);
-
-    let connection = endpoint.connect(addr, "localhost")?.await?;
-    log::info!("⛏️ QUIC connection established to {}", addr);
-
-    log::info!("⛏️ Opening bidirectional stream to node...");
-    let (mut send, recv) = connection.open_bi().await?;
-
-    // Send Ready message to establish the stream
-    write_message(&mut send, &MinerMessage::Ready).await?;
-    log::info!("⛏️ Bidirectional stream established");
-
-    Ok((connection, send, recv))
+    let result = quic_transport::connect(addr, auth_token, tls_cert_sha256).await?;
+    log::info!(
+        "⛏️ QUIC connection and bidirectional stream established to {}",
+        addr
+    );
+    Ok(result)
 }

 /// Helper to send a message while monitoring connection health.
@@ -116,11 +117,15 @@ async fn send_message_checked(
 }

 /// Handle an established connection, receiving jobs and sending results.
+///
+/// Sets `authenticated` when the first `NewJob` arrives (proof the node
+/// accepted our Ready token).
 async fn handle_connection(
    connection: quinn::Connection,
    mut send: quinn::SendStream,
    mut recv: quinn::RecvStream,
    worker_pool: &WorkerPool,
+    authenticated: &mut bool,
 ) -> anyhow::Result<()> {
    use crossbeam_channel::RecvTimeoutError;

@@ -220,12 +225,20 @@ async fn handle_connection(
            biased;

            reason = connection.closed() => {
-                return Err(anyhow::anyhow!("Connection closed: {}", reason));
+                let msg = reason.to_string();
+                if !*authenticated && msg.to_ascii_lowercase().contains("auth") {
+                    return Err(quic_transport::PermanentConnectError(format!(
+                        "node rejected miner auth ({msg}); check --auth-token / miner-auth-token"
+                    ))
+                    .into());
+                }
+                return Err(anyhow::anyhow!("Connection closed: {}", msg));
            }

            msg_result = read_message(&mut recv) => {
                match msg_result {
                    Ok(MinerMessage::NewJob(request)) => {
+                        *authenticated = true;
                        log::info!(
                            "⛏️ Received job: id={}, hash=0x{}",
                            request.job_id,
@@ -305,7 +318,7 @@ async fn handle_connection(
                    Ok(MinerMessage::JobResult(_)) => {
                        log::warn!("Received unexpected JobResult from node");
                    }
-                    Ok(MinerMessage::Ready) => {
+                    Ok(MinerMessage::Ready { .. }) => {
                        log::warn!("Received unexpected Ready from node");
                    }
                    Err(e) if e.kind() == std::io::ErrorKind::UnexpectedEof => {
@@ -322,20 +335,3 @@ async fn handle_connection(
        }
    }
 }
-
-/// Certificate verifier that accepts any certificate (for self-signed certs).
-struct InsecureCertVerifier;
-
-impl rustls::client::ServerCertVerifier for InsecureCertVerifier {
-    fn verify_server_cert(
-        &self,
-        _end_entity: &rustls::Certificate,
-        _intermediates: &[rustls::Certificate],
-        _server_name: &rustls::ServerName,
-        _scts: &mut dyn Iterator<Item = &[u8]>,
-        _ocsp_response: &[u8],
-        _now: std::time::SystemTime,
-    ) -> Result<ServerCertVerified, rustls::Error> {
-        Ok(ServerCertVerified::assertion())
-    }
-}
@@ -0,0 +1,32 @@
+[package]
+name = "pool-service"
+version.workspace = true
+edition.workspace = true
+publish = false
+description = "Captcha share pool: turns browser proof-of-work captcha solves into Quantus mining shares"
+
+[dependencies]
+anyhow = { workspace = true }
+clap = { workspace = true }
+env_logger = { workspace = true }
+hex = { workspace = true }
+log = { workspace = true }
+primitive-types = { workspace = true }
+rand = { workspace = true }
+serde = { workspace = true }
+serde_json = { workspace = true, features = ["std"] }
+tokio = { workspace = true }
+warp = { workspace = true }
+
+# QUIC transport (shared with miner-service)
+quinn = "0.10"
+quic-transport = { path = "../quic-transport" }
+
+# Constant-time comparison for the site secret
+constant_time_eq = "0.3"
+
+# Protocol types from the node repo
+quantus-miner-api = { workspace = true }
+
+# Local crates
+pow-core = { path = "../pow-core" }
@@ -0,0 +1,219 @@
+//! HTTP API for the captcha widget (session/share) and for protected sites
+//! (siteverify), plus optional static file serving for the demo page.
+
+use std::net::{IpAddr, SocketAddr};
+use std::path::PathBuf;
+use std::sync::Arc;
+
+use primitive_types::U512;
+use serde::{Deserialize, Serialize};
+use warp::Filter;
+
+use crate::rate_limit::SessionIssuerLimiter;
+use crate::state::{PoolState, ShareOutcome};
+
+#[derive(Serialize)]
+struct SessionResponse {
+    session_id: String,
+    /// Hex, 32 bytes, no 0x prefix.
+    header_hash: String,
+    /// Hex U512, no 0x prefix. Solver iterates nonce_start..nonce_end.
+    nonce_start: String,
+    nonce_end: String,
+    /// Share target as hex U512: a share is valid iff hash < share_target.
+    /// (Target form avoids the solver needing U512 division.)
+    share_target: String,
+    /// Expected number of hashes to find a share (= share difficulty), FYI/UX.
+    expected_hashes: u64,
+    expires_in_secs: u64,
+}
+
+#[derive(Deserialize)]
+struct ShareRequest {
+    session_id: String,
+    /// Hex U512 nonce, no 0x prefix.
+    nonce: String,
+}
+
+#[derive(Serialize)]
+struct ShareResponse {
+    success: bool,
+    #[serde(skip_serializing_if = "Option::is_none")]
+    token: Option<String>,
+    #[serde(skip_serializing_if = "Option::is_none")]
+    error: Option<&'static str>,
+    block_found: bool,
+}
+
+/// Request shape mirrors reCAPTCHA/Turnstile: `secret` + `response`.
+#[derive(Deserialize)]
+struct SiteVerifyRequest {
+    secret: String,
+    response: String,
+}
+
+#[derive(Serialize)]
+struct SiteVerifyResponse {
+    success: bool,
+    #[serde(rename = "challenge_ts", skip_serializing_if = "Option::is_none")]
+    challenge_ts: Option<u64>,
+    #[serde(rename = "error-codes", skip_serializing_if = "Vec::is_empty")]
+    error_codes: Vec<&'static str>,
+}
+
+pub async fn serve(
+    state: Arc<PoolState>,
+    limiter: Arc<SessionIssuerLimiter>,
+    addr: SocketAddr,
+    serve_dir: Option<PathBuf>,
+) {
+    let with_state = {
+        let state = state.clone();
+        warp::any().map(move || state.clone())
+    };
+    let with_limiter = {
+        let limiter = limiter.clone();
+        warp::any().map(move || limiter.clone())
+    };
+
+    // POST /api/session -> new captcha challenge
+    let session = warp::path!("api" / "session")
+        .and(warp::post())
+        .and(warp::addr::remote())
+        .and(with_limiter.clone())
+        .and(with_state.clone())
+        .map(
+            |remote: Option<SocketAddr>,
+             limiter: Arc<SessionIssuerLimiter>,
+             state: Arc<PoolState>| {
+                let ip = remote
+                    .map(|a| a.ip())
+                    .unwrap_or(IpAddr::from([127, 0, 0, 1]));
+                if let Err(reason) = limiter.check(ip) {
+                    return reply_error(warp::http::StatusCode::TOO_MANY_REQUESTS, reason);
+                }
+                match state.issue_session() {
+                    Err(reason) => reply_error(warp::http::StatusCode::SERVICE_UNAVAILABLE, reason),
+                    Ok(s) => {
+                        let target = U512::MAX / s.share_difficulty;
+                        let expected = s.share_difficulty.min(U512::from(u64::MAX)).as_u64();
+                        warp::reply::with_status(
+                            warp::reply::json(&SessionResponse {
+                                session_id: s.id,
+                                header_hash: hex::encode(s.job.header),
+                                nonce_start: format!("{:x}", s.nonce_start),
+                                nonce_end: format!("{:x}", s.nonce_end),
+                                share_target: format!("{:x}", target),
+                                expected_hashes: expected,
+                                expires_in_secs: crate::state::SESSION_TTL.as_secs(),
+                            }),
+                            warp::http::StatusCode::OK,
+                        )
+                    }
+                }
+            },
+        );
+
+    // POST /api/share -> verify a solved nonce, mint a token
+    let share = warp::path!("api" / "share")
+        .and(warp::post())
+        .and(warp::body::content_length_limit(4096))
+        .and(warp::body::json())
+        .and(with_state.clone())
+        .map(|req: ShareRequest, state: Arc<PoolState>| {
+            let nonce = match U512::from_str_radix(req.nonce.trim_start_matches("0x"), 16) {
+                Ok(n) => n,
+                Err(_) => {
+                    return warp::reply::with_status(
+                        warp::reply::json(&ShareResponse {
+                            success: false,
+                            token: None,
+                            error: Some("malformed_nonce"),
+                            block_found: false,
+                        }),
+                        warp::http::StatusCode::BAD_REQUEST,
+                    )
+                }
+            };
+            match state.submit_share(&req.session_id, nonce) {
+                ShareOutcome::Accepted { token, block_found } => warp::reply::with_status(
+                    warp::reply::json(&ShareResponse {
+                        success: true,
+                        token: Some(token),
+                        error: None,
+                        block_found,
+                    }),
+                    warp::http::StatusCode::OK,
+                ),
+                ShareOutcome::Rejected(reason) => warp::reply::with_status(
+                    warp::reply::json(&ShareResponse {
+                        success: false,
+                        token: None,
+                        error: Some(reason),
+                        block_found: false,
+                    }),
+                    warp::http::StatusCode::FORBIDDEN,
+                ),
+            }
+        });
+
+    // POST /siteverify -> protected site's backend redeems a token
+    let siteverify = warp::path!("siteverify")
+        .and(warp::post())
+        .and(warp::body::content_length_limit(4096))
+        // Accept both JSON and form encoding, like the incumbents do.
+        .and(
+            warp::body::json()
+                .or(warp::body::form())
+                .unify()
+                .map(|req: SiteVerifyRequest| req),
+        )
+        .and(with_state.clone())
+        .map(|req: SiteVerifyRequest, state: Arc<PoolState>| {
+            match state.verify_token(&req.secret, &req.response) {
+                Ok(ts) => warp::reply::json(&SiteVerifyResponse {
+                    success: true,
+                    challenge_ts: Some(ts),
+                    error_codes: vec![],
+                }),
+                Err(code) => warp::reply::json(&SiteVerifyResponse {
+                    success: false,
+                    challenge_ts: None,
+                    error_codes: vec![code],
+                }),
+            }
+        });
+
+    // GET /api/stats -> pool counters
+    let stats = warp::path!("api" / "stats")
+        .and(warp::get())
+        .and(with_state.clone())
+        .map(|state: Arc<PoolState>| warp::reply::json(&state.stats()));
+
+    let cors = warp::cors()
+        .allow_any_origin()
+        .allow_methods(vec!["GET", "POST", "OPTIONS"])
+        .allow_headers(vec!["content-type"]);
+
+    let api = session.or(share).or(siteverify).or(stats);
+
+    log::info!("HTTP API listening on {}", addr);
+    if let Some(dir) = serve_dir {
+        log::info!("Serving demo statics from {}", dir.display());
+        let statics = warp::get().and(warp::fs::dir(dir));
+        warp::serve(api.or(statics).with(cors)).run(addr).await;
+    } else {
+        warp::serve(api.with(cors)).run(addr).await;
+    }
+}
+
+fn reply_error(
+    status: warp::http::StatusCode,
+    error: &'static str,
+) -> warp::reply::WithStatus<warp::reply::Json> {
+    #[derive(Serialize)]
+    struct Err<'a> {
+        error: &'a str,
+    }
+    warp::reply::with_status(warp::reply::json(&Err { error }), status)
+}
@@ -0,0 +1,250 @@
+//! Quantus captcha share pool.
+//!
+//! Sits between a quantus-node (external miner protocol) and browser captcha
+//! solvers: hands out low-difficulty share challenges over the real block
+//! header, verifies solves, mints single-use tokens for site backends, and
+//! submits any share that happens to meet full network difficulty as a block.
+
+mod http;
+mod rate_limit;
+mod state;
+mod upstream;
+
+use std::net::SocketAddr;
+use std::path::PathBuf;
+use std::sync::Arc;
+use std::time::Duration;
+
+use clap::Parser;
+use primitive_types::U512;
+
+use crate::rate_limit::SessionRateLimit;
+
+#[derive(Parser, Debug)]
+#[command(name = "pool-service", about = "Quantus captcha share pool")]
+struct Args {
+    /// HTTP listen address for the captcha/siteverify API.
+    #[arg(long, default_value = "127.0.0.1:8787", env = "POOL_HTTP_ADDR")]
+    http_addr: SocketAddr,
+
+    /// Address of the quantus-node external-miner QUIC endpoint.
+    /// If omitted, runs in standalone mode with synthetic jobs (demo only).
+    #[arg(long, env = "POOL_NODE_ADDR")]
+    node_addr: Option<SocketAddr>,
+
+    /// Shared auth token from the node's `miner-auth-token` file.
+    /// Required when `--node-addr` is set. Prefer `--auth-token-file`.
+    #[arg(long, env = "POOL_AUTH_TOKEN", conflicts_with = "auth_token_file")]
+    auth_token: Option<String>,
+
+    /// Path to the node's `miner-auth-token` file (trimmed).
+    #[arg(long, env = "POOL_AUTH_TOKEN_FILE", conflicts_with = "auth_token")]
+    auth_token_file: Option<PathBuf>,
+
+    /// SHA-256 fingerprint of the node's miner TLS certificate (64 hex chars).
+    /// Required when `--node-addr` is set. Prefer `--tls-cert-sha256-file`.
+    #[arg(
+        long,
+        env = "POOL_TLS_CERT_SHA256",
+        conflicts_with = "tls_cert_sha256_file"
+    )]
+    tls_cert_sha256: Option<String>,
+
+    /// Path to the node's `miner-tls-cert-sha256` file.
+    #[arg(
+        long,
+        env = "POOL_TLS_CERT_SHA256_FILE",
+        conflicts_with = "tls_cert_sha256"
+    )]
+    tls_cert_sha256_file: Option<PathBuf>,
+
+    /// Share difficulty: expected number of hashes per captcha solve.
+    /// Measured browser WASM rate ≈ 120 kH/s on an M-series laptop, so
+    /// 50000 ≈ 0.4 s desktop / ~2 s phone. Raise for stronger rate limiting.
+    #[arg(long, default_value = "50000", env = "POOL_SHARE_DIFFICULTY")]
+    share_difficulty: u64,
+
+    /// Secret that protected-site backends must present to /siteverify.
+    #[arg(long, default_value = "dev-secret", env = "POOL_SITE_SECRET")]
+    site_secret: String,
+
+    /// Optional directory of static files to serve (demo page / widget).
+    #[arg(long, env = "POOL_SERVE_DIR")]
+    serve_dir: Option<PathBuf>,
+
+    /// Job rotation interval in standalone mode, seconds.
+    #[arg(long, default_value = "20")]
+    standalone_job_secs: u64,
+
+    /// Maximum live captcha sessions; /api/session returns 503 when full.
+    #[arg(long, default_value = "100000", env = "POOL_MAX_SESSIONS")]
+    max_sessions: usize,
+
+    /// Maximum live (unredeemed) share tokens; shares are refused when full.
+    #[arg(long, default_value = "100000", env = "POOL_MAX_TOKENS")]
+    max_tokens: usize,
+
+    /// Max /api/session issuances per client IP per minute.
+    #[arg(long, default_value = "60", env = "POOL_SESSIONS_PER_IP_PER_MIN")]
+    sessions_per_ip_per_min: u32,
+}
+
+#[tokio::main]
+async fn main() -> anyhow::Result<()> {
+    env_logger::Builder::from_env(env_logger::Env::default().default_filter_or("info")).init();
+    let args = Args::parse();
+
+    anyhow::ensure!(args.share_difficulty > 0, "--share-difficulty must be > 0");
+    if args.site_secret == "dev-secret" {
+        log::warn!("Using default site secret; set --site-secret in production");
+    }
+
+    // Resolve + validate node auth before starting HTTP / upstream tasks so a
+    // bad fingerprint or oversized token exits instead of retrying forever.
+    let node_upstream = match args.node_addr {
+        Some(addr) => {
+            let auth_token = resolve_auth_token(args.auth_token, args.auth_token_file)?;
+            let tls_cert_sha256 =
+                resolve_tls_cert_sha256(args.tls_cert_sha256, args.tls_cert_sha256_file)?;
+            quic_transport::validate_auth_config(&auth_token, &tls_cert_sha256)?;
+            Some((addr, auth_token, tls_cert_sha256))
+        }
+        None => {
+            if args.auth_token.is_some()
+                || args.auth_token_file.is_some()
+                || args.tls_cert_sha256.is_some()
+                || args.tls_cert_sha256_file.is_some()
+            {
+                log::warn!("Ignoring auth/TLS pin flags in standalone mode (no --node-addr)");
+            }
+            None
+        }
+    };
+
+    let (solution_tx, solution_rx) = tokio::sync::mpsc::channel(16);
+    let state = state::PoolState::new(
+        U512::from(args.share_difficulty),
+        args.site_secret.clone(),
+        solution_tx,
+        state::Limits {
+            max_sessions: args.max_sessions,
+            max_tokens: args.max_tokens,
+        },
+    );
+
+    let limiter = Arc::new(rate_limit::SessionIssuerLimiter::new(SessionRateLimit {
+        max_per_ip: args.sessions_per_ip_per_min,
+        window: Duration::from_secs(60),
+    }));
+
+    // Periodic cleanup of expired sessions/tokens and stale rate-limit windows.
+    {
+        let state = state.clone();
+        let limiter = limiter.clone();
+        tokio::spawn(async move {
+            loop {
+                tokio::time::sleep(Duration::from_secs(30)).await;
+                state.gc();
+                limiter.gc();
+            }
+        });
+    }
+
+    // Job source. Node mode reconnects across transient disconnects (node
+    // restart); only a PermanentConnectError ends the task — then we tear
+    // down HTTP so the process does not look healthy on a bad token/pin.
+    match node_upstream {
+        Some((addr, auth_token, tls_cert_sha256)) => {
+            log::info!("Upstream: quantus-node at {}", addr);
+            let state_for_upstream = state.clone();
+            let upstream = tokio::spawn(async move {
+                upstream::run_node_client(
+                    state_for_upstream,
+                    addr,
+                    auth_token,
+                    tls_cert_sha256,
+                    solution_rx,
+                )
+                .await
+            });
+
+            tokio::select! {
+                biased;
+                result = upstream => {
+                    match result {
+                        Ok(Ok(())) => anyhow::bail!(
+                            "upstream task exited unexpectedly (should reconnect forever)"
+                        ),
+                        Ok(Err(e)) => {
+                            log::error!("Upstream permanent failure; shutting down: {e}");
+                            Err(e.into())
+                        }
+                        Err(e) => Err(anyhow::anyhow!("upstream task panicked: {e}")),
+                    }
+                }
+                _ = http::serve(state, limiter, args.http_addr, args.serve_dir) => Ok(()),
+            }
+        }
+        None => {
+            log::warn!("No --node-addr given: running STANDALONE with synthetic jobs");
+            let state_for_upstream = state.clone();
+            tokio::spawn(async move {
+                upstream::run_standalone(
+                    state_for_upstream,
+                    Duration::from_secs(args.standalone_job_secs),
+                    solution_rx,
+                )
+                .await
+            });
+            http::serve(state, limiter, args.http_addr, args.serve_dir).await;
+            Ok(())
+        }
+    }
+}
+
+fn resolve_auth_token(
+    auth_token: Option<String>,
+    auth_token_file: Option<PathBuf>,
+) -> anyhow::Result<String> {
+    resolve_required_secret(
+        auth_token,
+        auth_token_file,
+        "--auth-token",
+        "--auth-token-file",
+        "when --node-addr is set, pass --auth-token-file <PATH> to the node's \
+         miner-auth-token file (or --auth-token <TOKEN>)",
+    )
+}
+
+fn resolve_tls_cert_sha256(value: Option<String>, file: Option<PathBuf>) -> anyhow::Result<String> {
+    resolve_required_secret(
+        value,
+        file,
+        "--tls-cert-sha256",
+        "--tls-cert-sha256-file",
+        "when --node-addr is set, pass --tls-cert-sha256-file <PATH> to the node's \
+         miner-tls-cert-sha256 file (or --tls-cert-sha256 <HEX>; also printed in node logs)",
+    )
+}
+
+fn resolve_required_secret(
+    value: Option<String>,
+    file: Option<PathBuf>,
+    value_flag: &str,
+    file_flag: &str,
+    missing_msg: &str,
+) -> anyhow::Result<String> {
+    if let Some(value) = value {
+        let value = value.trim().to_string();
+        anyhow::ensure!(!value.is_empty(), "{value_flag} is empty");
+        return Ok(value);
+    }
+    if let Some(path) = file {
+        let contents = std::fs::read_to_string(&path)
+            .map_err(|e| anyhow::anyhow!("failed to read {file_flag} {}: {}", path.display(), e))?;
+        let value = contents.trim().to_string();
+        anyhow::ensure!(!value.is_empty(), "{file_flag} {} is empty", path.display());
+        return Ok(value);
+    }
+    anyhow::bail!("{missing_msg}");
+}
@@ -0,0 +1,104 @@
+//! Per-client rate limiting for unauthenticated endpoints.
+//!
+//! `/api/session` is free, so a hard map cap alone still lets an attacker
+//! churn entries at the TTL boundary. A per-IP issuance ceiling bounds that
+//! even when the global cap has headroom.
+
+use std::collections::HashMap;
+use std::net::IpAddr;
+use std::sync::Mutex;
+use std::time::{Duration, Instant};
+
+#[derive(Debug, Clone, Copy)]
+pub struct SessionRateLimit {
+    /// Max session issuances per IP per window.
+    pub max_per_ip: u32,
+    pub window: Duration,
+}
+
+impl Default for SessionRateLimit {
+    fn default() -> Self {
+        SessionRateLimit {
+            // 60 sessions/min/IP is generous for real users (a page load is
+            // one) but stops a single host from dominating issuance.
+            max_per_ip: 60,
+            window: Duration::from_secs(60),
+        }
+    }
+}
+
+#[derive(Default)]
+struct IpWindow {
+    count: u32,
+    window_start: Option<Instant>,
+}
+
+pub struct SessionIssuerLimiter {
+    cfg: SessionRateLimit,
+    by_ip: Mutex<HashMap<IpAddr, IpWindow>>,
+}
+
+impl SessionIssuerLimiter {
+    pub fn new(cfg: SessionRateLimit) -> Self {
+        SessionIssuerLimiter {
+            cfg,
+            by_ip: Mutex::new(HashMap::new()),
+        }
+    }
+
+    /// Returns Ok(()) if this IP may issue another session, Err otherwise.
+    pub fn check(&self, ip: IpAddr) -> Result<(), &'static str> {
+        let now = Instant::now();
+        let mut map = self.by_ip.lock().unwrap();
+        let entry = map.entry(ip).or_default();
+        let window_start = entry.window_start.get_or_insert_with(Instant::now);
+        if now.duration_since(*window_start) >= self.cfg.window {
+            *window_start = now;
+            entry.count = 0;
+        }
+        if entry.count >= self.cfg.max_per_ip {
+            return Err("rate_limited");
+        }
+        entry.count += 1;
+        Ok(())
+    }
+
+    /// Drop stale IP windows so the side table doesn't grow forever.
+    pub fn gc(&self) {
+        let cutoff = self.cfg.window * 2;
+        let now = Instant::now();
+        self.by_ip.lock().unwrap().retain(|_, w| {
+            w.window_start
+                .map(|t| now.duration_since(t) < cutoff)
+                .unwrap_or(true)
+        });
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use std::net::{Ipv4Addr, Ipv6Addr};
+
+    #[test]
+    fn allows_up_to_cap_then_rejects() {
+        let lim = SessionIssuerLimiter::new(SessionRateLimit {
+            max_per_ip: 2,
+            window: Duration::from_secs(60),
+        });
+        let ip = IpAddr::V4(Ipv4Addr::new(10, 0, 0, 1));
+        assert!(lim.check(ip).is_ok());
+        assert!(lim.check(ip).is_ok());
+        assert_eq!(lim.check(ip), Err("rate_limited"));
+    }
+
+    #[test]
+    fn ips_are_independent() {
+        let lim = SessionIssuerLimiter::new(SessionRateLimit {
+            max_per_ip: 1,
+            window: Duration::from_secs(60),
+        });
+        assert!(lim.check(IpAddr::V4(Ipv4Addr::new(1, 2, 3, 4))).is_ok());
+        assert!(lim.check(IpAddr::V6(Ipv6Addr::LOCALHOST)).is_ok());
+    }
+}
@@ -0,0 +1,527 @@
+//! Core pool state: the current mining job, captcha sessions with disjoint
+//! nonce ranges, and single-use share tokens.
+
+use std::collections::HashMap;
+use std::sync::atomic::{AtomicU64, Ordering};
+use std::sync::{Arc, Mutex};
+use std::time::{Duration, Instant, SystemTime, UNIX_EPOCH};
+
+use primitive_types::U512;
+use rand::RngCore;
+
+/// How long a captcha session stays solvable after issuance.
+pub const SESSION_TTL: Duration = Duration::from_secs(120);
+/// How long a share token stays redeemable via /siteverify.
+pub const TOKEN_TTL: Duration = Duration::from_secs(300);
+/// Each session gets a disjoint slice of the nonce space this many nonces wide.
+/// 2^64 nonces is unreachable within a session TTL at browser hash rates.
+pub const SESSION_RANGE_BITS: u32 = 64;
+
+/// Caps on live entries in the in-memory maps. `/api/session` is free and
+/// unauthenticated, so without a ceiling a request flood becomes a memory-DoS
+/// on the anti-abuse service itself. At the ~300 bytes/entry ballpark the
+/// defaults bound each map to tens of MB.
+#[derive(Debug, Clone, Copy)]
+pub struct Limits {
+    pub max_sessions: usize,
+    pub max_tokens: usize,
+}
+
+impl Default for Limits {
+    fn default() -> Self {
+        Limits {
+            max_sessions: 100_000,
+            max_tokens: 100_000,
+        }
+    }
+}
+
+/// The job the pool is currently handing out to captcha sessions.
+#[derive(Debug, Clone)]
+pub struct Job {
+    /// Job id assigned by the upstream node (or synthesized in standalone mode).
+    pub job_id: String,
+    /// Block header hash to mine over.
+    pub header: [u8; 32],
+    /// Full network difficulty for this job (block found if a share meets it).
+    pub network_difficulty: U512,
+}
+
+/// A captcha session: one issued challenge, bound to a disjoint nonce range.
+#[derive(Debug, Clone)]
+pub struct Session {
+    pub id: String,
+    /// Snapshot of the job at issuance time. Shares are verified against this
+    /// header even if the pool has since moved to a newer job (grace period =
+    /// session TTL); only network-difficulty solutions for the *current* job
+    /// are submitted upstream.
+    pub job: Job,
+    pub nonce_start: U512,
+    pub nonce_end: U512, // exclusive
+    pub share_difficulty: U512,
+    pub issued_at: Instant,
+    pub solved: bool,
+}
+
+/// A verified share, redeemable exactly once by the protected site's backend.
+#[derive(Debug, Clone)]
+pub struct ShareToken {
+    pub created_at: Instant,
+    pub created_unix: u64,
+    pub consumed: bool,
+}
+
+#[derive(Debug, Default, Clone, serde::Serialize)]
+pub struct PoolStats {
+    pub sessions_issued: u64,
+    pub shares_accepted: u64,
+    pub shares_rejected: u64,
+    pub tokens_verified: u64,
+    pub blocks_found: u64,
+}
+
+pub struct PoolState {
+    /// Current job, None until the first job arrives from the job source.
+    current_job: Mutex<Option<Job>>,
+    sessions: Mutex<HashMap<String, Session>>,
+    tokens: Mutex<HashMap<String, ShareToken>>,
+    stats: Mutex<PoolStats>,
+    /// Monotonic counter used to carve disjoint nonce ranges out of U512 space.
+    range_counter: AtomicU64,
+    /// Difficulty for captcha shares (expected hashes per solve).
+    pub share_difficulty: U512,
+    /// Site secret checked by /siteverify.
+    ///
+    /// NOTE: single-tenant by design for now. All tokens are interchangeable
+    /// across any backend holding this secret; per-sitekey secrets land with
+    /// the host-registration follow-up.
+    pub site_secret: String,
+    limits: Limits,
+    /// Sender used to push network-difficulty solutions upstream.
+    pub solution_tx: tokio::sync::mpsc::Sender<FoundBlock>,
+}
+
+/// A share that met full network difficulty: a real block solution.
+#[derive(Debug, Clone)]
+pub struct FoundBlock {
+    pub job_id: String,
+    pub nonce: U512,
+}
+
+pub enum ShareOutcome {
+    /// Share accepted; token the client can hand to the protected site.
+    Accepted {
+        token: String,
+        block_found: bool,
+    },
+    Rejected(&'static str),
+}
+
+fn random_id() -> String {
+    let mut bytes = [0u8; 24];
+    rand::rng().fill_bytes(&mut bytes);
+    hex::encode(bytes)
+}
+
+fn unix_now() -> u64 {
+    SystemTime::now()
+        .duration_since(UNIX_EPOCH)
+        .map(|d| d.as_secs())
+        .unwrap_or(0)
+}
+
+impl PoolState {
+    pub fn new(
+        share_difficulty: U512,
+        site_secret: String,
+        solution_tx: tokio::sync::mpsc::Sender<FoundBlock>,
+        limits: Limits,
+    ) -> Arc<Self> {
+        assert!(
+            !share_difficulty.is_zero(),
+            "share difficulty must be non-zero"
+        );
+        Arc::new(PoolState {
+            current_job: Mutex::new(None),
+            sessions: Mutex::new(HashMap::new()),
+            tokens: Mutex::new(HashMap::new()),
+            stats: Mutex::new(PoolStats::default()),
+            range_counter: AtomicU64::new(0),
+            share_difficulty,
+            site_secret,
+            limits,
+            solution_tx,
+        })
+    }
+
+    /// Install a new job (from the upstream node or the standalone generator).
+    /// Existing sessions keep their job snapshot and remain solvable until TTL.
+    pub fn set_job(&self, job: Job) {
+        log::info!(
+            "New job {}: header=0x{} network_difficulty={}",
+            job.job_id,
+            hex::encode(job.header),
+            pow_core::format_u512(job.network_difficulty)
+        );
+        *self.current_job.lock().unwrap() = Some(job);
+    }
+
+    pub fn current_job(&self) -> Option<Job> {
+        self.current_job.lock().unwrap().clone()
+    }
+
+    /// Drop the current job (e.g. upstream disconnected). New sessions will
+    /// see `no_job_available` until the next `set_job`.
+    pub fn clear_job(&self) {
+        if self.current_job.lock().unwrap().take().is_some() {
+            log::info!("Cleared current job (upstream unavailable)");
+        }
+    }
+
+    /// Issue a new captcha session with a disjoint nonce range.
+    pub fn issue_session(&self) -> Result<Session, &'static str> {
+        let job = self.current_job().ok_or("no_job_available")?;
+
+        if self.sessions.lock().unwrap().len() >= self.limits.max_sessions {
+            return Err("at_capacity");
+        }
+
+        // Carve the next disjoint 2^64-nonce slice. The counter is 64 bits and
+        // slices are 2^64 wide, so slices live in the bottom 2^128 of the U512
+        // nonce space and never collide with each other.
+        let index = self.range_counter.fetch_add(1, Ordering::Relaxed);
+        let nonce_start = U512::from(index) << SESSION_RANGE_BITS;
+        let nonce_end = nonce_start + (U512::one() << SESSION_RANGE_BITS);
+
+        let session = Session {
+            id: random_id(),
+            job,
+            nonce_start,
+            nonce_end,
+            share_difficulty: self.share_difficulty,
+            issued_at: Instant::now(),
+            solved: false,
+        };
+
+        self.sessions
+            .lock()
+            .unwrap()
+            .insert(session.id.clone(), session.clone());
+        self.stats.lock().unwrap().sessions_issued += 1;
+        Ok(session)
+    }
+
+    /// Verify a submitted share and, if valid, mint a single-use token.
+    pub fn submit_share(&self, session_id: &str, nonce: U512) -> ShareOutcome {
+        let session = {
+            let mut sessions = self.sessions.lock().unwrap();
+            match sessions.get_mut(session_id) {
+                None => return self.reject("unknown_session"),
+                Some(s) => {
+                    if s.solved {
+                        return self.reject("session_already_solved");
+                    }
+                    if s.issued_at.elapsed() > SESSION_TTL {
+                        sessions.remove(session_id);
+                        return self.reject("session_expired");
+                    }
+                    // Mark solved optimistically; unmark on verification failure
+                    // would allow retries, but a failed share means a buggy or
+                    // malicious client, so we burn the session either way.
+                    s.solved = true;
+                    s.clone()
+                }
+            }
+        };
+
+        if nonce < session.nonce_start || nonce >= session.nonce_end {
+            return self.reject("nonce_out_of_range");
+        }
+
+        let (share_valid, hash) = pow_core::is_valid_nonce(
+            session.job.header,
+            nonce.to_big_endian(),
+            session.share_difficulty,
+        );
+        if !share_valid {
+            return self.reject("share_below_target");
+        }
+
+        // Jackpot check: does this share meet full network difficulty for the
+        // job it was issued against? Only submit if that job is still current.
+        // Target derivation delegates to pow-core's canonical JobContext.
+        let network_target =
+            pow_core::JobContext::new(session.job.header, session.job.network_difficulty).target;
+        let mut block_found = false;
+        if hash < network_target {
+            let still_current = self
+                .current_job()
+                .map(|j| j.job_id == session.job.job_id)
+                .unwrap_or(false);
+            if still_current {
+                log::info!(
+                    "BLOCK FOUND by captcha share! job={} nonce={:x}",
+                    session.job.job_id,
+                    nonce
+                );
+                // This is the revenue event: count it and be loud if the
+                // upstream queue can't take it.
+                match self.solution_tx.try_send(FoundBlock {
+                    job_id: session.job.job_id.clone(),
+                    nonce,
+                }) {
+                    Ok(()) => {
+                        block_found = true;
+                        self.stats.lock().unwrap().blocks_found += 1;
+                    }
+                    Err(e) => {
+                        log::error!(
+                            "DROPPED block solution for job {}: upstream queue unavailable ({})",
+                            session.job.job_id,
+                            e
+                        );
+                    }
+                }
+            }
+        }
+
+        let token = random_id();
+        {
+            let mut tokens = self.tokens.lock().unwrap();
+            // A full token map means someone is farming shares faster than
+            // TTL expiry; refuse rather than grow without bound. The client
+            // paid hashes for this, so make the refusal explicit.
+            if tokens.len() >= self.limits.max_tokens {
+                return self.reject("at_capacity");
+            }
+            tokens.insert(
+                token.clone(),
+                ShareToken {
+                    created_at: Instant::now(),
+                    created_unix: unix_now(),
+                    consumed: false,
+                },
+            );
+        }
+        self.stats.lock().unwrap().shares_accepted += 1;
+        ShareOutcome::Accepted { token, block_found }
+    }
+
+    /// Redeem a share token (the reCAPTCHA/Turnstile "siteverify" semantics):
+    /// valid exactly once, only with the correct site secret.
+    pub fn verify_token(&self, secret: &str, token: &str) -> Result<u64, &'static str> {
+        if !constant_time_eq::constant_time_eq(secret.as_bytes(), self.site_secret.as_bytes()) {
+            return Err("invalid_secret");
+        }
+        let mut tokens = self.tokens.lock().unwrap();
+        match tokens.get_mut(token) {
+            None => Err("invalid_token"),
+            Some(t) if t.consumed => Err("token_already_consumed"),
+            Some(t) if t.created_at.elapsed() > TOKEN_TTL => {
+                tokens.remove(token);
+                Err("token_expired")
+            }
+            Some(t) => {
+                t.consumed = true;
+                let ts = t.created_unix;
+                self.stats.lock().unwrap().tokens_verified += 1;
+                Ok(ts)
+            }
+        }
+    }
+
+    pub fn stats(&self) -> PoolStats {
+        self.stats.lock().unwrap().clone()
+    }
+
+    /// Drop expired sessions and tokens. Called periodically.
+    pub fn gc(&self) {
+        self.sessions
+            .lock()
+            .unwrap()
+            .retain(|_, s| s.issued_at.elapsed() <= SESSION_TTL);
+        self.tokens
+            .lock()
+            .unwrap()
+            .retain(|_, t| t.created_at.elapsed() <= TOKEN_TTL);
+    }
+
+    fn reject(&self, reason: &'static str) -> ShareOutcome {
+        self.stats.lock().unwrap().shares_rejected += 1;
+        ShareOutcome::Rejected(reason)
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn test_state(
+        share_difficulty: u64,
+    ) -> (Arc<PoolState>, tokio::sync::mpsc::Receiver<FoundBlock>) {
+        let (tx, rx) = tokio::sync::mpsc::channel(4);
+        let state = PoolState::new(
+            U512::from(share_difficulty),
+            "secret".into(),
+            tx,
+            Limits::default(),
+        );
+        state.set_job(Job {
+            job_id: "1".into(),
+            header: [7u8; 32],
+            network_difficulty: U512::MAX, // effectively unreachable
+        });
+        (state, rx)
+    }
+
+    /// Brute-force a valid share within the session's range.
+    fn solve(session: &Session) -> U512 {
+        let mut nonce = session.nonce_start;
+        loop {
+            let (ok, _) = pow_core::is_valid_nonce(
+                session.job.header,
+                nonce.to_big_endian(),
+                session.share_difficulty,
+            );
+            if ok {
+                return nonce;
+            }
+            nonce += U512::one();
+        }
+    }
+
+    #[test]
+    fn accepts_valid_share_and_verifies_token_once() {
+        let (state, _rx) = test_state(2); // ~2 hashes per solve
+        let session = state.issue_session().unwrap();
+        let nonce = solve(&session);
+
+        let token = match state.submit_share(&session.id, nonce) {
+            ShareOutcome::Accepted { token, block_found } => {
+                assert!(!block_found);
+                token
+            }
+            ShareOutcome::Rejected(r) => panic!("rejected: {r}"),
+        };
+
+        assert!(state.verify_token("secret", &token).is_ok());
+        assert_eq!(
+            state.verify_token("secret", &token),
+            Err("token_already_consumed")
+        );
+        assert_eq!(state.verify_token("wrong", &token), Err("invalid_secret"));
+    }
+
+    #[test]
+    fn rejects_out_of_range_nonce() {
+        let (state, _rx) = test_state(1); // difficulty 1: every nonce is a valid hash
+        let session = state.issue_session().unwrap();
+        // A nonce outside the assigned slice must be rejected even though its
+        // hash trivially meets the share target.
+        let outside = session.nonce_end;
+        match state.submit_share(&session.id, outside) {
+            ShareOutcome::Rejected("nonce_out_of_range") => {}
+            other => panic!(
+                "expected out-of-range rejection, got {:?}",
+                matches_str(&other)
+            ),
+        }
+    }
+
+    #[test]
+    fn burns_session_after_one_submission() {
+        let (state, _rx) = test_state(1);
+        let session = state.issue_session().unwrap();
+        let nonce = session.nonce_start;
+        assert!(matches!(
+            state.submit_share(&session.id, nonce),
+            ShareOutcome::Accepted { .. }
+        ));
+        match state.submit_share(&session.id, nonce) {
+            ShareOutcome::Rejected("session_already_solved") => {}
+            other => panic!("expected burned session, got {:?}", matches_str(&other)),
+        }
+    }
+
+    #[test]
+    fn disjoint_ranges() {
+        let (state, _rx) = test_state(1);
+        let a = state.issue_session().unwrap();
+        let b = state.issue_session().unwrap();
+        assert!(a.nonce_end <= b.nonce_start || b.nonce_end <= a.nonce_start);
+    }
+
+    #[test]
+    fn block_found_signalled_when_share_meets_network_difficulty() {
+        let (tx, mut rx) = tokio::sync::mpsc::channel(4);
+        // Network difficulty == share difficulty == 2: any valid share is a block.
+        let state = PoolState::new(U512::from(2u64), "secret".into(), tx, Limits::default());
+        state.set_job(Job {
+            job_id: "1".into(),
+            header: [9u8; 32],
+            network_difficulty: U512::from(2u64),
+        });
+        let session = state.issue_session().unwrap();
+        let nonce = solve(&session);
+        match state.submit_share(&session.id, nonce) {
+            ShareOutcome::Accepted { block_found, .. } => assert!(block_found),
+            ShareOutcome::Rejected(r) => panic!("rejected: {r}"),
+        }
+        let block = rx
+            .try_recv()
+            .expect("block solution should be queued upstream");
+        assert_eq!(block.job_id, "1");
+        assert_eq!(block.nonce, nonce);
+    }
+
+    #[test]
+    fn session_cap_rejects_flood() {
+        let (tx, _rx) = tokio::sync::mpsc::channel(4);
+        let state = PoolState::new(
+            U512::from(1u64),
+            "secret".into(),
+            tx,
+            Limits {
+                max_sessions: 2,
+                max_tokens: 2,
+            },
+        );
+        state.set_job(Job {
+            job_id: "1".into(),
+            header: [7u8; 32],
+            network_difficulty: U512::MAX,
+        });
+        assert!(state.issue_session().is_ok());
+        assert!(state.issue_session().is_ok());
+        assert!(matches!(state.issue_session(), Err("at_capacity")));
+    }
+
+    #[test]
+    fn dropped_block_does_not_count_or_flag() {
+        // Zero-capacity channel: try_send always fails, simulating a stalled
+        // upstream. The share must still be accepted (the user did the work),
+        // but block_found must be false and blocks_found must stay 0.
+        let (tx, rx) = tokio::sync::mpsc::channel(1);
+        drop(rx); // closed channel -> try_send errors
+        let state = PoolState::new(U512::from(2u64), "secret".into(), tx, Limits::default());
+        state.set_job(Job {
+            job_id: "1".into(),
+            header: [9u8; 32],
+            network_difficulty: U512::from(2u64),
+        });
+        let session = state.issue_session().unwrap();
+        let nonce = solve(&session);
+        match state.submit_share(&session.id, nonce) {
+            ShareOutcome::Accepted { block_found, .. } => assert!(!block_found),
+            ShareOutcome::Rejected(r) => panic!("rejected: {r}"),
+        }
+        assert_eq!(state.stats().blocks_found, 0);
+    }
+
+    fn matches_str(o: &ShareOutcome) -> &'static str {
+        match o {
+            ShareOutcome::Accepted { .. } => "accepted",
+            ShareOutcome::Rejected(r) => r,
+        }
+    }
+}
@@ -0,0 +1,192 @@
+//! Job sources: either a QUIC connection to a real quantus-node (speaking the
+//! external miner protocol), or a standalone generator for demos/tests.
+
+use std::net::SocketAddr;
+use std::sync::Arc;
+use std::time::Duration;
+
+use primitive_types::U512;
+use quantus_miner_api::{
+    read_message, write_message, ApiResponseStatus, MinerMessage, MiningResult,
+};
+use tokio::sync::mpsc::Receiver;
+
+use crate::state::{FoundBlock, Job, PoolState};
+
+/// Connect to a node as an external miner and keep the pool's job current.
+/// Solutions arriving on `solutions` are pushed upstream as JobResults.
+///
+/// Transient disconnects (node restart, network blip) clear the current job
+/// and reconnect with backoff — the pool process stays up. Only
+/// [`quic_transport::PermanentConnectError`] (bad pin / rejected auth) returns
+/// `Err` so `main` can shut down the HTTP side.
+pub async fn run_node_client(
+    state: Arc<PoolState>,
+    node_addr: SocketAddr,
+    auth_token: String,
+    tls_cert_sha256: String,
+    mut solutions: Receiver<FoundBlock>,
+) -> Result<(), quic_transport::PermanentConnectError> {
+    let mut reconnect_delay = Duration::from_secs(1);
+    const MAX_RECONNECT_DELAY: Duration = Duration::from_secs(30);
+
+    // A block taken off the channel but not yet acknowledged by a successful
+    // upstream write. Kept across reconnects so a failed write doesn't drop
+    // the block (the node may still accept it after reconnection if the job
+    // hasn't moved on).
+    let mut pending: Option<FoundBlock> = None;
+
+    loop {
+        log::info!("Connecting to node at {}...", node_addr);
+        match quic_transport::connect(node_addr, &auth_token, &tls_cert_sha256).await {
+            Ok((connection, mut send, mut recv)) => {
+                log::info!("Connected to node at {}", node_addr);
+                let mut authenticated = false;
+
+                loop {
+                    // Retry any unsubmitted block before doing anything else.
+                    if let Some(block) = pending.take() {
+                        match submit_block(&mut send, &block).await {
+                            Ok(()) => {
+                                log::info!(
+                                    "Submitted block solution to node (job {})",
+                                    block.job_id
+                                );
+                            }
+                            Err(e) => {
+                                log::error!("Failed to submit block solution, will retry: {}", e);
+                                pending = Some(block);
+                                break; // reconnect and retry
+                            }
+                        }
+                    }
+
+                    tokio::select! {
+                        biased;
+
+                        reason = connection.closed() => {
+                            let msg = reason.to_string();
+                            if !authenticated && msg.to_ascii_lowercase().contains("auth") {
+                                state.clear_job();
+                                return Err(quic_transport::PermanentConnectError(format!(
+                                    "node rejected miner auth ({msg}); check --auth-token / miner-auth-token"
+                                )));
+                            }
+                            log::warn!("Node connection closed: {}", msg);
+                            break;
+                        }
+
+                        // A captcha share met network difficulty: submit it.
+                        block = solutions.recv() => {
+                            if let Some(block) = block {
+                                // Park it in `pending`; the top of the loop
+                                // submits it and keeps it on write failure.
+                                pending = Some(block);
+                            }
+                        }
+
+                        msg = read_message(&mut recv) => {
+                            match msg {
+                                Ok(MinerMessage::NewJob(request)) => {
+                                    authenticated = true;
+                                    match parse_job(&request.job_id, &request.mining_hash, &request.difficulty) {
+                                        Ok(job) => state.set_job(job),
+                                        Err(e) => log::warn!("Ignoring malformed job from node: {}", e),
+                                    }
+                                }
+                                Ok(other) => {
+                                    log::warn!("Unexpected message from node: {:?}", other);
+                                }
+                                Err(e) => {
+                                    log::warn!("Read error from node: {}", e);
+                                    break;
+                                }
+                            }
+                        }
+                    }
+                }
+
+                // Node gone / stream dropped: stop handing out captchas on a
+                // stale header until we reconnect and get a fresh NewJob.
+                state.clear_job();
+                if authenticated {
+                    reconnect_delay = Duration::from_secs(1);
+                }
+            }
+            Err(e) => {
+                if let Some(perm) = e.downcast_ref::<quic_transport::PermanentConnectError>() {
+                    state.clear_job();
+                    return Err(quic_transport::PermanentConnectError(perm.0.clone()));
+                }
+                log::warn!("Failed to connect to node: {}", e);
+                state.clear_job();
+            }
+        }
+
+        log::info!("Reconnecting in {:?}...", reconnect_delay);
+        tokio::time::sleep(reconnect_delay).await;
+        reconnect_delay = (reconnect_delay * 2).min(MAX_RECONNECT_DELAY);
+    }
+}
+
+async fn submit_block(send: &mut quinn::SendStream, block: &FoundBlock) -> std::io::Result<()> {
+    let result = MiningResult {
+        status: ApiResponseStatus::Completed,
+        job_id: block.job_id.clone(),
+        nonce: Some(format!("{:x}", block.nonce)),
+        work: Some(hex::encode(block.nonce.to_big_endian())),
+        hash_count: 0,
+        elapsed_time: 0.0,
+        miner_id: None,
+    };
+    write_message(send, &MinerMessage::JobResult(result)).await
+}
+
+fn parse_job(job_id: &str, mining_hash: &str, difficulty: &str) -> anyhow::Result<Job> {
+    let bytes = hex::decode(mining_hash)?;
+    let header: [u8; 32] = bytes
+        .try_into()
+        .map_err(|_| anyhow::anyhow!("mining_hash must be 32 bytes"))?;
+    let network_difficulty =
+        U512::from_dec_str(difficulty).map_err(|e| anyhow::anyhow!("bad difficulty: {:?}", e))?;
+    if network_difficulty.is_zero() {
+        anyhow::bail!("zero difficulty");
+    }
+    Ok(Job {
+        job_id: job_id.to_string(),
+        header,
+        network_difficulty,
+    })
+}
+
+/// Standalone mode: synthesize a fresh random job every `interval` so the
+/// captcha works without a running node. Solutions are logged and dropped.
+pub async fn run_standalone(
+    state: Arc<PoolState>,
+    interval: Duration,
+    mut solutions: Receiver<FoundBlock>,
+) {
+    use rand::RngCore;
+    let mut counter: u64 = 0;
+
+    loop {
+        let mut header = [0u8; 32];
+        rand::rng().fill_bytes(&mut header);
+        counter += 1;
+        state.set_job(Job {
+            job_id: format!("standalone-{}", counter),
+            header,
+            // Unreachable "network" difficulty: standalone mode never finds blocks.
+            network_difficulty: U512::MAX,
+        });
+
+        tokio::select! {
+            _ = tokio::time::sleep(interval) => {}
+            block = solutions.recv() => {
+                if let Some(block) = block {
+                    log::info!("(standalone) would submit block for job {}", block.job_id);
+                }
+            }
+        }
+    }
+}
@@ -34,6 +34,7 @@ gpu-interop = []

 [dependencies]
 primitive-types = { workspace = true }
+qp-poseidon-core = { workspace = true }
 qpow-math = { workspace = true }

 [dev-dependencies]
@@ -1,7 +1,28 @@
 use primitive_types::U512;
+use qp_poseidon_core::{Goldilocks, Poseidon2};

+pub use qp_poseidon_core::SPONGE_WIDTH;
 pub use qpow_math::{get_nonce_hash, is_valid_nonce, mine_range};

+/// Sponge state (canonical u64 felts) after absorbing the 32-byte header and the
+/// high 32 bytes of the big-endian nonce — the first two of the five Poseidon2
+/// permutations of `get_nonce_hash`. This state is identical for every nonce in a
+/// batch as long as incrementing the nonce never carries into its high 256 bits,
+/// so GPU kernels can resume the sponge from here and skip 2 of 5 permutations.
+pub fn mining_midstate(header: [u8; 32], nonce_high_be: [u8; 32]) -> [u64; SPONGE_WIDTH] {
+    let poseidon2 = Poseidon2::new();
+    let mut state = [Goldilocks::ZERO; SPONGE_WIDTH];
+    for (i, chunk) in header.chunks_exact(4).enumerate() {
+        state[i] += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64);
+    }
+    poseidon2.permute_mut(&mut state);
+    for (i, chunk) in nonce_high_be.chunks_exact(4).enumerate() {
+        state[i] += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64);
+    }
+    poseidon2.permute_mut(&mut state);
+    state.map(|g| g.as_canonical_u64())
+}
+
 /// Format a U512 in a human-readable way (scientific notation for large numbers).
 pub fn format_u512(n: U512) -> String {
    if n.is_zero() {
@@ -222,6 +243,39 @@ mod tests {
        assert!(result.is_none());
    }

+    #[test]
+    fn test_midstate_resumes_to_full_hash() {
+        let header = [7u8; 32];
+        let nonce = (U512::from(0xdeadbeefcafeu64) << 300) | U512::from(0x1234567890abcdefu64);
+        let nonce_be = nonce.to_big_endian();
+
+        let mut state =
+            mining_midstate(header, nonce_be[..32].try_into().unwrap()).map(Goldilocks::from_u64);
+        let poseidon2 = Poseidon2::new();
+        for (i, chunk) in nonce_be[32..].chunks_exact(4).enumerate() {
+            state[i] += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64);
+        }
+        poseidon2.permute_mut(&mut state);
+        state[0] += Goldilocks::ONE;
+        state[1] += Goldilocks::ONE;
+        poseidon2.permute_mut(&mut state);
+
+        let mut hash = [0u8; 64];
+        for i in 0..4 {
+            hash[i * 8..(i + 1) * 8].copy_from_slice(&state[i].as_canonical_u64().to_le_bytes());
+        }
+        poseidon2.permute_mut(&mut state);
+        for i in 0..4 {
+            hash[32 + i * 8..32 + (i + 1) * 8]
+                .copy_from_slice(&state[i].as_canonical_u64().to_le_bytes());
+        }
+
+        assert_eq!(
+            U512::from_big_endian(&hash),
+            qpow_math::get_nonce_hash(header, nonce_be)
+        );
+    }
+
    #[test]
    #[should_panic(expected = "division by zero")]
    fn test_zero_difficulty_panics() {
@@ -0,0 +1,15 @@
+[package]
+name = "quic-transport"
+version.workspace = true
+edition.workspace = true
+publish = false
+description = "Shared QUIC client transport for connecting to a quantus-node's external-miner endpoint"
+
+[dependencies]
+anyhow = { workspace = true }
+quantus-miner-api = { workspace = true }
+quinn = "0.10"
+rustls = { version = "0.21", default-features = false, features = ["dangerous_configuration", "quic"] }
+serde_json = { workspace = true }
+sha2 = "0.10"
+tokio = { workspace = true, features = ["macros", "time"] }
@@ -0,0 +1,218 @@
+//! Shared QUIC client used by miner-service and pool-service to connect to a
+//! quantus-node's external-miner endpoint.
+//!
+//! The node uses a persisted self-signed certificate. Callers must supply the
+//! expected SHA-256 fingerprint of that certificate (lowercase hex); the
+//! transport rejects any other server cert.
+
+use std::net::SocketAddr;
+use std::sync::Arc;
+use std::time::Duration;
+
+use quantus_miner_api::{write_message, MinerMessage, MAX_MESSAGE_SIZE};
+use quinn::{ClientConfig, Endpoint};
+use rustls::client::ServerCertVerified;
+use sha2::{Digest, Sha256};
+
+/// Permanent misconfiguration / auth failure: callers must not reconnect-loop.
+#[derive(Debug)]
+pub struct PermanentConnectError(pub String);
+
+impl std::fmt::Display for PermanentConnectError {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        write!(f, "{}", self.0)
+    }
+}
+
+impl std::error::Error for PermanentConnectError {}
+
+/// Parse and validate miner auth inputs that never become valid by retrying.
+///
+/// Call this at process startup (before workers / HTTP) so a bad fingerprint
+/// or oversized token exits immediately instead of looping forever.
+pub fn validate_auth_config(auth_token: &str, tls_cert_sha256_hex: &str) -> anyhow::Result<()> {
+    parse_fingerprint(tls_cert_sha256_hex)?;
+    validate_ready_frame(auth_token)?;
+    Ok(())
+}
+
+/// Parse a 64-hex (optional whitespace/`:` separators) SHA-256 fingerprint.
+pub fn parse_fingerprint(tls_cert_sha256_hex: &str) -> anyhow::Result<[u8; 32]> {
+    normalize_fingerprint(tls_cert_sha256_hex).ok_or_else(|| {
+        PermanentConnectError("invalid TLS cert SHA-256 fingerprint (expected 64 hex chars)".into())
+            .into()
+    })
+}
+
+fn validate_ready_frame(auth_token: &str) -> anyhow::Result<()> {
+    let frame = serde_json::to_vec(&MinerMessage::Ready {
+        token: auth_token.to_string(),
+    })
+    .map_err(|e| PermanentConnectError(format!("failed to serialize Ready frame: {e}")))?;
+    if frame.len() > MAX_MESSAGE_SIZE as usize {
+        return Err(PermanentConnectError(format!(
+            "auth token serializes to a {}-byte Ready frame; max is {} \
+             (shorten the token or avoid characters that require JSON escaping)",
+            frame.len(),
+            MAX_MESSAGE_SIZE
+        ))
+        .into());
+    }
+    Ok(())
+}
+
+/// Establish a QUIC connection to the node's external-miner endpoint, open
+/// the bidirectional stream, and send the initial `Ready { token }` message.
+///
+/// `tls_cert_sha256_hex` is the lowercase hex SHA-256 of the node's miner TLS
+/// certificate DER (contents of the node's `miner-tls-cert-sha256` file).
+///
+/// Wrong tokens are rejected by the node closing with "auth failed"; that is
+/// returned as [`PermanentConnectError`] so callers do not reconnect forever.
+pub async fn connect(
+    addr: SocketAddr,
+    auth_token: &str,
+    tls_cert_sha256_hex: &str,
+) -> anyhow::Result<(quinn::Connection, quinn::SendStream, quinn::RecvStream)> {
+    // Re-check here so library callers that skip startup validation still fail
+    // closed instead of looping on a deterministic parse error.
+    let expected = parse_fingerprint(tls_cert_sha256_hex)?;
+    validate_ready_frame(auth_token)?;
+
+    let mut crypto = rustls::ClientConfig::builder()
+        .with_safe_defaults()
+        .with_custom_certificate_verifier(Arc::new(PinnedCertVerifier { expected }))
+        .with_no_client_auth();
+    // Versioned with the wire protocol (node side: MINER_ALPN). `/2` = the
+    // authenticated `Ready { token }` protocol; a mismatched node/miner pair
+    // fails at the TLS handshake with "no application protocol" instead of an
+    // opaque auth/deserialize error.
+    crypto.alpn_protocols = vec![b"quantus-miner/2".to_vec()];
+
+    let mut client_config = ClientConfig::new(Arc::new(crypto));
+    let mut transport_config = quinn::TransportConfig::default();
+    transport_config.keep_alive_interval(Some(Duration::from_secs(5)));
+    transport_config.max_idle_timeout(Some(Duration::from_secs(15).try_into().unwrap()));
+    client_config.transport_config(Arc::new(transport_config));
+
+    let mut endpoint = Endpoint::client("0.0.0.0:0".parse().unwrap())?;
+    endpoint.set_default_client_config(client_config);
+
+    let connection = match endpoint.connect(addr, "localhost")?.await {
+        Ok(c) => c,
+        Err(e) => {
+            let msg = e.to_string();
+            // Wrong pin fails the TLS handshake deterministically.
+            if msg.to_ascii_lowercase().contains("fingerprint") {
+                return Err(PermanentConnectError(format!(
+                    "TLS certificate pin rejected ({msg}); check --tls-cert-sha256 / \
+                     miner-tls-cert-sha256"
+                ))
+                .into());
+            }
+            return Err(e.into());
+        }
+    };
+    let (mut send, recv) = connection.open_bi().await?;
+    write_message(
+        &mut send,
+        &MinerMessage::Ready {
+            token: auth_token.to_string(),
+        },
+    )
+    .await?;
+
+    // Ready is fire-and-forget on the wire; the node closes immediately on a
+    // bad token. Wait briefly so that surfaces as a permanent error here
+    // instead of "connected" + 1s reconnect churn.
+    const AUTH_REJECT_GRACE: Duration = Duration::from_secs(2);
+    tokio::select! {
+        reason = connection.closed() => {
+            let msg = reason.to_string();
+            if msg.to_ascii_lowercase().contains("auth") {
+                return Err(PermanentConnectError(format!(
+                    "node rejected miner auth ({msg}); check --auth-token / miner-auth-token"
+                ))
+                .into());
+            }
+            return Err(anyhow::anyhow!("connection closed during auth handshake: {msg}"));
+        }
+        _ = tokio::time::sleep(AUTH_REJECT_GRACE) => {}
+    }
+
+    Ok((connection, send, recv))
+}
+
+fn normalize_fingerprint(input: &str) -> Option<[u8; 32]> {
+    let hex: String = input
+        .chars()
+        .filter(|c| !c.is_whitespace() && *c != ':')
+        .map(|c| c.to_ascii_lowercase())
+        .collect();
+    if hex.len() != 64 || !hex.chars().all(|c| c.is_ascii_hexdigit()) {
+        return None;
+    }
+    let mut out = [0u8; 32];
+    for i in 0..32 {
+        out[i] = u8::from_str_radix(&hex[i * 2..i * 2 + 2], 16).ok()?;
+    }
+    Some(out)
+}
+
+/// Accepts only a server certificate whose SHA-256 fingerprint matches.
+struct PinnedCertVerifier {
+    expected: [u8; 32],
+}
+
+impl rustls::client::ServerCertVerifier for PinnedCertVerifier {
+    fn verify_server_cert(
+        &self,
+        end_entity: &rustls::Certificate,
+        _intermediates: &[rustls::Certificate],
+        _server_name: &rustls::ServerName,
+        _scts: &mut dyn Iterator<Item = &[u8]>,
+        _ocsp_response: &[u8],
+        _now: std::time::SystemTime,
+    ) -> Result<ServerCertVerified, rustls::Error> {
+        let actual = Sha256::digest(&end_entity.0);
+        if actual.as_slice() != self.expected.as_slice() {
+            return Err(rustls::Error::General(
+                "server certificate SHA-256 fingerprint mismatch".into(),
+            ));
+        }
+        Ok(ServerCertVerified::assertion())
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn rejects_short_fingerprint() {
+        let err = validate_auth_config("token", "nope").unwrap_err();
+        assert!(
+            err.downcast_ref::<PermanentConnectError>().is_some(),
+            "{err}"
+        );
+    }
+
+    #[test]
+    fn accepts_valid_fingerprint_and_token() {
+        let fp = "a".repeat(64);
+        validate_auth_config("deadbeef", &fp).unwrap();
+    }
+
+    #[test]
+    fn rejects_token_whose_ready_frame_exceeds_limit() {
+        let fp = "a".repeat(64);
+        // Escape-heavy token: raw length under a naive cap can still blow the frame.
+        let token = "\"".repeat(600);
+        let err = validate_auth_config(&token, &fp).unwrap_err();
+        assert!(
+            err.downcast_ref::<PermanentConnectError>().is_some(),
+            "{err}"
+        );
+        assert!(err.to_string().contains("Ready frame"), "{err}");
+    }
+}
@@ -0,0 +1,13 @@
+[package]
+name = "solver-wasm"
+version.workspace = true
+edition.workspace = true
+publish = false
+description = "Browser WASM solver for Quantus captcha shares (raw C ABI, no wasm-bindgen)"
+
+[lib]
+crate-type = ["cdylib", "rlib"]
+
+[dependencies]
+primitive-types = { workspace = true }
+qpow-math = { workspace = true }
@@ -0,0 +1,101 @@
+//! Browser captcha solver: grinds Poseidon2 nonce hashes over a block header.
+//!
+//! Exposed as a raw C-ABI WASM module (no wasm-bindgen) so it can be built
+//! with a bare `cargo build --target wasm32-unknown-unknown` and loaded with
+//! minimal JS glue. All I/O goes through one static buffer:
+//!
+//! ```text
+//! offset   0..32   header hash                     (in)
+//! offset  32..96   nonce, 64-byte big-endian       (in: start; out: found/next)
+//! offset  96..160  share target, 64-byte BE        (in)
+//! offset 160..224  winning hash, 64-byte BE        (out, only when found)
+//! ```
+//!
+//! JS calls `solve(iters)` in chunks (e.g. 5–20k iterations) to keep the
+//! worker responsive; a return of 0 means "keep going" (the nonce field has
+//! been advanced), 1 means the nonce field now holds a valid share.
+
+use primitive_types::U512;
+
+const HEADER_OFF: usize = 0;
+const NONCE_OFF: usize = 32;
+const TARGET_OFF: usize = 96;
+const HASH_OFF: usize = 160;
+const IO_SIZE: usize = 224;
+
+static mut IO: [u8; IO_SIZE] = [0u8; IO_SIZE];
+
+/// Pointer to the shared I/O buffer in WASM linear memory.
+#[no_mangle]
+pub extern "C" fn io_ptr() -> *mut u8 {
+    core::ptr::addr_of_mut!(IO) as *mut u8
+}
+
+/// Size of the shared I/O buffer, for JS-side sanity checks.
+#[no_mangle]
+pub extern "C" fn io_len() -> u32 {
+    IO_SIZE as u32
+}
+
+/// Grind up to `iters` nonces. Returns 1 if a share was found.
+#[no_mangle]
+pub extern "C" fn solve(iters: u32) -> u32 {
+    let io = unsafe { &mut *core::ptr::addr_of_mut!(IO) };
+
+    let mut header = [0u8; 32];
+    header.copy_from_slice(&io[HEADER_OFF..HEADER_OFF + 32]);
+    let mut nonce = U512::from_big_endian(&io[NONCE_OFF..NONCE_OFF + 64]);
+    let target = U512::from_big_endian(&io[TARGET_OFF..TARGET_OFF + 64]);
+
+    for _ in 0..iters {
+        let nonce_bytes = nonce.to_big_endian();
+        let hash = qpow_math::get_nonce_hash(header, nonce_bytes);
+        if hash < target {
+            io[NONCE_OFF..NONCE_OFF + 64].copy_from_slice(&nonce_bytes);
+            io[HASH_OFF..HASH_OFF + 64].copy_from_slice(&hash.to_big_endian());
+            return 1;
+        }
+        nonce = nonce.saturating_add(U512::one());
+    }
+
+    // Not found: persist progress so the next chunk resumes where we stopped.
+    io[NONCE_OFF..NONCE_OFF + 64].copy_from_slice(&nonce.to_big_endian());
+    0
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn solve_matches_pool_verification() {
+        // Difficulty 4 -> target = MAX/4; expect a solution within a few iters.
+        let header = [7u8; 32];
+        let difficulty = U512::from(4u64);
+        let target = U512::MAX / difficulty;
+
+        let io = unsafe { &mut *core::ptr::addr_of_mut!(IO) };
+        io[HEADER_OFF..HEADER_OFF + 32].copy_from_slice(&header);
+        let start = U512::from(1u64) << 64; // arbitrary "session range" start
+        io[NONCE_OFF..NONCE_OFF + 64].copy_from_slice(&start.to_big_endian());
+        io[TARGET_OFF..TARGET_OFF + 64].copy_from_slice(&target.to_big_endian());
+
+        let mut found = 0;
+        for _ in 0..64 {
+            found = solve(256);
+            if found == 1 {
+                break;
+            }
+        }
+        assert_eq!(found, 1, "should find a share at difficulty 4");
+
+        let nonce = U512::from_big_endian(&io[NONCE_OFF..NONCE_OFF + 64]);
+        let (valid, hash) = qpow_math::is_valid_nonce(header, nonce.to_big_endian(), difficulty);
+        assert!(
+            valid,
+            "pool-side verification must accept the solver's nonce"
+        );
+        let reported = U512::from_big_endian(&io[HASH_OFF..HASH_OFF + 64]);
+        assert_eq!(hash, reported);
+    }
+}
Author	SHA1	Message	Date
illuzenandn13	37f7837f76	bump version to v4.0.0 (#85 ) Co-authored-by: n13 <n13@users.noreply.github.com>	2026-08-13 19:34:16 +08:00
illuzenandCursor	2c8d530de4	Send miner auth token in Ready handshake. (#78 ) * Send miner auth token in Ready handshake. Require --auth-token or --auth-token-file when connecting to a node so the miner matches the node's shared-secret QUIC auth. Co-authored-by: Cursor <cursoragent@cursor.com> * Pin the node's miner TLS certificate by SHA-256 fingerprint. Require --tls-cert-sha256 or --tls-cert-sha256-file so miners reject MITM certs instead of accepting any self-signed server certificate. Co-authored-by: Cursor <cursoragent@cursor.com> * fmt * Docs: prefer reading miner auth token from file, not node logs. Co-authored-by: Cursor <cursoragent@cursor.com> * Version the miner ALPN as quantus-miner/2 to match the node. A node/miner protocol mismatch now fails cleanly at the TLS handshake with "no application protocol" instead of an opaque auth error. Co-authored-by: Cursor <cursoragent@cursor.com> * point at git * Depend on published quantus-miner-api 0.3.0 from crates.io. The temporary git branch pin is no longer needed now that 0.3.0 is published. Co-authored-by: Cursor <cursoragent@cursor.com> * Fail closed on permanent auth/TLS misconfig instead of reconnect looping. Validate the fingerprint and Ready frame size at miner/pool startup before metrics, workers, or HTTP come up. Treat node auth rejection and pin mismatch as PermanentConnectError (no retry), and only reset reconnect backoff after the first NewJob proves auth succeeded. Co-authored-by: Cursor <cursoragent@cursor.com> * Point miner auth docs and examples at miner-auth-token / TLS pin files. The node never logs the auth token; CLI/pool help and every serve/Docker example now require the chain config files (or env vars) instead of implying credentials can be copied from logs. Co-authored-by: Cursor <cursoragent@cursor.com> * fmt * Remove the miner Dockerfile and Docker packaging docs/CI. Container GPU access is a poor fit for this miner, and the Docker examples were unusable with SocketAddr (hostnames rejected). Run the native binary instead. Co-authored-by: Cursor <cursoragent@cursor.com> * Shut down the pool on permanent upstream auth failure; keep reconnecting otherwise. Supervise run_node_client alongside HTTP so a bad token/pin ends the process instead of leaving a live API on a dead upstream. Node restarts still reconnect with backoff; the current job is cleared while disconnected so captchas are not issued against a stale header. Co-authored-by: Cursor <cursoragent@cursor.com> * Replace obsolete EXTERNAL_MINER_PROTOCOL.md with a pointer to MINING.md. The local copy still described the pre-auth protocol (no Ready token, no pin, old ALPN, 16 MB frames). Canonical docs live in quantus-miner-api and the node's MINING.md. Co-authored-by: Cursor <cursoragent@cursor.com> * Delete EXTERNAL_MINER_PROTOCOL.md; README already points at MINING.md. Co-authored-by: Cursor <cursoragent@cursor.com> --------- Co-authored-by: Cursor <cursoragent@cursor.com>	2026-08-13 09:50:03 +08:00
Nikolaus Heger	2b966e780a	Revert "Fuzz GPU/CPU parity at scale and CPU-verify every seal before submission" This reverts commit `b0cfc3005e`.	2026-08-12 15:41:38 +08:00
Nikolaus HegerandClaude Fable 5	b0cfc3005e	Fuzz GPU/CPU parity at scale and CPU-verify every seal before submission Investigating reported seal submission failures. Root cause analysis: the miner's job-id staleness chain (worker check, quic-loop filter, node-side job_id compare) cannot mislabel a candidate — internal and node job ids are updated atomically together. The remaining benign race is node-side: a block-template rebuild (new pre_hash, same parent) does not bump the node's job counter, so an in-flight seal mined on the old template passes the job-id check and fails seal verification, logging "Failed to submit seal". Fixing that belongs in the chain repo (bump the job counter on template rebuild). Miner-side hardening in this commit: - gpu_cpu_parity is now a two-phase fuzzer. Phase 1 dumps FULL 512-bit hashes for randomized (header, nonce) pairs via the real midstate-resume datapath and verifies each against the canonical CPU implementation, multi-threaded (defaults to 1M hashes, ~5s; 10M runs in ~15s). Phase 2 fuzzes whole mining jobs with five profiles: random, carry-edge starts (saturated low limbs), 2^256-boundary crossings, impossible difficulty with exact hash-count assertion, and CPU-known solutions to catch false negatives. Seeded and reproducible; also asserts the submitted work bytes match the nonce. Verified clean: 10M bulk hashes + ~1,500 seals across 6 seeds. - miner-service re-verifies every candidate on the CPU (one hash) before sending JobResult; an engine-produced invalid seal is never submitted and is logged loudly instead. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>	2026-08-12 14:24:05 +08:00
Nikolaus Heger	9b16df1dff	Merge pull request #81 from Quantus-Network/gpu-poseidon2-optimization GPU mining 4.1x faster on Apple Silicon: native-u64 Poseidon2, midstate precompute, lazy second squeeze	2026-08-12 12:10:58 +08:00
Nikolaus Heger	7e449c0a5f	Benchmark: honor --gpu-batch-size/--cpu-batch-size and advance nonce ranges The benchmark fed each worker a fixed range (10k CPU / 1M GPU nonces), capping GPU dispatches at 1M regardless of --gpu-batch-size, and re-searched the same range on every iteration. Worker ranges now step through the nonce space and are sized from the batch flags, floored at the old widths so tiny flags keep engine-internal cancellation batching instead of turning per-call harness overhead into the measured quantity. Batch-size flags are validated non-zero at parse time (serve and benchmark) and echoed in the benchmark banner.	2026-08-12 12:04:44 +08:00
illuzen	22d87d4772	fix benches	2026-08-12 11:16:26 +08:00
Nikolaus HegerandClaude Fable 5	8aa81ec853	Disable naga runtime shader checks for the trusted mining shaders (+8-9%) Exploring a fully native Metal port of the mining kernel showed it runs ~16% faster than the wgpu path at equal batch size. Decomposing that gap: ~9% is naga's injected runtime bounds checks and forced loop bounding, ~1.5% is mulhi()-based multiplication (inexpressible in WGSL), <1% is the constant address space; the rest is binding/codegen residue. The dominant term is recoverable inside wgpu on every backend: the engine now builds its pipelines with create_shader_module_trusted + ShaderRuntimeChecks::unchecked(). This is sound for our shaders: sources are compiled into the binary, every buffer access is a constant-bounded loop index into fixed-size bindings the engine allocates itself, all loops have static bounds, and both variants remain covered by the dual-shader component suite and the CPU parity example. Apple M4: 11.39 -> 12.34 MH/s at the default 1M batch (+8.3%), 11.74 -> 12.73 MH/s at 8M (+8.4%). Cumulative vs main: 4.4x. New examples/trusted_hashrate.rs measures checked vs trusted on the same shader and CPU-verifies results. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>	2026-08-12 00:55:15 +08:00
Nikolaus HegerandClaude Fable 5	58615ed3fa	Adopt PR #80 micro-optimizations that measure positive on the u64 kernel Evaluated all ideas from #80 against the native-u64 kernel. Its three main wins (lazy field arithmetic, precomputed sponge state, conditional second squeeze) were already present in stronger form; two micro-opts carried over and measured positive at large batches (+1.3% at 16M, noise-level at the 1M default): - Hoist midstate/target/start-nonce storage reads out of the nonce loop. - Produce byte-swapped hash words on demand during the target compare, so the reject path skips building the swapped hash entirely. Forcing nonces_per_thread=2 was also tested and regressed; one thread per nonce stays. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>	2026-08-12 00:01:39 +08:00
Nikolaus HegerandClaude Fable 5	d55b970ab6	Optimize GPU Poseidon2 mining 4.1x on Apple Silicon (u64 shader, midstate, lazy squeeze) Three stacked optimizations, each verified bit-exact against the CPU reference (qp-poseidon-core) by the dual-shader component suite and a new GPU/CPU parity example: - mining_u64.wgsl: native 64-bit Goldilocks arithmetic with plonky2-style lazy reduction, used automatically when the adapter supports SHADER_INT64 (all Apple Silicon, NVIDIA, modern AMD). mining.wgsl remains the 32-bit fallback. (+50%) - Midstate precompute: the first two of five sponge permutations absorb the header and high nonce half, both constant per batch; they are now computed once on the CPU (pow_core::mining_midstate) and resumed on the GPU. Batches are clamped so nonces never carry past 2^256. (+77%) - Lazy second squeeze: the first squeeze yields the most significant 256 bits of the hash, which decide hash-vs-target unless exactly equal to the target's high half, so the common reject path skips the final permutation. (+56%) Apple M4 throughput: 2.81 -> 11.4 MH/s at the default 1M batch (criterion large_range_1m/gpu: 355.8ms -> 86.3ms, -75.7%). Also: - Fix criterion GPU bench crash: thread-local worker resources are now tagged with an engine id so multiple GpuEngines per process never mix devices. - Fix stale end-to-end test harness (6-binding layout) and run the full component suite against both shader variants. - New examples: hashrate, gpu_cpu_parity (25 jobs incl. 2^256-boundary crossing), gpu_features. - Add Apple M5 family GPU tiers. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>	2026-08-11 23:40:12 +08:00
Nikolaus Heger	0135e65278	Merge pull request #76 from Quantus-Network/add-security-policy Add security policy	2026-08-11 22:08:43 +08:00
Nikolaus Heger	2eb0341a70	chore: bump qp-poseidon-core to 3.1.0 (#77 ) Update workspace pin from 2.1.0 to 3.1.0 for sponge wipe hardening. (qpow-math from the pinned chain tag may still pull 2.1.0 transitively.)	2026-08-08 15:47:42 +08:00
Nikolaus HegerandClaude Fable 5	0c61e3da34	Add security policy Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>	2026-08-07 12:11:38 +08:00
illuzenandCursor	71781a0a80	Add captcha share pool service and browser WASM solver (#74 ) * Add captcha share pool service and browser WASM solver pool-service connects to a node over the external-miner QUIC protocol and turns browser captcha solves into real mining shares: sessions with disjoint nonce ranges over the current header, single-use share tokens, a reCAPTCHA-shaped /siteverify endpoint, and upstream block submission when a share meets full network difficulty. solver-wasm is a raw C-ABI wasm32 build of the Poseidon2 nonce grinder (no wasm-bindgen) for the embeddable widget. Co-authored-by: Cursor <cursoragent@cursor.com> * fmt * Re-queue block solution when upstream write fails A network-difficulty share taken off the solution channel was permanently lost if the JobResult write to the node failed: the connection loop broke to reconnect but the block was no longer anywhere. Park the block in a pending slot that survives reconnection and is retried before any other work. Co-authored-by: Cursor <cursoragent@cursor.com> * fmt * Address PR review: capacity limits, loud block drops, shared QUIC transport, constant-time secret compare - Cap live sessions and tokens (configurable via --max-sessions/--max-tokens, default 100k each); /api/session returns 503 at_capacity so a request flood can no longer grow the maps without bound. - A block solution that fails to queue upstream is now logged at error level, and blocks_found/block_found only report solutions actually queued. - Extract the node QUIC connect + insecure-verifier pattern into a shared quic-transport crate used by both miner-service and pool-service; derive the network target via pow_core::JobContext instead of re-deriving it. - Compare the site secret in constant time (constant_time_eq). - Note single-tenant token semantics on site_secret for the host-registration follow-up. Co-authored-by: Cursor <cursoragent@cursor.com> * Add per-IP rate limiting on /api/session issuance The global session cap stops unbounded map growth, but a single client could still churn entries at the TTL boundary. Limit each IP to 60 issuances per minute (configurable via --sessions-per-ip-per-min); excess requests get 429 rate_limited. Stale IP windows are GC'd alongside sessions/tokens. Co-authored-by: Cursor <cursoragent@cursor.com> * clippy --------- Co-authored-by: Cursor <cursoragent@cursor.com>	2026-07-07 17:52:38 +08:00
illuzen	15edf4a806	bump version to v3.3.1 (#73 )	2026-06-30 16:34:53 +08:00
illuzen	a47db3d1a2	Fix windows again (#72 ) * handle lost device correctly * Update lib.rs * allow-integrated flag otherwise exclude * DeviceLost is a status * Integrated GPU skipped when discrete init fails * nits	2026-06-30 16:30:11 +08:00
illuzen	8359892406	bump version to v3.3.0 (#71 )	2026-06-30 13:47:10 +08:00
Nikolaus Hegerandilluzen	d99be29a64	Select GPU adapters by backend to fix multi-backend OOM (#67 ) On Windows wgpu enumerates each physical GPU once per backend (Vulkan + DX12) plus a CPU-emulated fallback ("Microsoft Basic Render Driver"). Building a mining context for every entry causes VRAM contention and OOMs the process during benchmark/serve startup (#61). Instead of deduplicating by (vendor, device) PCI IDs - which would collapse rigs with multiple identical cards into a single context - drop CPU-emulated adapters and keep all adapters from the highest-ranked backend present (Vulkan/Metal, then DX12). Within a single backend each physical GPU appears exactly once, so identical cards are preserved by construction and no physical-ID matching is needed. Selected adapters are ordered discrete-first so `--gpu-devices 1` picks the discrete card on hybrid laptops. Skipped adapters are logged at info level; if nothing usable remains, init fails with an explicit error. Selection logic is a pure index-based function unit-tested against the exact enumeration reported in #61, identical multi-GPU rigs, DX12-only machines, and software-only environments. Co-authored-by: illuzen <illuzen@users.noreply.github.com>	2026-06-30 13:13:45 +08:00
illuzen	7296c003bb	Improve Windows gpu support (#70 ) * support AMD GPUs on windows * add support for more GPUs * remove shader dump from logs * handle gpu initialization failure better * fmt * re-order matching to avoid mobile / desktop false positives * harden gpu matching * add tests to engine-gpu module * handle dead device and failed mapping correctly * handle timeout properly with tokio * nits * fix benchmarks * fmt	2026-06-30 12:55:35 +08:00
I Dewa Gede Bisma Mahendra	d37f1dc0e5	Merge pull request #69 from Quantus-Network/beast/add-use-main-in-workflow feat: make dockerfile use main an option	2026-06-24 17:43:08 +08:00
Beast	4d53fff54b	feat: make dockerfile use main an option	2026-06-24 17:16:00 +08:00
I Dewa Gede Bisma Mahendra	18e8539fa3	Merge pull request #68 from Quantus-Network/beast/fix-docker-build fix: docker build	2026-06-24 17:00:40 +08:00
Beast	4ab6914331	fix: docker build	2026-06-24 16:44:49 +08:00