wai-quantum 0.4.4

//! The fabric's CUDA member: the gate kernel of
//! [`quantum_fabric::cuda`](crate::quantum_fabric::cuda) on an NVIDIA GPU,
//! compiled at run time for the device's own architecture. See that module
//! for the kernel and the argument that it is bit-exact.
//!
//! - **Discovery** ([`discover`]): the driver and the runtime compiler are
//!   loaded when the process first asks for them, so a build with this feature
//!   runs on a host with no CUDA at all and finds no member there
//!   ([`Discovery::unavailable`] says why). Every device the driver lists is
//!   opened, the kernel is compiled for it, and its arithmetic is probed
//!   against [`lanes`] in this process. A device that cannot host the kernel,
//!   or whose arithmetic disagrees, is listed with the reason ([`Refusal`]);
//!   one that passes is a [`CudaMember`], present as [`Presence::Detected`].
//! - **Running** ([`CudaMember::run_program`]): the state stays on the device;
//!   one launch per gate, in op-log order, on one stream; then the report word
//!   and the state come back. The kernel reports exactly when the reference's
//!   result leaves 32 bits; the member then returns
//!   [`CudaError::OutOfRange`], never a state.
//! - **Certification** ([`CudaMember::certify`], [`CudaMember::member`]): the
//!   member's own run of a corpus, folded case by case into a
//!   [`Certification`]. The corpus digest covers inputs and outputs only, so a
//!   bit-exact member reproduces the CPU's pinned digests; a device that runs
//!   nothing, skips a case or gets one bit wrong cannot.
//! - **Placement, receipts and metering** are as for the WebGPU member:
//!   [`place`](crate::quantum_fabric::place),
//!   [`FabricReceipt`](crate::quantum_fabric::FabricReceipt) and a caller's
//!   [`RunMeter`], whose window runs from the first upload to the last
//!   readback.
//!
//! Every call is synchronous; the device is waited on before a call returns.
//! No type of the CUDA bindings appears in this module's public API.
//!
//! # Honest boundaries
//!
//! - **A member is what this process measured.** The probe and the
//!   certificates are earned per process; the member id binds the device, its
//!   UUID, the driver, the runtime compiler and the architecture.
//! - **Plans are per gate.** [`Config::PER_GATE`]: a tile kernel in shared
//!   memory, as the WebGPU member has, is not here yet.
//! - **Unmetered by default**, as for every member.

use crate::quantum::{Amp, Circuit, Gateset, MAX_QUBITS, ONE, QuantumError, StateVector};
use crate::quantum_energy::EnergyClass;
use crate::quantum_fabric::cuda::{self, DeviceIdentity};
use crate::quantum_fabric::lanes::{self, GateWords, Lane};
use crate::quantum_fabric::plan::{self, Config, LaneRun, Program, Step};
use crate::quantum_fabric::{
    Certificate, Certification, Corpus, EnergyCalibration, FabricClass, KernelEntry, Member, MemberError, Presence, RunFuture, RunMeter, Runner, Unmetered, Window,
    apply_kernel, pair_updates,
};
use cudarc::driver::{CudaContext, CudaFunction, CudaSlice, CudaStream, DriverError, LaunchConfig, PushKernelArg};
use cudarc::nvrtc::{CompileOptions, compile_ptx_with_opts};
use std::sync::{Arc, Mutex, OnceLock};

/// One probe case's device results: [`lanes::mul_wide`] and
/// [`lanes::fxmul_wide`], each as `(high, low)` words.
pub type ProbeOutput = ((u32, u32), (u32, u32));

/// A run's lanes and report, with the meter's accepted reading (zero and
/// `None` when unmetered).
type MeteredRun = (LaneRun, (u64, Option<EnergyClass>));

/// Why a device is not a member, or a run has no state.
#[derive(Debug, PartialEq, Eq)]
#[non_exhaustive]
pub enum CudaError {
    /// The driver or the runtime compiler could not be used, and why.
    Unavailable(String),
    /// The runtime compiler refused the kernel for this device.
    Compile(String),
    /// The driver reported an error, or the device was lost.
    Device(String),
    /// Out of memory on the device or the host.
    Memory(String),
    /// Exactly what the CPU reference refuses.
    Circuit(QuantumError),
    /// More qubits than this member's memory holds.
    TooWide { n_qubits: u8, max: u8 },
    /// The reference's result leaves 32 bits at this op; the member returns
    /// no state.
    OutOfRange { gate: u32 },
    /// A program the host cannot lay out (malformed words, a start outside the
    /// state, a step this member does not run): a defect of the host, never a
    /// result.
    Internal(&'static str),
    /// The device's arithmetic disagreed with the CPU's on this pair.
    Probe { a: i32, b: i32 },
}

impl std::fmt::Display for CudaError {
    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
        match self {
            CudaError::Unavailable(why) => write!(f, "CUDA is unavailable: {why}"),
            CudaError::Compile(e) => write!(f, "the kernel did not compile: {e}"),
            CudaError::Device(e) => write!(f, "device: {e}"),
            CudaError::Memory(e) => write!(f, "out of memory: {e}"),
            CudaError::Circuit(e) => write!(f, "{e}"),
            CudaError::TooWide { n_qubits, max } => write!(f, "{n_qubits} qubits; this member holds at most {max}"),
            CudaError::OutOfRange { gate } => write!(f, "the result of op {gate} leaves 32 bits"),
            CudaError::Internal(why) => write!(f, "host defect: {why}"),
            CudaError::Probe { a, b } => write!(f, "the device's arithmetic differs from the CPU's at ({a}, {b})"),
        }
    }
}

impl std::error::Error for CudaError {}

impl From<CudaError> for MemberError {
    fn from(e: CudaError) -> MemberError {
        match e {
            CudaError::TooWide { n_qubits, max } => MemberError::TooWide { n_qubits, max },
            CudaError::OutOfRange { gate } => MemberError::OutOfRange { gate },
            CudaError::Memory(why) => MemberError::Memory(why),
            CudaError::Device(why) => MemberError::Device(why),
            CudaError::Circuit(e) => MemberError::Circuit(e),
            other => MemberError::Other(other.to_string()),
        }
    }
}

/// A driver error as this module reports it: out of memory as
/// [`CudaError::Memory`], anything else as [`CudaError::Device`].
fn driver(e: DriverError) -> CudaError {
    if e.0 == cudarc::driver::sys::cudaError_enum::CUDA_ERROR_OUT_OF_MEMORY {
        CudaError::Memory(e.to_string())
    } else {
        CudaError::Device(e.to_string())
    }
}

/// A device that is not a member, and why.
#[derive(Debug)]
pub struct Refusal {
    /// The device's name, or its ordinal when the driver gave none.
    pub device: String,
    /// Its CI key ([`DeviceIdentity::key`]), when the device could be read.
    pub key: String,
    pub error: CudaError,
}

/// What [`discover`] found.
#[derive(Debug)]
pub struct Discovery {
    /// Devices that compiled the kernel and passed the probe, by ordinal.
    pub found: Vec<CudaMember>,
    /// Devices that did not, with why.
    pub refused: Vec<Refusal>,
    /// Why no device was even listed: no driver, no runtime compiler, or the
    /// driver failed to start. `None` when the driver listed its devices.
    pub unavailable: Option<String>,
}

/// Whether the driver and the runtime compiler can be loaded in this process,
/// or why not.
pub fn available() -> Result<(), String> {
    // SAFETY: each only tries to open a shared library by its known names.
    if !unsafe { cudarc::driver::sys::is_culib_present() } {
        return Err("no CUDA driver library on this host".into());
    }
    // SAFETY: as above.
    if !unsafe { cudarc::nvrtc::sys::is_culib_present() } {
        return Err("no CUDA runtime compiler (NVRTC) library on this host".into());
    }
    Ok(())
}

/// Every device the driver lists, each opened, compiled for and probed in this
/// process before it is listed. On a host without the driver or the runtime
/// compiler, nothing is found and [`Discovery::unavailable`] says why.
pub fn discover() -> Discovery {
    let mut d = Discovery { found: Vec::new(), refused: Vec::new(), unavailable: None };
    if let Err(why) = available() {
        d.unavailable = Some(why);
        return d;
    }
    let count = match CudaContext::device_count() {
        Ok(n) => usize::try_from(n).unwrap_or(0),
        Err(e) => {
            d.unavailable = Some(format!("the driver did not start: {e}"));
            return d;
        }
    };
    for ordinal in 0..count {
        match open(ordinal) {
            Ok(m) => d.found.push(m),
            Err(r) => d.refused.push(r),
        }
    }
    d
}

/// The target architecture as the runtime compiler's options take it: a
/// `&'static str`, one per distinct architecture for the life of the process.
fn arch_option(arch: String) -> &'static str {
    static ARCHES: OnceLock<Mutex<Vec<&'static str>>> = OnceLock::new();
    let mut arches = ARCHES.get_or_init(|| Mutex::new(Vec::new())).lock().unwrap_or_else(|p| p.into_inner());
    if let Some(a) = arches.iter().find(|a| **a == arch) {
        return a;
    }
    let leaked: &'static str = Box::leak(arch.into_boxed_str());
    arches.push(leaked);
    leaked
}

/// The driver's CUDA version (`major·1000 + minor·10`), or 0.
fn driver_version() -> u32 {
    let mut v: std::ffi::c_int = 0;
    // SAFETY: writes one integer; the driver library is loaded.
    let ok = unsafe { cudarc::driver::sys::cuDriverGetVersion(&mut v) } == cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS;
    if ok { u32::try_from(v).unwrap_or(0) } else { 0 }
}

/// The runtime compiler's version, or `(0, 0)`.
fn compiler_version() -> (u32, u32) {
    let (mut major, mut minor): (std::ffi::c_int, std::ffi::c_int) = (0, 0);
    // SAFETY: writes two integers; the compiler library is loaded.
    let ok = unsafe { cudarc::nvrtc::sys::nvrtcVersion(&mut major, &mut minor) } == cudarc::nvrtc::sys::nvrtcResult::NVRTC_SUCCESS;
    if ok { (u32::try_from(major).unwrap_or(0), u32::try_from(minor).unwrap_or(0)) } else { (0, 0) }
}

/// The widest state of `2^n` interleaved 8-byte lanes that half of `free`
/// bytes holds, at most [`MAX_QUBITS`].
fn qubits_for(free: u64) -> u8 {
    let lanes = free / 2 / 8;
    if lanes < 2 { 0 } else { (lanes.ilog2() as u8).min(MAX_QUBITS) }
}

/// Reads the device at `ordinal`, compiles the kernel for it and probes it.
fn open(ordinal: usize) -> Result<CudaMember, Refusal> {
    let unread = |error: CudaError| Refusal { device: format!("device {ordinal}"), key: String::new(), error };
    let ctx = CudaContext::new(ordinal).map_err(|e| unread(driver(e)))?;
    let name = ctx.name().map_err(|e| unread(driver(e)))?;
    let (major, minor) = ctx.compute_capability().map_err(|e| unread(driver(e)))?;
    let uuid = ctx.uuid().map_err(|e| unread(driver(e)))?;
    let attr = |a| ctx.attribute(a).map(|v| u32::try_from(v).unwrap_or(0));
    use cudarc::driver::sys::CUdevice_attribute::*;
    let pci = match (attr(CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID), attr(CU_DEVICE_ATTRIBUTE_PCI_BUS_ID), attr(CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID)) {
        (Ok(d), Ok(b), Ok(v)) => format!("{d:04x}:{b:02x}:{v:02x}"),
        _ => String::new(),
    };
    let memory = ctx.total_mem().map_err(|e| unread(driver(e)))? as u64;
    let identity = DeviceIdentity {
        name: name.clone(),
        uuid: uuid.bytes.iter().map(|b| format!("{:02x}", *b as u8)).collect(),
        compute: (u32::try_from(major).unwrap_or(0), u32::try_from(minor).unwrap_or(0)),
        pci,
        driver: driver_version(),
        compiler: compiler_version(),
        memory,
    };
    let key = identity.key();
    let refuse = |error: CudaError| Refusal { device: name.clone(), key: key.clone(), error };
    let kernel = Kernel::compile(&ctx, &identity, cuda::SOURCE).map_err(refuse)?;
    let (free, _) = ctx.mem_get_info().map_err(|e| refuse(driver(e)))?;
    let max_qubits = qubits_for(free as u64);
    let detail = format!(
        "cuda; device {name:?}; sm_{major}{minor}; uuid {}; pci {}; driver {}; runtime compiler {}.{}; memory {memory} bytes, {free} free at discovery; up to {max_qubits} qubits; a metered run's window runs from the first upload to the last readback",
        identity.uuid, identity.pci, identity.driver, identity.compiler.0, identity.compiler.1,
    );
    let member = CudaMember {
        stream: ctx.default_stream(),
        ctx,
        kernel,
        id: identity.member_id(),
        key: key.clone(),
        implementation: cuda::implementation_id(),
        detail,
        max_qubits,
        identity,
    };
    let pairs = lanes::probe_pairs(4096, 0x5eed);
    match member.probe(&pairs) {
        Ok(out) => match lanes::first_mismatch(&pairs, &out) {
            None => Ok(member),
            Some((a, b)) => Err(refuse(CudaError::Probe { a, b })),
        },
        Err(e) => Err(refuse(e)),
    }
}

/// The kernel compiled for one device: its two entry points.
#[derive(Clone)]
struct Kernel {
    apply: CudaFunction,
    probe: CudaFunction,
}

impl Kernel {
    fn compile(ctx: &Arc<CudaContext>, identity: &DeviceIdentity, source: &str) -> Result<Kernel, CudaError> {
        let opts = CompileOptions {
            arch: Some(arch_option(identity.arch())),
            options: cuda::COMPILE_OPTIONS.iter().map(|s| s.to_string()).collect(),
            ..Default::default()
        };
        let ptx = compile_ptx_with_opts(source, opts).map_err(|e| CudaError::Compile(e.to_string()))?;
        let module = ctx.load_module(ptx).map_err(driver)?;
        Ok(Kernel { apply: module.load_function("apply_gate").map_err(driver)?, probe: module.load_function("probe").map_err(driver)? })
    }
}

/// The launch shape for `work` items: [`cuda::BLOCK`] threads per block and
/// at most [`cuda::GRID_MAX`] blocks, the kernel striding over the rest.
fn launch(work: u64) -> LaunchConfig {
    let blocks = work.div_ceil(u64::from(cuda::BLOCK)).clamp(1, u64::from(cuda::GRID_MAX)) as u32;
    LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (cuda::BLOCK, 1, 1), shared_mem_bytes: 0 }
}

/// One CUDA device as a fabric member: a context on it, the kernel compiled
/// for it, and the arithmetic probe passed in this process.
#[derive(Clone)]
pub struct CudaMember {
    ctx: Arc<CudaContext>,
    stream: Arc<CudaStream>,
    kernel: Kernel,
    identity: DeviceIdentity,
    id: String,
    key: String,
    implementation: String,
    detail: String,
    max_qubits: u8,
}

impl std::fmt::Debug for CudaMember {
    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
        f.debug_struct("CudaMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).field("max_qubits", &self.max_qubits).finish()
    }
}

impl CudaMember {
    /// The member id ([`DeviceIdentity::member_id`]).
    pub fn id(&self) -> &str {
        &self.id
    }

    /// The key a runner declares for it ([`DeviceIdentity::key`]).
    pub fn key(&self) -> &str {
        &self.key
    }

    /// Always [`FabricClass::Gpu`].
    pub fn class(&self) -> FabricClass {
        self.identity.class()
    }

    /// What discovery found: device, architecture, driver, compiler, memory,
    /// and what a metered run's window covers. Not signed.
    pub fn detail(&self) -> &str {
        &self.detail
    }

    /// The most qubits a run holds on this member.
    pub fn max_qubits(&self) -> u8 {
        self.max_qubits
    }

    /// The implementation id ([`cuda::implementation_id`], or a mutant's).
    pub fn implementation(&self) -> &str {
        &self.implementation
    }

    /// What the member id hashes.
    pub fn identity(&self) -> &DeviceIdentity {
        &self.identity
    }

    /// The device's [`lanes::mul_wide`] and [`lanes::fxmul_wide`] of each
    /// pair, in order.
    pub fn probe(&self, pairs: &[(i32, i32)]) -> Result<Vec<ProbeOutput>, CudaError> {
        if pairs.is_empty() {
            return Ok(Vec::new());
        }
        let flat: Vec<i32> = pairs.iter().flat_map(|&(a, b)| [a, b]).collect();
        let cases = self.stream.clone_htod(&flat).map_err(driver)?;
        let mut out: CudaSlice<u32> = self.stream.alloc_zeros(pairs.len() * 4).map_err(driver)?;
        let count = pairs.len() as u64;
        let mut args = self.stream.launch_builder(&self.kernel.probe);
        args.arg(&cases).arg(&mut out).arg(&count);
        // SAFETY: `cases` holds `count` int2 cases and `out` `count` uint4
        // outputs, the kernel's parameters in order.
        unsafe { args.launch(launch(count)) }.map_err(driver)?;
        let words = self.stream.clone_dtoh(&out).map_err(driver)?;
        Ok(words.as_chunks::<4>().0.iter().map(|w| ((w[0], w[1]), (w[2], w[3]))).collect())
    }

    /// The planning this member runs: [`Config::PER_GATE`].
    pub fn config(&self) -> Config {
        Config::PER_GATE
    }

    /// Runs `p` from its basis state: the lanes, and the first op whose exact
    /// result left 32 bits. A program another config planned (one with tile
    /// steps) is [`CudaError::Internal`]: this member runs gate steps only.
    pub fn run_program(&self, p: &Program) -> Result<LaneRun, CudaError> {
        Ok(self.run_program_metered(p, &mut Unmetered)?.0)
    }

    fn run_program_metered<M: RunMeter + ?Sized>(&self, p: &Program, meter: &mut M) -> Result<MeteredRun, CudaError> {
        let n = p.n_qubits;
        if n == 0 {
            return Ok((LaneRun { amps: vec![[ONE as i32, 0]], out_of_range: None }, (0, None)));
        }
        if n > self.max_qubits {
            return Err(CudaError::TooWide { n_qubits: n, max: self.max_qubits });
        }
        if n > MAX_QUBITS || p.start >= 1usize << n {
            return Err(CudaError::Internal("the start is outside the state"));
        }
        let len = 1usize << n;
        let mut gates = Vec::with_capacity(p.steps.len());
        for step in &p.steps {
            let Step::Gate { gate } = *step else { return Err(CudaError::Internal("a tile step; this member runs gate steps only")) };
            let g = p.gates.get(gate as usize).ok_or(CudaError::Internal("a step names a gate the program does not hold"))?;
            if !g.is_well_formed() {
                return Err(CudaError::Internal("malformed gate words for this state"));
            }
            gates.push((&g.words, g.index));
        }
        let mut init: Vec<i32> = Vec::new();
        init.try_reserve_exact(2 * len).map_err(|e| CudaError::Memory(e.to_string()))?;
        init.resize(2 * len, 0);
        init[2 * p.start] = ONE as i32;
        self.execute(init, &gates, meter)
    }

    /// Gates in order from `state`, as [`lanes::run_words`] runs them: the
    /// lanes, and the lowest position in `gates` of a gate whose exact result
    /// left 32 bits. A state that is not `2^n` lanes (`n ≥ 1`) or words
    /// [`lanes::apply_words`] refuses for it are [`CudaError::Internal`].
    pub fn run_lanes(&self, state: &[Lane], gates: &[GateWords]) -> Result<LaneRun, CudaError> {
        let len = state.len();
        if len < 2 || !len.is_power_of_two() {
            return Err(CudaError::Internal("a state that is not 2^n lanes, n at least 1"));
        }
        let n = len.trailing_zeros() as u8;
        if n > self.max_qubits {
            return Err(CudaError::TooWide { n_qubits: n, max: self.max_qubits });
        }
        let planned: Vec<(&GateWords, u32)> = gates.iter().zip(0u32..).collect();
        let mut init: Vec<i32> = Vec::new();
        init.try_reserve_exact(2 * len).map_err(|e| CudaError::Memory(e.to_string()))?;
        init.extend(state.iter().flat_map(|l| [l[0], l[1]]));
        Ok(self.execute(init, &planned, &mut Unmetered)?.0)
    }

    /// Uploads `init` (interleaved lanes, `2·2^n` words), runs `gates` in order
    /// inside a window of `meter`, and reads the report and the state back.
    /// Every gate's words are checked against the state first: the words
    /// [`lanes::apply_words`] refuses, the host refuses.
    fn execute<M: RunMeter + ?Sized>(&self, init: Vec<i32>, gates: &[(&GateWords, u32)], meter: &mut M) -> Result<MeteredRun, CudaError> {
        let len = init.len() / 2;
        for (w, _) in gates {
            let (tbit, ctrl) = (w.tbit as usize, w.ctrl as usize);
            if tbit == 0 || !tbit.is_power_of_two() || tbit > len / 2 || ctrl & tbit != 0 || ctrl & !(len - 1) != 0 {
                return Err(CudaError::Internal("malformed gate words for this state"));
            }
        }
        let pairs = (len / 2) as u64;
        let mut window = Window::new(meter);
        window.open();
        let mut amps = self.stream.clone_htod(&init).map_err(driver)?;
        drop(init);
        let mut report = self.stream.clone_htod(&[u32::MAX]).map_err(driver)?;
        for &(w, index) in gates {
            let m: [i32; 8] = w.m.map(|x| x as i32);
            let mut args = self.stream.launch_builder(&self.kernel.apply);
            args.arg(&mut amps).arg(&pairs).arg(&w.tbit).arg(&w.ctrl);
            for v in &m {
                args.arg(v);
            }
            args.arg(&index).arg(&mut report);
            // SAFETY: `amps` holds `2·pairs` int2 lanes and `report` one word,
            // and the words were checked against the state above.
            unsafe { args.launch(launch(pairs)) }.map_err(driver)?;
        }
        let bad = self.stream.clone_dtoh(&report).map_err(driver)?;
        let words = self.stream.clone_dtoh(&amps).map_err(driver)?;
        let reading = window.close();
        let out_of_range = bad.first().copied().filter(|&k| k != u32::MAX);
        let mut lanes_out: Vec<Lane> = Vec::new();
        lanes_out.try_reserve_exact(len).map_err(|e| CudaError::Memory(e.to_string()))?;
        lanes_out.extend(words.as_chunks::<2>().0.iter().copied());
        Ok((LaneRun { amps: lanes_out, out_of_range }, reading))
    }

    /// Equal to `c.simulate_from(start)`, or a typed refusal: what the CPU
    /// reference refuses ([`CudaError::Circuit`]), more qubits than the member
    /// holds, or a result that leaves 32 bits.
    pub fn simulate_from(&self, c: &Circuit, start: usize) -> Result<StateVector, CudaError> {
        Ok(self.simulate_with(c, start, &mut Unmetered)?.0)
    }

    /// [`Self::simulate_from`] inside a window of `meter`: the state, and the
    /// microjoules and class the meter read, which are zero and `None` unless
    /// the crate's label check accepts them. The window opens just before the
    /// first upload and closes after the last readback; planning and the
    /// conversion to a state vector are outside it.
    pub fn simulate_metered(&self, c: &Circuit, start: usize, meter: &mut dyn RunMeter) -> Result<(StateVector, u64, Option<EnergyClass>), CudaError> {
        self.simulate_with(c, start, meter)
    }

    fn simulate_with<M: RunMeter + ?Sized>(&self, c: &Circuit, start: usize, meter: &mut M) -> Result<(StateVector, u64, Option<EnergyClass>), CudaError> {
        let p = plan::compile_with(c, start, &self.config()).map_err(CudaError::Circuit)?;
        if p.n_qubits > self.max_qubits {
            return Err(CudaError::TooWide { n_qubits: p.n_qubits, max: self.max_qubits });
        }
        let (run, (joules_micro, class)) = self.run_program_metered(&p, meter)?;
        if let Some(gate) = run.out_of_range {
            return Err(CudaError::OutOfRange { gate });
        }
        let mut amps = Vec::new();
        amps.try_reserve_exact(run.amps.len()).map_err(|e| CudaError::Memory(e.to_string()))?;
        amps.extend(run.amps.iter().map(|l| Amp { re: i64::from(l[0]), im: i64::from(l[1]) }));
        Ok((StateVector { n_qubits: p.n_qubits, amps }, joules_micro, class))
    }

    /// [`Self::simulate_from`] the all-zeros state.
    pub fn simulate(&self, c: &Circuit) -> Result<StateVector, CudaError> {
        self.simulate_from(c, 0)
    }

    /// The certificate of [`apply_kernel`]`(gs)` as [`Self::implementation`]
    /// on this member, over `corpus`: each circuit run from `|0…0⟩`, against
    /// the native simulator. A range report inside a corpus is an error: its
    /// circuits are valid and within the horizon.
    pub fn certify(&self, gs: Gateset, corpus: Corpus) -> Result<Certificate, CudaError> {
        Ok(self.certify_metered(gs, corpus, &mut Unmetered)?.0)
    }

    fn certify_metered<M: RunMeter + ?Sized>(&self, gs: Gateset, corpus: Corpus, meter: &mut M) -> Result<(Certificate, Option<(u64, EnergyClass)>), CudaError> {
        let mut cert = Certification::new();
        let mut calibration = EnergyCalibration::new();
        for c in corpus.circuits(gs) {
            let reference = lanes::reference_output(&c).map_err(CudaError::Circuit)?;
            let (run, (joules_micro, class)) = self.run_program_metered(&plan::compile_with(&c, 0, &self.config()).map_err(CudaError::Circuit)?, &mut *meter)?;
            if let Some(gate) = run.out_of_range {
                return Err(CudaError::OutOfRange { gate });
            }
            cert.add(&lanes::circuit_bytes(&c), &reference, &lanes::output(&run.amps));
            calibration.add(pair_updates(&c), joules_micro, class.as_ref());
        }
        Ok((cert.finish(apply_kernel(gs), &self.implementation, &self.id, &corpus.generator(gs)), calibration.finish()))
    }

    /// A fabric [`Member`]: present as [`Presence::Detected`], with one entry
    /// per gateset, each certified in this call on `corpus`, and no energy
    /// figure.
    pub fn member(&self, corpus: Corpus) -> Result<Member, CudaError> {
        self.member_with(corpus, &mut Unmetered)
    }

    /// [`Self::member`], each corpus run inside its own window of `meter`:
    /// every entry also carries the energy figure its runs earn
    /// ([`EnergyCalibration`]), or none if any run went unmetered.
    pub fn member_metered(&self, corpus: Corpus, meter: &mut dyn RunMeter) -> Result<Member, CudaError> {
        self.member_with(corpus, meter)
    }

    fn member_with<M: RunMeter + ?Sized>(&self, corpus: Corpus, meter: &mut M) -> Result<Member, CudaError> {
        let mut kernels = Vec::new();
        for gs in Gateset::ALL {
            let (certificate, energy) = self.certify_metered(gs, corpus, &mut *meter)?;
            kernels.push(KernelEntry { kernel: apply_kernel(gs).into(), implementation: self.implementation.clone(), certificate, energy });
        }
        Ok(Member { id: self.id.clone(), class: self.class(), presence: Presence::Detected, kernels, detail: self.detail.clone() })
    }

    /// The same member holding at most `max_qubits` qubits (no more than it
    /// holds), so that tests can make it refuse a circuit as too wide.
    #[doc(hidden)]
    pub fn with_max_qubits(&self, max_qubits: u8) -> CudaMember {
        CudaMember { max_qubits: max_qubits.min(self.max_qubits), ..self.clone() }
    }

    /// The same device running `source` in place of the kernel, under the
    /// implementation id that text earns: for the mutation tests, which
    /// require every mutant to be caught.
    #[doc(hidden)]
    pub fn with_source(&self, source: &str) -> Result<CudaMember, CudaError> {
        let kernel = Kernel::compile(&self.ctx, &self.identity, source)?;
        Ok(CudaMember { kernel, implementation: cuda::implementation_id_of(source), ..self.clone() })
    }
}

/// The member on the fabric ([`Fabric`](crate::quantum_fabric::Fabric)): its
/// own implementation, through [`CudaMember::simulate_metered`]. Pair it with
/// a [`Member`] whose entries name that implementation, such as
/// [`CudaMember::member`]'s; another implementation is
/// [`MemberError::Other`].
impl Runner for CudaMember {
    fn member_id(&self) -> String {
        self.id.clone()
    }

    fn run<'r>(&'r self, implementation: &'r str, c: &'r Circuit, start: usize, meter: &'r mut dyn RunMeter) -> RunFuture<'r> {
        Box::pin(async move {
            if implementation != self.implementation {
                return Err(MemberError::Other(format!("member {} runs implementation {}, not {implementation:?}", self.id, self.implementation)));
            }
            Ok(self.simulate_with(c, start, meter)?)
        })
    }
}

#[cfg(test)]
mod tests {
    use super::*;

    #[test]
    fn a_driver_out_of_memory_is_a_memory_error_and_every_other_a_device_error() {
        use cudarc::driver::sys::cudaError_enum::*;
        assert!(matches!(driver(DriverError(CUDA_ERROR_OUT_OF_MEMORY)), CudaError::Memory(_)));
        assert!(matches!(driver(DriverError(CUDA_ERROR_LAUNCH_FAILED)), CudaError::Device(_)));
        assert!(MemberError::from(CudaError::Memory("x".into())).falls_back());
        assert!(MemberError::from(CudaError::OutOfRange { gate: 3 }).falls_back());
        assert!(!MemberError::from(CudaError::Probe { a: 1, b: 2 }).falls_back());
        assert!(!MemberError::from(CudaError::Internal("x")).falls_back());
    }

    #[test]
    fn the_state_a_member_holds_is_half_its_free_memory_at_most() {
        assert_eq!(qubits_for(0), 0);
        assert_eq!(qubits_for(31), 0);
        assert_eq!(qubits_for(32), 1);
        // 6 GiB free: 3 GiB of 8-byte lanes is 2^28.58 lanes, so 28 qubits, capped.
        assert_eq!(qubits_for(6 << 30), MAX_QUBITS.min(28));
        assert_eq!(qubits_for(u64::MAX), MAX_QUBITS);
    }

    #[test]
    fn the_launch_shape_covers_the_work_and_caps_the_blocks() {
        assert_eq!(launch(0).grid_dim.0, 1);
        assert_eq!(launch(1).grid_dim.0, 1);
        assert_eq!(launch(u64::from(cuda::BLOCK) + 1).grid_dim.0, 2);
        assert_eq!(launch(1 << 40).grid_dim.0, cuda::GRID_MAX);
        assert_eq!(launch(7).block_dim, (cuda::BLOCK, 1, 1));
    }

    #[test]
    fn an_architecture_is_leaked_once_per_process() {
        let a = arch_option("compute_89".into());
        let b = arch_option("compute_89".into());
        assert!(std::ptr::eq(a, b));
        assert_eq!(arch_option("compute_90".into()), "compute_90");
    }
}