aprender-serve 0.65.1

//! CUDA-accelerated quantized model
//!
//! This module provides GPU-accelerated inference for quantized models
//! using NVIDIA CUDA.
//!
//! # Architecture
//!
//! `OwnedQuantizedModelCuda` wraps an `OwnedQuantizedModel` with a CUDA executor
//! for GPU-accelerated matrix operations. Key features:
//!
//! - GPU-resident KV cache (avoids CPU→GPU transfer per token)
//! - Fused attention kernels
//! - Pre-cached quantized weights
//! - Batch generation support
//!
//! # Module Structure
//!
//! - `backend.rs`: CUDA kernel configuration and PTX generation (CudaBackend)
//! - `forward.rs`: Forward pass methods (single token, cached, GPU-resident)
//! - `generation.rs`: Token generation loops (basic, cached, streaming, batch)
//! - `speculative.rs`: Speculative decoding (self-speculative, draft model)
//! - `weights.rs`: Weight management (pre-caching, GPU upload)
//!
//! # Example
//!
//! ```rust,ignore
//! use realizar::gguf::{OwnedQuantizedModel, OwnedQuantizedModelCuda};
//!
//! let model = OwnedQuantizedModel::from_mapped(&mapped)?;
//! let mut cuda_model = OwnedQuantizedModelCuda::new(model, 0)?; // GPU 0
//!
//! // GPU-accelerated forward pass
//! let logits = cuda_model.forward_cuda(&tokens)?;
//! ```

mod backend;
mod forward;
mod generation;
mod speculative;
mod weights;

// Re-export types for public API
pub use backend::CudaBackend;
// PMAT-072: Step-wise batched decode state for lock-releasing scheduler
pub use generation::BatchedDecodeState;

use crate::error::{RealizarError, Result};

// Import types from peer modules (parent of cuda/)
use super::model::OwnedQuantizedModel;
use super::quantized::{OwnedQKVWeights, OwnedQuantizedTensor};
use super::runtime::{OwnedQuantizedKVCache, QuantizedGenerateConfig};
use super::utils::verbose;

// =============================================================================
// IMP-800: CUDA-Accelerated Model Wrapper
// =============================================================================

/// Error from CUDA model initialization that preserves the unconsumed model.
///
/// When `OwnedQuantizedModelCuda::new()` fails, the model is returned inside this error
/// so callers can fall back to CPU without an expensive 1GB clone.
pub struct CudaInitError {
    /// The initialization error
    pub error: RealizarError,
    /// The unconsumed model, returned for CPU fallback.
    ///
    /// Boxed to keep the `Err` variant small (`clippy::result_large_err`):
    /// `OwnedQuantizedModel` is ~8 KB by value, which would bloat every
    /// `Result<_, CudaInitError>` on the stack. Boxing is a single move (no
    /// clone), preserving the cheap CPU-fallback recovery semantics.
    model: Box<OwnedQuantizedModel>,
}

impl CudaInitError {
    /// Extract the unconsumed model for CPU fallback
    #[must_use]
    pub fn into_model(self) -> OwnedQuantizedModel {
        *self.model
    }
}

impl std::fmt::Display for CudaInitError {
    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
        write!(f, "{}", self.error)
    }
}

impl std::fmt::Debug for CudaInitError {
    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
        write!(f, "CudaInitError({:?})", self.error)
    }
}

/// CUDA-accelerated wrapper for `OwnedQuantizedModel` (IMP-800a)
///
/// Provides GPU-accelerated forward pass using NVIDIA CUDA via trueno-gpu.
/// Caches the CudaExecutor to avoid initialization overhead (~50ms) per call.
///
/// # Example
///
/// ```rust,ignore
/// use realizar::gguf::{OwnedQuantizedModel, OwnedQuantizedModelCuda};
///
/// let model = OwnedQuantizedModel::from_mapped(&mapped)?;
/// let mut cuda_model = OwnedQuantizedModelCuda::new(model, 0)?; // GPU 0
///
/// // GPU-accelerated forward pass
/// let logits = cuda_model.forward_cuda(&tokens)?;
/// ```
pub struct OwnedQuantizedModelCuda {
    /// Inner model
    pub(crate) model: OwnedQuantizedModel,
    /// Cached CUDA executor
    pub(crate) executor: crate::cuda::CudaExecutor,
    /// GPU device name
    device_name: String,
    /// GPU memory (free, total) in bytes
    memory_info: (usize, usize),
    /// PAR-083: Pre-allocated embedding buffer to eliminate per-token heap allocation.
    /// Five-Whys root cause: embed() allocates Vec<f32> per token (~14KB for 7B).
    /// Fix: Reuse this buffer with embed_into().
    embed_buf: Vec<f32>,
    /// realizr#199 (PMAT-450): Prefix cache for prompt KV reuse.
    /// Stores GPU KV cache snapshots keyed by prompt tokens.
    /// On cache hit, skip prefill entirely (TTFT ~900ms → ~5ms).
    #[cfg(feature = "gpu")]
    prefix_cache: crate::gguf::batch_scheduler::PrefixCache,
    /// PP-LLAMA-001 §10: every input the `max_batch` ceiling was sized from,
    /// captured BEFORE `CUDA_MAX_BATCH` is written, plus which of the two
    /// sources produced it. `None` on a model built before sizing ran.
    max_batch_sizing: Option<crate::cuda::gpu_profile::MaxBatchSizing>,
    /// §9 #7: free VRAM measured AFTER weights and every warmed cache are
    /// resident. With `memory_info` (taken before) this gives the measured cost
    /// of preload — the bulk of the 9.5 GB the accounting is missing.
    free_after_preload: Option<usize>,
    /// §3: the phase split of the LAST request this model served on the
    /// serialized (write-lock) path. Valid only while that write lock is held,
    /// which is exactly when the handler reads it.
    last_phase_timings: crate::api::PhaseTimings,
}

impl OwnedQuantizedModelCuda {
    /// Create a new CUDA-accelerated model wrapper
    ///
    /// # Arguments
    ///
    /// * `model` - The quantized model to wrap
    /// * `device_ordinal` - GPU device index (0 for first GPU)
    ///
    /// # Errors
    ///
    /// Returns error if CUDA is not available or device doesn't exist.
    /// Create a CUDA model wrapper. Model is consumed on failure.
    /// Use `with_max_seq_len` for recoverable errors (returns model on failure).
    pub fn new(model: OwnedQuantizedModel, device_ordinal: i32) -> Result<Self> {
        Self::with_max_seq_len(model, device_ordinal, 2048).map_err(|e| e.error)
    }

    /// Create a new CUDA-accelerated model wrapper with custom max sequence length
    ///
    /// # Arguments
    ///
    /// * `model` - The quantized model to wrap
    /// * `device_ordinal` - GPU device index (0 for first GPU)
    /// * `max_seq_len` - Maximum sequence length for GPU KV cache (PAR-018)
    ///
    /// # Errors
    ///
    /// Returns `CudaInitError` containing both the error and the unconsumed model,
    /// allowing callers to recover the model for CPU fallback without cloning.
    /// GH-280: Check GPU capability before allocation.
    /// Extracted to reduce cognitive complexity of `with_max_seq_len`.
    fn check_gpu_capability(
        model: OwnedQuantizedModel,
    ) -> std::result::Result<OwnedQuantizedModel, CudaInitError> {
        // PMAT-824: use the MODEL-AWARE op set. `required_ops(constraints)` alone
        // cannot distinguish Gemma v1 from Gemma2/Gemma3 (the arch-constraints
        // contract collapses all three onto one alias row), so it would report
        // them all GPU-supported (GatedMlp→SwiGLU reads as "supported"). The
        // model-aware variant adds the Gemma2/Gemma3 softcap + post-attn/post-FFN
        // norm ops the CUDA forward does NOT implement, routing those models to
        // CPU at the CAPABILITY layer (belt-and-suspenders with the runtime
        // cosine parity gate) instead of relying solely on the parity check.
        let required = crate::capability::required_ops_for_model(
            &model.config.constraints,
            &model.config.architecture,
        );
        let supported = crate::capability::gpu_supported_ops();
        if let Err(missing) = crate::capability::check_capability(&required, &supported) {
            let missing_names: Vec<String> = missing.iter().map(ToString::to_string).collect();
            return Err(CudaInitError {
                error: RealizarError::CapabilityMismatch {
                    architecture: model.config.architecture.clone(),
                    missing_ops: missing_names.join(", "),
                    suggestion: "Model will use CPU inference. To add GPU support, implement the missing kernels in trueno.".to_string(),
                },
                model: Box::new(model),
            });
        }
        Ok(model)
    }

    /// PMAT-785: Quant-eligibility gate — refuse GPU residency for any model
    /// carrying a GGML quant type without a verified GPU GEMV kernel.
    ///
    /// `OwnedQuantizedModel::has_gpu_unsupported_quant` inspects every projection
    /// tensor (lm_head, QKV, attn output, FFN gate/up/down) against the
    /// `gpu_unsupported_quant_qtype` whitelist {F32, Q4_0, Q4_1, Q5_0, Q8_0,
    /// Q4_K, Q5_K, Q6_K}. Anything outside it (F16, Q5_1, Q8_1, Q2_K, Q3_K,
    /// Q8_K, IQ*, BF16, unknown) would hit the GPU weight upload's
    /// `resolve_qtype().unwrap_or(Q4K)` and be SILENTLY decoded as Q4_K → garbage
    /// logits (PMAT-781/783). The model is returned unconsumed so the caller can
    /// fall back to CPU (loud) without a 1 GB clone.
    ///
    /// Centralizes the fail-closed gate at the single CUDA-model constructor that
    /// every serve `generate_gpu_resident` entry point reaches, closing the
    /// bypass on the serve chat/completions/batch HTTP handlers (which consume
    /// the shared `state.cuda_model()` built here) and the `apr run --gpu`
    /// GGUF/APR CLI paths.
    fn check_quant_gpu_capability(
        model: OwnedQuantizedModel,
    ) -> std::result::Result<OwnedQuantizedModel, CudaInitError> {
        if model.has_gpu_unsupported_quant() {
            return Err(CudaInitError {
                error: RealizarError::CapabilityMismatch {
                    architecture: model.config.architecture.clone(),
                    missing_ops: "GPU GEMV kernel for the model's quantization type".to_string(),
                    suggestion: "Model carries a quant type without a verified GPU kernel \
                                 (e.g. Q5_1/Q8_1/Q2_K/Q3_K/Q8_K/F16/IQ*). It will use CPU \
                                 inference to avoid silent Q4_K-decode garbage (PMAT-785). \
                                 Requantize to a GPU-supported type (Q4_0/Q4_1/Q5_0/Q8_0/\
                                 Q4_K/Q5_K/Q6_K) for GPU acceleration."
                        .to_string(),
                },
                model: Box::new(model),
            });
        }
        Ok(model)
    }

    /// GH-199/PARITY-GATE: Preload GPU weights and verify correctness.
    /// Extracted to reduce cognitive complexity of `with_max_seq_len`.
    fn preload_and_verify(mut self) -> std::result::Result<Self, CudaInitError> {
        if !self.supports_gpu_resident() {
            return Ok(self);
        }

        // GH-199 ROOT CAUSE B: Eagerly preload weights for GPU-resident path.
        if let Err(e) = self.preload_weights_gpu() {
            return Err(CudaInitError {
                error: e,
                model: Box::new(self.into_model()),
            });
        }

        // PMAT-037: Eagerly warm FP16 weight cache for HGEMM prefill.
        // Five-Whys root cause: FP16 cache lazily populated on first inference
        // request (303ms cold-start), inflating TTFT P50 from 25ms to 50.9ms.
        // Pre-populating at model init moves cost to startup (one-time).
        // PMAT-067: Skip FP16 cache when FP8 is active — saves ~1.5 GB VRAM.
        // PMAT-400: Skip FP16 cache on unified memory (cc>=120) — saves 61 GB for 32B model.
        // PMAT-409: Override with FORCE_FP16_CACHE=1 for 7B on GB10 (2.9 GB, restores prefill perf).
        let skip_fp16_unified = self.executor.gpu_profile.cc >= 120
            && std::env::var("FORCE_FP16_CACHE").as_deref() != Ok("1");
        if std::env::var("HGEMM_PREFILL").as_deref() != Ok("0")
            && !self.executor.gpu_profile.fp8_prefill
            && !skip_fp16_unified
        {
            let num_layers = self.model.config.num_layers;
            let hidden_dim = self.model.config.hidden_dim as u32;
            let intermediate_dim = self.model.config.intermediate_dim as u32;
            let vocab_size = self.model.config.vocab_size as u32;
            if let Err(e) = self.executor.ensure_cublas() {
                eprintln!("[PMAT-037] cuBLAS init failed (non-fatal): {e}");
            } else if let Err(e) = self.executor.warmup_hgemm_cache(
                num_layers,
                hidden_dim,
                intermediate_dim,
                vocab_size,
            ) {
                eprintln!("[PMAT-037] FP16 cache warmup failed (non-fatal): {e}");
            }
        }

        // PMAT-053/067: FP8 E4M3 weight cache warmup (auto-enabled on sm_89+)
        // GH-286: Skip if --no-fp8-cache (saves ~1.5 GB RSS)
        let no_fp8 = std::env::var("REALIZR_NO_FP8_CACHE").as_deref() == Ok("1");
        if self.executor.gpu_profile.fp8_prefill && !no_fp8 {
            let num_layers = self.model.config.num_layers;
            let hidden_dim = self.model.config.hidden_dim as u32;
            let intermediate_dim = self.model.config.intermediate_dim as u32;
            let vocab_size = self.model.config.vocab_size as u32;
            if let Err(e) =
                self.executor
                    .warmup_fp8_cache(num_layers, hidden_dim, intermediate_dim, vocab_size)
            {
                eprintln!("[PMAT-053] FP8 cache warmup failed (non-fatal): {e}");
            }
        }

        // PMAT-091: Interleaved Q4K weight cache warmup (W4A16 WMMA GEMM)
        if self.executor.gpu_profile.w4a16_interleaved {
            let num_layers = self.model.config.num_layers;
            let hidden_dim = self.model.config.hidden_dim as u32;
            let intermediate_dim = self.model.config.intermediate_dim as u32;
            let vocab_size = self.model.config.vocab_size as u32;
            if let Err(e) = self.executor.warmup_interleaved_cache(
                num_layers,
                hidden_dim,
                intermediate_dim,
                vocab_size,
            ) {
                eprintln!("[PMAT-091] Interleaved cache warmup failed (non-fatal): {e}");
            }
        }

        // GH-181: Reinitialize workspace after cache warmup (FP8/FP16/interleaved).
        // Cache allocations can relocate workspace buffers, causing the parity
        // gate's forward pass to read stale pointers (cosine -0.28 on RTX 4060).
        {
            let hidden_dim = self.model.config.hidden_dim;
            let intermediate_dim = self.model.config.intermediate_dim;
            self.executor.force_workspace_reinit();
            if let Err(e) = self.executor.init_workspace(hidden_dim, intermediate_dim) {
                eprintln!("[GH-181] Workspace reinit failed (non-fatal): {e}");
            }
        }

        // PARITY-GATE: Jidoka — stop-the-line if GPU diverges from CPU.
        // Run ONE token through both backends and compare logits.
        // If cosine similarity < 0.99, refuse to construct.
        // Skip gate if SKIP_PARITY_GATE=1 (for debugging the gate itself)
        //
        // M-GPU-MOE-1.3 (qwen3-moe-forward-gpu-v1 v1.3.0): Skip the
        // parity gate for MoE architectures. The gate calls
        // `forward_single_with_cache` (CPU dense) and
        // `forward_gpu_resident` (GPU dense) — both dense forward
        // paths that index `layer.ffn_up_weight` directly. For MoE
        // those weights are `dense_ffn_placeholder` (byte_size=0)
        // because the real expert tensors live in `moe_layers`, so
        // the dense forward panics in `fused_matmul_f32` at the
        // empty data slice.
        //
        // The MoE-specific parity gate (FALSIFY-QW3-MOE-GPU-PARITY-001)
        // is exercised by `qwen3_moe_gpu_parity.rs` which uses the
        // MoE forward methods directly, bypassing the dense path
        // this gate runs.
        let skip_gate = std::env::var("SKIP_PARITY_GATE")
            .map(|v| v == "1")
            .unwrap_or(false);

        if !skip_gate && !self.model.config.constraints.is_moe {
            if let Err(e) = parity_gate(&mut self) {
                return Err(CudaInitError {
                    error: e,
                    model: Box::new(self.into_model()),
                });
            }
        }

        // §9 #7: the ONE snapshot that makes preload's VRAM cost accountable.
        // `memory_info` was taken before `init_kv_cache_gpu`, weight upload and
        // every cache warmup, so on its own it says nothing about what this
        // process is holding. The delta between the two is measured, not
        // estimated, and it is what the effective-config endpoint reports.
        if let Some((free, _total, _peak)) = self.executor.sample_vram_used() {
            self.free_after_preload = Some(free);
        }

        Ok(self)
    }

    /// Create a GPU-accelerated inference engine with a custom maximum sequence length.
    pub fn with_max_seq_len(
        model: OwnedQuantizedModel,
        device_ordinal: i32,
        max_seq_len: usize,
    ) -> std::result::Result<Self, CudaInitError> {
        use crate::cuda::CudaExecutor;

        // GH-279: Contract gate — validate architecture and dimensions before CUDA init
        if let Err(e) = crate::contract_gate::validate_model_load_basic(
            &model.config.architecture,
            model.config.num_layers,
            model.config.hidden_dim,
            model.config.num_heads,
            model.config.num_kv_heads,
            model.config.intermediate_dim,
            model.config.vocab_size,
        ) {
            return Err(CudaInitError {
                error: crate::contract_gate::gate_error(e),
                model: Box::new(model),
            });
        }

        // GH-280: Capability gate — refuse GPU if model requires unsupported ops.
        let model = Self::check_gpu_capability(model)?;

        // PMAT-785: Quant gate — refuse GPU for any quant type lacking a verified
        // GPU GEMV kernel, so it cannot be silently decoded as Q4_K → garbage.
        // Centralizes the PMAT-781/783 fail-closed gate at the single CUDA-model
        // constructor that EVERY serve `generate_gpu_resident` entry point funnels
        // through (serve chat/completions/batch via the shared state model, and
        // the `apr run --gpu` GGUF/APR CLI paths). Returns the model unconsumed so
        // callers fall back to CPU (loud) instead of shipping GPU garbage.
        let model = Self::check_quant_gpu_capability(model)?;

        let mut executor = match CudaExecutor::new(device_ordinal) {
            Ok(e) => e,
            Err(e) => {
                return Err(CudaInitError {
                    error: RealizarError::UnsupportedOperation {
                        operation: "CudaExecutor::new".to_string(),
                        reason: format!("CUDA initialization failed: {e}"),
                    },
                    model: Box::new(model),
                });
            },
        };

        let device_name = executor
            .device_name()
            .unwrap_or_else(|_| "Unknown GPU".to_string());
        let memory_info = executor.memory_info().unwrap_or((0, 0));

        // PAR-018: Initialize GPU-resident KV cache for attention acceleration
        // This avoids ~66 MB CPU→GPU transfer per token for TinyLlama
        let num_layers = model.layers.len();
        let num_heads = model.config.num_heads;
        let num_kv_heads = model.config.num_kv_heads; // PAR-021 GQA support
        let head_dim = model.config.hidden_dim / num_heads;

        if let Err(e) =
            executor.init_kv_cache_gpu(num_layers, num_heads, num_kv_heads, head_dim, max_seq_len)
        {
            return Err(CudaInitError {
                error: RealizarError::UnsupportedOperation {
                    operation: "init_kv_cache_gpu".to_string(),
                    reason: format!("GPU KV cache initialization failed: {e}"),
                },
                model: Box::new(model),
            });
        }

        // PAR-118: Initialize Flash Decoding for split-K attention acceleration.
        // Five-Whys: batched_incremental_attention uses Grid=(num_heads,M,1) Block=(32,1,1)
        // = only 896 threads on RTX 4090 for 7B (28 heads). Flash Decoding splits KV cache
        // into chunks processed in parallel, achieving 1.5-2x decode speedup.
        // NOTE: flash_decode_enabled is used by BOTH the graphed decode path
        // (flash_decoding_graphed in attention.rs) and the batched path (batched.rs).
        // The batched path uses threshold 1024 to avoid triggering during normal prefill.
        if let Err(e) = executor.init_flash_decoding(num_heads, head_dim, max_seq_len, 1) {
            if verbose() {
                eprintln!(
                    "[PAR-118] Flash Decoding init failed: {e}, falling back to sequential attention"
                );
            }
            // Non-fatal: sequential attention still works, just slower
        }

        // PAR-060: Set RoPE theta for position embeddings
        if verbose() {
            eprintln!(
                "[PAR-060] Setting rope_theta = {} for GPU path",
                model.config.rope_theta
            );
        }
        executor.set_rope_theta(model.config.rope_theta);

        // CORRECTNESS-011: Set rope_type for correct RoPE style (NORM vs NEOX)
        if verbose() {
            eprintln!(
                "[CORRECTNESS-011] Setting rope_type = {} for GPU path (0=NORM, 2=NEOX)",
                model.config.rope_type
            );
        }
        executor.set_rope_type(model.config.rope_type);

        // GH-129: Pre-load ALL kernel modules BEFORE heavy GPU allocations.
        // Five-Whys root cause: PTX JIT compilation requires GPU memory for the
        // JIT compiler itself. On Jetson (unified memory), weight upload consumes
        // ~1 GB, leaving less for JIT. Moving kernel preload here (before weight
        // upload) ensures JIT runs with maximum available GPU memory.
        {
            let hidden_dim = model.config.hidden_dim as u32;
            let intermediate_dim = model.config.intermediate_dim as u32;
            let vocab_size = model.config.vocab_size as u32;
            match executor.preload_modules_for_capture(
                num_layers,
                hidden_dim,
                intermediate_dim,
                vocab_size,
            ) {
                Ok(()) => eprintln!(
                    "[GH-129] Early kernel preload: {} modules compiled",
                    executor.module_count()
                ),
                Err(e) => {
                    eprintln!("[GH-129] Early kernel preload failed: {e}");
                },
            }
        }

        // PAR-083: Pre-allocate embedding buffer (hidden_dim f32s) to avoid per-token malloc.
        let embed_buf = vec![0.0f32; model.config.hidden_dim];

        let cuda_model = Self {
            model,
            executor,
            device_name,
            memory_info,
            embed_buf,
            #[cfg(feature = "gpu")]
            prefix_cache: crate::gguf::batch_scheduler::PrefixCache::new(16),
            max_batch_sizing: None,
            free_after_preload: None,
            last_phase_timings: crate::api::PhaseTimings::default(),
        };

        // GH-199 ROOT CAUSE B + PARITY-GATE: preload weights and verify GPU correctness.
        let cuda_model = cuda_model.preload_and_verify()?;

        // PMAT-399 / X1-KVALLOC (#2774): auto-size max_batch AFTER the weights
        // and the FP8/FP16 prefill cache are resident.
        //
        // This ran BEFORE `preload_and_verify`, where `memory_info()` still
        // reports an essentially empty card, so the budget overstated free VRAM
        // by the entire model. Measured on lambda-4090 with
        // `qwen2.5-coder-7b-instruct-q4_k_m` at `--context-length 4096`: 4.4 GB
        // of Q4_K weights plus a 6.7 GB FP8 prefill cache are invisible at the
        // old call site, so `compute_max_batch_for_memory` returned its clamp
        // ceiling of 32 and the scheduler advertised `max_batch=32`. Measured
        // after this move, the same server advertises `max_batch=12` and serves
        // c=4/8/16 at 100% 200s with a 20471 MiB peak of 24564 MiB.
        // `CUDA_MAX_BATCH` is the ADMISSION ceiling,
        // and X1-KVALLOC's allocation floor is the admitted batch -- so an
        // admission ceiling computed against fictional VRAM is still an OOM,
        // one clamp further out. Both ends have to be measured against the
        // same device state.
        //
        // PP-LLAMA-001 §10 / §12 kill criterion: the sizing INPUTS are captured
        // BEFORE `set_var`, and the source is recorded. The env-var transport
        // is what erased the distinction — after load, an operator ceiling and
        // a loader-computed one were the same string in the same variable, so
        // no receipt could say which one the server enforced.
        let mut cuda_model = cuda_model;
        let mut sizing = cuda_model.executor.compute_max_batch_sizing(
            num_layers,
            num_kv_heads,
            head_dim,
            max_seq_len,
        );
        match std::env::var("CUDA_MAX_BATCH")
            .ok()
            .and_then(|v| v.parse::<usize>().ok())
        {
            Some(operator_ceiling) => {
                // The measured inputs are kept: they are what makes an operator
                // ceiling of 4 over a card that would have sized 11 visible.
                sizing.resolved = operator_ceiling;
                sizing.source = crate::cuda::gpu_profile::MAX_BATCH_SOURCE_ENV;
            },
            None => {
                // GH-611: Suppressed — was noisy in non-verbose mode
                // Store for scheduler to pick up
                std::env::set_var("CUDA_MAX_BATCH", sizing.resolved.to_string());
            },
        }
        cuda_model.max_batch_sizing = Some(sizing);

        Ok(cuda_model)
    }

    /// GH-129: Free CPU projection weight copies after GPU preload.
    ///
    /// On Jetson (unified memory), the CPU `Vec<u8>` copies are redundant after
    /// `preload_weights_gpu()` uploads them to GPU buffers. Frees ~1 GB for 1.5B.
    pub fn free_cpu_weights(&mut self) {
        let mut freed = 0usize;
        for layer in &mut self.model.layers {
            freed += layer.qkv_weight.data_bytes();
            freed += layer.attn_output_weight.data.len();
            freed += layer.ffn_up_weight.data.len();
            freed += layer.ffn_down_weight.data.len();
            if let Some(ref gate) = layer.ffn_gate_weight {
                freed += gate.data.len();
            }
            layer.free_projection_weights();
        }
        freed += self.model.lm_head_weight.data.len();
        self.model.lm_head_weight.data = Vec::new();
        eprintln!(
            "[GH-129] Freed {:.1} MB CPU weight copies (GPU-resident path active)",
            freed as f64 / (1024.0 * 1024.0)
        );
    }

    /// Check if CUDA is available
    #[must_use]
    pub fn is_available() -> bool {
        crate::cuda::CudaExecutor::is_available()
    }

    /// Get number of CUDA devices
    #[must_use]
    pub fn num_devices() -> usize {
        crate::cuda::CudaExecutor::num_devices()
    }

    /// Get GPU device name
    #[must_use]
    pub fn device_name(&self) -> &str {
        &self.device_name
    }

    /// Get GPU memory info (free, total) in bytes.
    ///
    /// NOTE the provenance: this snapshot is taken BEFORE the KV cache, the
    /// weights and the FP8/FP16 prefill cache are uploaded. `vram_report()`
    /// labels it `free_at_load_bytes` for that reason and reports the
    /// post-preload measurement beside it.
    #[must_use]
    pub fn memory_info(&self) -> (usize, usize) {
        self.memory_info
    }

    /// PP-LLAMA-001 §5.2: the CUDA executor, for reporting its resolved
    /// configuration under a READ lock.
    #[must_use]
    pub fn executor(&self) -> &crate::cuda::CudaExecutor {
        &self.executor
    }

    /// §9 #1: the prefill path this model will run, resolved once at profile
    /// detection and shared with the engine and the multi-prompt guard.
    #[must_use]
    pub fn prefill_path(&self) -> crate::cuda::gpu_profile::PrefillPathChoice {
        self.executor.gpu_profile().prefill_path()
    }

    /// §10: the `max_batch` decision and every input it was made from.
    #[must_use]
    pub fn max_batch_sizing(&self) -> Option<crate::cuda::gpu_profile::MaxBatchSizing> {
        self.max_batch_sizing
    }

    /// §3: take the phase split of the request that just finished, clearing it.
    ///
    /// TAKING rather than reading is the poka-yoke: the caller holds the same
    /// exclusive lock the request ran under, so what it takes belongs to that
    /// request — and a later request that skipped prefill (prefix-cache hit)
    /// finds the slot empty instead of inheriting its predecessor's numbers.
    pub fn take_phase_timings(&mut self) -> crate::api::PhaseTimings {
        std::mem::take(&mut self.last_phase_timings)
    }

    /// §5.2 / §9 #7: what this process can honestly say about its VRAM.
    ///
    /// Every field states its provenance — see
    /// [`VramReport`](crate::cuda::gpu_profile::VramReport). Nothing here is
    /// estimated: a driver query that fails leaves the corresponding field
    /// `None`.
    #[must_use]
    pub fn vram_report(&self) -> crate::cuda::gpu_profile::VramReport {
        let (free_at_load, total_at_load) = self.memory_info;
        let sample = self.executor.sample_vram_used();
        let (free_now, total_now, used_peak) = match sample {
            Some((free, total, peak)) => (Some(free), total, Some(peak)),
            None => (None, total_at_load, self.executor.vram_used_peak()),
        };
        let num_layers = self.model.config.num_layers;
        let kv_per_slot_bytes = self.executor.batched_kv_bytes_per_slot(num_layers);
        let kv_slots_allocated = self.executor.batched_kv_allocated_slots();
        let kv_single_seq_bytes = self.executor.kv_single_sequence_bytes(num_layers);
        crate::cuda::gpu_profile::VramReport {
            device_name: self.device_name.clone(),
            total_bytes: if total_now > 0 {
                total_now
            } else {
                total_at_load
            },
            free_at_load_bytes: free_at_load,
            free_after_preload_bytes: self.free_after_preload,
            preload_delta_bytes: self
                .free_after_preload
                .map(|after| free_at_load.saturating_sub(after)),
            free_now_bytes: free_now,
            used_now_bytes: free_now.map(|free| total_now.saturating_sub(free)),
            used_peak_bytes: used_peak,
            recorded_alloc_peak_bytes: self.executor.pool_stats().peak_usage,
            kv_single_seq_bytes,
            kv_per_slot_bytes,
            kv_slots_allocated,
            kv_slots_max: crate::cuda::CudaExecutor::MAX_BATCH_CLAMP.1,
            kv_bytes_reserved: kv_single_seq_bytes
                .saturating_add(kv_per_slot_bytes.saturating_mul(kv_slots_allocated)),
            kv_blocks_total: None,
            kv_layout: crate::cuda::gpu_profile::KV_LAYOUT,
        }
    }

    /// Get VRAM usage in MB
    #[must_use]
    pub fn vram_mb(&self) -> u64 {
        (self.memory_info.1 / (1024 * 1024)) as u64
    }

    // ========================================================================
    // PAR-073: BrickProfiler API for per-brick timing
    // ========================================================================

    /// Enable per-brick profiling for real timing measurements.
    ///
    /// When enabled, each brick operation is timed individually using
    /// `std::time::Instant` with CUDA sync for accurate GPU timing.
    pub fn enable_profiling(&mut self) {
        self.executor.enable_profiling();
    }

    /// Disable per-brick profiling (default state).
    pub fn disable_profiling(&mut self) {
        self.executor.disable_profiling();
    }

    /// Check if profiling is enabled.
    #[must_use]
    pub fn is_profiling_enabled(&self) -> bool {
        self.executor.is_profiling_enabled()
    }

    /// Get the brick profiler for reading statistics.
    #[must_use]
    pub fn profiler(&self) -> &trueno::BrickProfiler {
        self.executor.profiler()
    }

    /// Reset profiler statistics.
    pub fn reset_profiler(&mut self) {
        self.executor.reset_profiler();
    }

    /// Get profiler summary report.
    #[must_use]
    pub fn profiler_summary(&self) -> String {
        self.executor.profiler_summary()
    }

    /// Get reference to inner model
    #[must_use]
    pub fn model(&self) -> &OwnedQuantizedModel {
        &self.model
    }

    /// Consume CUDA wrapper and return inner model (for CPU fallback)
    #[must_use]
    pub fn into_model(self) -> OwnedQuantizedModel {
        self.model
    }

    /// PAR-111: Get mutable reference to CUDA executor
    ///
    /// Allows direct access for batched forward path and workspace initialization.
    #[must_use]
    pub fn executor_mut(&mut self) -> &mut crate::cuda::CudaExecutor {
        &mut self.executor
    }

    /// Synchronize CUDA stream (wait for all GPU operations to complete)
    pub fn synchronize(&self) -> Result<()> {
        self.executor
            .synchronize()
            .map_err(|e| RealizarError::UnsupportedOperation {
                operation: "CudaExecutor::synchronize".to_string(),
                reason: format!("CUDA sync failed: {e}"),
            })
    }
}

// =============================================================================
// PARITY GATE: Load-time mathematical proof of GPU/CPU equivalence
// =============================================================================
//
// Toyota Way: Jidoka (自働化) — stop-the-line on defect.
//
// Just as build.rs refuses to compile if ALG-001 through ALG-009 fail,
// this gate refuses to construct `OwnedQuantizedModelCuda` if GPU and CPU
// compute different functions.
//
// An `OwnedQuantizedModelCuda` that passes this gate is PROVEN to produce
// the same output as CPU. One that fails CANNOT be constructed.
//
// Contract: layer-parity-v1.yaml
// Tolerance: cosine_similarity ≥ 0.98 on first-token logits

/// Minimum cosine similarity for parity gate to pass.
/// 0.98 accommodates DP4A integer dot-product kernels which use int8 arithmetic
/// with different rounding than CPU float dequantization. On Qwen 1.5B Q4K:
///   - Non-DP4A (float GEMV): cosine ~0.9999
///   - DP4A (int8 dot product): cosine ~0.9887
/// Real bugs (corrupted weights, wrong kernel) produce cosine < 0.5.
const PARITY_GATE_COSINE_MIN: f32 = 0.98;

include!("mod_parity_gate.rs");