← All proposals

LP-132: QuasarGPU Execution Adapter

Abstract

QuasarGPU is **CEVM's GPU-side execution adapter for Quasar-certified

rounds**. It executes the round on GPU — Block-STM, EVM fibers, batched

crypto, root construction — and produces the artifacts the Quasar

consensus engine certifies: block_hash, state_root, receipts_root,

execution_root, mode_root. It optionally aggregates per-lane

certificates from inbound vote/share/proof artifacts.

Critically, this is an execution adapter, not a replacement for

luxfi/consensus. The boundary is:

lux/consensus    orders, votes, certifies, finalizes
cevm             executes, validates, roots, receipts
quasar/gpu/      GPU-side execution adapter for Quasar rounds

> One sentence: **consensus decides what must be executed; cevm proves

> what execution produced; Quasar certifies the resulting commitment.**

Architecture

Wave-tick scheduler (no host phase orchestration)

Classical GPU EVM models dispatch a separate kernel per phase

(exec → wait → validate → wait → commit → wait). QuasarGPU replaces

all of that with one bounded kernel re-launched once per wave tick:

host:                                  GPU wave-tick kernel:
─────                                  ──────────────────────
begin_round(QuasarRoundDescriptor) → │ 12 service drains run in parallel:
push_txs(...)                       │
push_certs(...) (P/Q/Z lanes)       │   gid 0: Ingress    → Decode
poll_state_requests() → host LSM    │   gid 1: Decode     → Crypto/StateRequest
push_state_pages(...)               │   gid 2: Crypto     → Commit/DagReady/Exec
poll_quasar_certs() ← from CertOut  │   gid 3: DagReady   → Exec
                                    │   gid 4: Exec       → Validate (EVM fibers)
run_wave_tick() — re-dispatch ──────┤   gid 5: Validate   → Commit/Repair (Block-STM)
                                    │   gid 6: Repair     → Exec (incarnation++)
poll_round_result() ← QuasarRoundResult │   gid 7: Commit  → root material
                                    │   gid 8: StateRequest (host poll)
                                    │   gid 9: StateResp  → resume Crypto
                                    │   gid 10: CertLane  → batch verify + aggregate
                                    │   gid 11: CertOut   (host poll)

No persistent hot-spinning kernel. Each wave tick is a bounded

dispatch; workgroups exit when their service ring is empty or their

budget is exhausted. The GPU scheduler is free to interleave other

kernels between ticks.

Service map (3.0: 12 services; 4.0: 16 services)

enum class ServiceId : uint32_t {
    // 3.0 services (12) — substrate, shipped 2025-12-25
    Ingress           = 0,   // host tx blobs
    Decode            = 1,   // sender recovery / admission gate
    Crypto            = 2,   // sig verify, route to Commit / DagReady / Exec
    DagReady          = 3,   // MVCC ready set (Nebula mode)
    Exec              = 4,   // EVM fiber VM
    Validate          = 5,   // Block-STM read-set check
    Repair            = 6,   // re-execute conflicting txs
    Commit            = 7,   // commit + per-tx receipt keccak
    StateRequest      = 8,   // GPU → host page faults (out)
    StateResp         = 9,   // host → GPU page replies (in)
    CertLane          = 10,  // BLS / Pulsar / MLDSAGroth16 lane artifacts
    CertOut           = 11,  // GPU-emitted per-lane QuasarCert commitments

    // 4.0 services (4 new) — production, shipped 2026-02-14 (v0.43, v0.47, v0.48)
    AdaptiveSchedule  = 12,  // ConflictSpec admission, predictor consult (A)
    BridgeAttest      = 13,  // cross-chain attestation drain (B)
    MarketAuction     = 14,  // auction-rule batched matching for hot DEX lanes (M)
    FiberCheckpoint   = 15,  // per CALL-frame checkpoint emit / GC (F)

    Count             = 16
};

The 4.0 service set is referred to as A/B/M/F for the four new

drains. Each new service uses a dedicated ring with the same RingHeader

+ items_arena layout as the 3.0 services and the same wave-tick budget

contract. The 4.0 activation height (2026-02-14) is when Count flips

from 12 to 16 in the layout type.

Every ring is a fixed-capacity device buffer with a RingHeader at the

front (head/tail/capacity/mask/items_ofs/pushed/consumed) and items

laid out back-to-back in a per-round items_arena. Cross-workgroup

visibility uses the relaxed-atomic + threadgroup_barrier(mem_device)

pattern proven in v0.30 (V3 wave-dispatch).

Layout types (host/device shared)

Inline keccak-f[1600]

The kernel ships an inline FIPS-202 keccak permutation (24 rounds,

RC table, rotation table). Used for:

• block_hash — finalization digest
• state_root — accumulated commit chain
• receipts_root — per-tx receipt hash chain
• execution_root — Block-STM trace commitment
• per-cert subject re-derivation (replay protection)

Zero CPU keccak fallback on the round path. The metal::keccak256_cpu

helper in lib/evm/gpu/metal/keccak_host.mm is reserved for legacy

non-Quasar consumers.

EVM fiber VM (LP-009 § Fiber model)

drain_exec runs the in-kernel EVM interpreter — a per-tx fiber

with:

• pc, sp, gas, status
• stack[64] of 256-bit values (4 × 64-bit limbs each = 13.6 MB for 4096 fibers)
• memory[1024] per-fiber scratch
• pending_key_*[4] cold-state suspend slot
• rw[8] Block-STM read/write set

Currently: 118 opcodes (full arithmetic incl. ADDMOD / MULMOD /

EXP, all signed/unsigned compares, bitwise incl. SAR/BYTE,

KECCAK256, env opcodes, PUSH0..PUSH32, DUP1..16, SWAP1..16,

MLOAD / MSTORE / MSTORE8 / MSIZE, SLOAD / SSTORE with cold-

miss suspend, JUMP / JUMPI / PC / JUMPDEST / GAS, STOP /

RETURN / REVERT / INVALID).

Cold-miss SLOAD → suspend (status=2, key in pending_key_*, push

StateRequest, restore stack, return). drain_state_resp resumes by

re-injecting and stamping the MVCC slot's `last_writer_tx |=

0x80000000` as the "loaded" sentinel.

Deferred to v0.40+: CALL family, CREATE/CREATE2, LOGn,

EXTCODE*, RETURNDATA*, TLOAD/TSTORE, MCOPY, BLOBHASH/BLOBBASEFEE.

Block-STM (LP-010 § QuasarSTM)

drain_exec records read+write entries into the tx's RW set.

drain_validate runs MVCC version-check (mvcc_check_consistent); on

mismatch → conflict → repair queue with bumped incarnation. On match →

mvcc_apply_writes (atomic version bump) → push CommitItem.

Telemetry on QuasarRoundResult: conflict_count, repair_count,

fibers_suspended, fibers_resumed. Real measured behavior: 16 same-

key contending txs produce **120 conflicts → 120 repairs → 16

commits**. That's textbook Block-STM running entirely on GPU.

Async cold-state page faults

The host services GPU-emitted state requests via its LSM/cache/disk

path. The GPU never blocks on host I/O:

host loop:
    while round in progress:
        run_wave_tick()
        reqs = poll_state_requests()
        if reqs.empty(): continue
        pages = state_db.batch_get(reqs)
        push_state_pages(pages)

A faulted tx exits Decode into StateRequest and the workgroup moves on.

Other txs continue through the fast path while the slow tx awaits host

service. Finalization gate (commit.consumed == ingress.pushed) holds

across both fast and slow lanes.

Per-lane cert verification (LP-020 § 3.0)

drain_cert_lane reads inbound QuasarCertIngress items, dispatches

to a per-lane verifier, and accumulates stake into per-lane counters.

At 2/3 quorum threshold, a QuasarCert is emitted onto CertOut.

Three verifiers (LP-132 §§drain_cert_lane), shipped state per lane:

The 3.0 launch (2025-12-25) shipped HMAC-keccak placeholders across all

three lanes — real cryptographic verification (one-way with a master

secret; cross-lane domain tags reject replay), structured so the swap

to real BLS / Module-LWE (Corona) / Groth16 was a single function-pointer change.

That swap landed under QuasarSTM 4.0 (LP-135) in v0.44 (BLS) and v0.45

(Pulsar + Groth16) and activated on 2026-02-14. The HMAC-keccak

path is preserved as a development-only mode for deterministic test

vectors and is gated behind EVM_DEV_HMAC_VERIFIER=1. See

LP-010-quasar-stm-4 (4.0 paper, 2026-02-14) for the migration.

Public API

namespace quasar::gpu {

class QuasarGPUEngine {
public:
    static std::unique_ptr<QuasarGPUEngine> create();

    QuasarRoundHandle begin_round(const QuasarRoundDescriptor&);
    void              push_txs(QuasarRoundHandle, std::span<const HostTxBlob>);
    void              push_certs(QuasarRoundHandle, std::span<const HostCertIngress>);
    void              push_state_pages(QuasarRoundHandle, std::span<const HostStatePage>);

    std::vector<HostStateRequest> poll_state_requests(QuasarRoundHandle);
    std::vector<HostQuasarCert>   poll_quasar_certs(QuasarRoundHandle);

    QuasarRoundResult run_wave_tick(QuasarRoundHandle);
    QuasarRoundResult run_until_done(QuasarRoundHandle, std::size_t max_ticks);
    QuasarRoundResult poll_round_result(QuasarRoundHandle);

    void request_close(QuasarRoundHandle);
    void end_round(QuasarRoundHandle);

    // ...
};

}  // namespace quasar::gpu

Backends

The Metal and CUDA kernels share the layout types (quasar_gpu_layout.hpp)

byte-for-byte; structural drift is caught by static_assert at compile

time on both sides.

Scaling Formula

T ≈ min(
    E_gpu,                                    // GPU event execution capacity
    B_net / bytes_per_event,                  // network bandwidth
    B_state / bytes_per_state_access,         // hot-state memory bandwidth
    L_independent / conflict_amplification,   // lane independence
    C_cert / cert_cost_per_root,              // certificate throughput
    D_da / data_availability_bytes            // data availability
)

QuasarGPU's design thesis:

> **Increase T by maximizing L_independent and minimizing

> cert_cost_per_root.**

L_independent comes from lane partitioning (LP-010 § Lanes).

cert_cost_per_root is constant in TPS because validators sign roots,

not individual txs.

Throughput claim ladder (honest)

The pitch:

> Billion-event throughput, EVM-settled, Quasar-certified.

Not "1 B EVM TPS on a single shared state machine" — that's bounded by

contention and data availability, not GPU compute.

Test surface

The substrate ships with a comprehensive test surface

(test/unittests/quasar_gpu_engine_test.mm):

Plus the broader 19-binary GPU test surface (parity, modes, host-bridge,

pipeline, unified, gpu-state, opcodes, dispatch, refund, access-list,

stack-depth, storage-overflow, etc.). All 19 PASS on Apple M1 Max.

Forbidden patterns

QuasarGPU must not use:

• one global STM clock
• one global version map lock
• retry-until-success GPU loops
• nondeterministic conflict victim selection
• opcode-level STM (it's tx-level)
• CPU compute on the round path (host is doorbell + I/O only)
• per-version malloc (use the per-round arena)
• global hot-key spinlocks

Implementation Plan (3.0 substrate + 4.0 production)

3.0 substrate (shipped 2025-12-25)

4.0 production train (shipped 2026-02-14, activation height)

The QuasarSTM 4.0 production spec is LP-135. The 4.0 production paper

(LP-010-quasar-stm-4, 2026-02-14) is the canonical reference.

References

Changelog

• 2025-12-15 — 3.0 substrate spec, 12 service IDs, three HMAC-keccak

cert verifiers, 118-opcode EVM fiber, single-GPU; activated 2025-12-25.

• 2026-02-14 — QuasarSTM 4.0 activation: 16 service IDs (12 + 4

A/B/M/F), real BLS12-381 / Corona Module-LWE / Groth16 verifiers

(v0.44/v0.45), 175-opcode EVM with full CALL/CREATE/LOG/EXTCODE/

RETURNDATA/TLOAD/TSTORE/MCOPY/BLOBHASH/BLOBBASEFEE coverage (v0.46),

predictive scheduling + Chiron hint roots (v0.47), Motor VersionBlock

(v0.48), CSMV scaffold + formal CPU reference + cross-backend

determinism harness (v0.49). Wire schema unchanged. See LP-135 and

LP-010-quasar-stm-4.

Copyright

Copyright (C) 2025-2026, Lux Partners Limited. All rights reserved.