bevy_spark 0.2.0

//! Render-side glue: pipeline, bind groups, queue, draw command.
//!
//! No mesh asset is involved — the shader emits a quad from `vertex_index`
//! and instancing is the per-splat dimension.

#![allow(clippy::too_many_arguments, clippy::type_complexity)]

pub mod sort;

use std::collections::HashMap;
use std::sync::{
    Arc, Mutex, Once,
    atomic::{AtomicU64, Ordering},
};
use std::time::Instant;

use bevy::asset::{AssetEvent, AssetId, AssetServer, Handle, embedded_asset};
use bevy::camera::visibility::{self, VisibilityClass};
use bevy::core_pipeline::core_3d::Transparent3d;
use bevy::ecs::query::QueryItem;
use bevy::ecs::system::SystemParamItem;
use bevy::ecs::system::lifetimeless::Read;
use bevy::math::{Mat4, Vec3};
use bevy::prelude::*;
use bevy::render::extract_component::{ExtractComponent, ExtractComponentPlugin};
use bevy::render::extract_resource::{ExtractResource, ExtractResourcePlugin};
use bevy::render::render_phase::{
    AddRenderCommand, DrawFunctions, PhaseItemExtraIndex, RenderCommand, RenderCommandResult,
    SetItemPipeline, TrackedRenderPass, ViewSortedRenderPhases,
};
use bevy::render::render_resource::*;
use bevy::render::renderer::{RenderDevice, RenderQueue};
use bevy::render::settings::WgpuLimits;
use bevy::render::sync_world::MainEntity;
use bevy::render::view::{ExtractedView, ViewUniformOffset, ViewUniforms};
use bevy::render::{Render, RenderApp, RenderStartup, RenderSystems};
use bevy::shader::Shader;
use bevy::tasks::{AsyncComputeTaskPool, Task, block_on, futures_lite::future};
use bevy::transform::components::GlobalTransform;

use bytemuck::{Pod, Zeroable};

use crate::splats::{GpuSplat, GpuSplatSh, SplatCoordinateConvention, Splats};

const SHADER_PATH: &str = "embedded://bevy_spark/render/splat.wgsl";
static NEXT_SPLAT_UPLOAD_ID: AtomicU64 = AtomicU64::new(1);

const TEXTURE_BACKEND_TEXELS_PER_ROW: u32 = 2048;
const RGBA32UI_TEXEL_BYTES: usize = 16;
const R32UI_TEXEL_BYTES: usize = 4;
const SORT_FORCED_DRIFT_METERS: f32 = 0.75;
const SORT_FORCED_DIR_DOT: f32 = 0.95;
const SORT_STABLE_DRIFT_METERS: f32 = 0.15;
const SORT_STABLE_DIR_DOT: f32 = 0.995;

/// Marker + handle for a splat cloud entity. Spawn alongside [`Transform`].
#[derive(Component, Clone, Default)]
#[require(Transform, Visibility, VisibilityClass)]
#[component(on_add = visibility::add_visibility_class::<SplatCloud>)]
pub struct SplatCloud {
    pub handle: Handle<Splats>,
}

#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub enum SplatFalloffProfile {
    SparkJs,
    EdgeNormalized,
}

impl SplatFalloffProfile {
    fn shader_code(self) -> u32 {
        match self {
            SplatFalloffProfile::SparkJs => 0,
            SplatFalloffProfile::EdgeNormalized => 1,
        }
    }
}

#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub enum SplatHighAlphaProfile {
    SparkJs,
    Bounded,
}

impl SplatHighAlphaProfile {
    fn shader_code(self) -> u32 {
        match self {
            SplatHighAlphaProfile::SparkJs => 0,
            SplatHighAlphaProfile::Bounded => 1,
        }
    }
}

#[derive(Clone, Copy, Debug, PartialEq)]
pub struct SplatQualitySettings {
    pub max_stddev: f32,
    pub min_alpha: f32,
    pub min_pixel_radius: f32,
    pub max_pixel_radius: f32,
    pub falloff_profile: SplatFalloffProfile,
    pub high_alpha_profile: SplatHighAlphaProfile,
}

impl Default for SplatQualitySettings {
    fn default() -> Self {
        Self {
            max_stddev: 8.0_f32.sqrt(),
            min_alpha: 0.5 / 255.0,
            min_pixel_radius: 0.0,
            max_pixel_radius: 512.0,
            falloff_profile: SplatFalloffProfile::SparkJs,
            high_alpha_profile: SplatHighAlphaProfile::SparkJs,
        }
    }
}

/// Typed LOD controls for SPZ clouds.
#[derive(Clone, Copy, Debug, PartialEq)]
pub struct SplatLodSettings {
    pub enabled: bool,
    pub pixel_radius: f32,
    pub center_density_scale: f32,
    pub peripheral_density_scale: f32,
    pub behind_density_scale: f32,
    pub debug_counters: bool,
}

impl Default for SplatLodSettings {
    fn default() -> Self {
        Self {
            enabled: true,
            pixel_radius: 1.0,
            center_density_scale: 1.0,
            peripheral_density_scale: 0.5,
            behind_density_scale: 0.25,
            debug_counters: false,
        }
    }
}

#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub enum SplatSortMode {
    Radial,
    Depth,
}

#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub enum SplatCpuCullMode {
    Radius,
    Center,
}

#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub enum SplatSortBackend {
    Cpu,
    ExperimentalGpu,
    None,
}

#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub struct SplatSortSettings {
    pub backend: SplatSortBackend,
    pub mode: SplatSortMode,
    pub cpu_cull_mode: SplatCpuCullMode,
}

impl Default for SplatSortSettings {
    fn default() -> Self {
        Self {
            backend: SplatSortBackend::Cpu,
            mode: SplatSortMode::Radial,
            cpu_cull_mode: SplatCpuCullMode::Radius,
        }
    }
}

#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub enum SplatMultiCloudMode {
    SingleCloudExact,
    MultiCloudApproximate,
}

#[derive(Clone, Copy, Debug, Default, Eq, PartialEq)]
pub enum SparkQualityPreset {
    Reference,
    #[default]
    Balanced,
    Performance,
    Capture,
}

#[derive(Resource, Clone, Copy, Debug, ExtractResource)]
pub struct SparkSettings {
    pub quality_preset: SparkQualityPreset,
    pub quality: SplatQualitySettings,
    pub sort: SplatSortSettings,
    pub lod: SplatLodSettings,
    pub multi_cloud_mode: SplatMultiCloudMode,
    pub upload_chunk_bytes: usize,
}

#[derive(Resource, Clone, Copy, Debug, Default, Eq, PartialEq)]
pub enum SparkConfigSource {
    ResourceOnly,
    #[default]
    ResourceWithEnvOverrides,
}

impl Default for SparkSettings {
    fn default() -> Self {
        SparkSettings::from_preset(SparkQualityPreset::Balanced)
    }
}

impl SparkSettings {
    pub fn from_preset(preset: SparkQualityPreset) -> Self {
        let mut settings = SparkSettings {
            quality_preset: preset,
            quality: SplatQualitySettings::default(),
            sort: SplatSortSettings::default(),
            lod: SplatLodSettings::default(),
            multi_cloud_mode: SplatMultiCloudMode::MultiCloudApproximate,
            upload_chunk_bytes: 8 * 1024 * 1024,
        };
        settings.apply_preset(preset);
        settings
    }

    pub fn apply_preset(&mut self, preset: SparkQualityPreset) {
        *self = match preset {
            SparkQualityPreset::Reference => SparkSettings {
                quality_preset: preset,
                quality: SplatQualitySettings::default(),
                sort: SplatSortSettings::default(),
                lod: SplatLodSettings::default(),
                multi_cloud_mode: SplatMultiCloudMode::MultiCloudApproximate,
                upload_chunk_bytes: self.upload_chunk_bytes,
            },
            SparkQualityPreset::Balanced => SparkSettings {
                quality_preset: preset,
                quality: SplatQualitySettings {
                    max_stddev: 2.0,
                    min_pixel_radius: 0.25,
                    max_pixel_radius: 256.0,
                    falloff_profile: SplatFalloffProfile::EdgeNormalized,
                    high_alpha_profile: SplatHighAlphaProfile::Bounded,
                    ..Default::default()
                },
                sort: SplatSortSettings::default(),
                lod: SplatLodSettings::default(),
                multi_cloud_mode: SplatMultiCloudMode::MultiCloudApproximate,
                upload_chunk_bytes: self.upload_chunk_bytes,
            },
            SparkQualityPreset::Performance => SparkSettings {
                quality_preset: preset,
                quality: SplatQualitySettings {
                    max_stddev: 2.0,
                    min_alpha: 0.015,
                    min_pixel_radius: 0.25,
                    max_pixel_radius: 128.0,
                    falloff_profile: SplatFalloffProfile::EdgeNormalized,
                    high_alpha_profile: SplatHighAlphaProfile::Bounded,
                },
                sort: SplatSortSettings::default(),
                lod: SplatLodSettings {
                    pixel_radius: 2.0,
                    center_density_scale: 0.75,
                    peripheral_density_scale: 0.35,
                    behind_density_scale: 0.1,
                    ..Default::default()
                },
                multi_cloud_mode: SplatMultiCloudMode::MultiCloudApproximate,
                upload_chunk_bytes: self.upload_chunk_bytes,
            },
            SparkQualityPreset::Capture => SparkSettings {
                quality_preset: preset,
                quality: SplatQualitySettings::default(),
                sort: SplatSortSettings::default(),
                lod: SplatLodSettings {
                    enabled: false,
                    debug_counters: false,
                    ..Default::default()
                },
                multi_cloud_mode: SplatMultiCloudMode::SingleCloudExact,
                upload_chunk_bytes: self.upload_chunk_bytes,
            },
        };
    }

    pub fn with_env_overrides(mut self) -> Self {
        self.apply_env_overrides();
        self
    }

    pub fn apply_env_overrides(&mut self) {
        if let Some(value) = env_lower("BEVY_SPARK_QUALITY_PRESET") {
            let preset = match value.as_str() {
                "balanced" | "desktop" => Some(SparkQualityPreset::Balanced),
                "performance" | "perf" => Some(SparkQualityPreset::Performance),
                "capture" | "deterministic" | "screenshot" => Some(SparkQualityPreset::Capture),
                "reference" | "ref" => Some(SparkQualityPreset::Reference),
                _ => None,
            };
            if let Some(preset) = preset {
                self.apply_preset(preset);
            }
        }
        self.quality.max_stddev = env_f32("BEVY_SPARK_MAX_STDDEV", self.quality.max_stddev);
        self.quality.min_alpha = env_f32("BEVY_SPARK_MIN_ALPHA", self.quality.min_alpha);
        self.quality.min_pixel_radius =
            env_f32("BEVY_SPARK_MIN_PIXEL_RADIUS", self.quality.min_pixel_radius);
        self.quality.max_pixel_radius =
            env_f32("BEVY_SPARK_MAX_PIXEL_RADIUS", self.quality.max_pixel_radius);

        if let Some(value) = env_lower("BEVY_SPARK_FALLOFF_PROFILE") {
            self.quality.falloff_profile = match value.as_str() {
                "edge" | "edge-normalized" | "edge_normalized" | "normalized" | "performance" => {
                    SplatFalloffProfile::EdgeNormalized
                }
                _ => SplatFalloffProfile::SparkJs,
            };
        }
        if let Some(value) = env_lower("BEVY_SPARK_HIGH_ALPHA_PROFILE") {
            self.quality.high_alpha_profile = match value.as_str() {
                "bounded" | "legacy" | "performance" => SplatHighAlphaProfile::Bounded,
                _ => SplatHighAlphaProfile::SparkJs,
            };
        }
        if let Some(value) = env_lower("BEVY_SPARK_LOD") {
            self.lod.enabled = !matches!(
                value.as_str(),
                "0" | "off" | "false" | "full" | "full-quality" | "full_quality"
            );
        }
        self.lod.pixel_radius = env_f32("BEVY_SPARK_LOD_PIXEL_RADIUS", self.lod.pixel_radius);
        self.lod.center_density_scale = env_f32(
            "BEVY_SPARK_LOD_CENTER_DENSITY",
            self.lod.center_density_scale,
        );
        self.lod.peripheral_density_scale = env_f32(
            "BEVY_SPARK_LOD_PERIPHERAL_DENSITY",
            self.lod.peripheral_density_scale,
        );
        self.lod.behind_density_scale = env_f32(
            "BEVY_SPARK_LOD_BEHIND_DENSITY",
            self.lod.behind_density_scale,
        );
        self.lod.debug_counters = env_bool("BEVY_SPARK_LOD_DEBUG", self.lod.debug_counters);

        if let Some(value) = env_lower("BEVY_SPARK_SORT_MODE") {
            self.sort.mode = match value.as_str() {
                "depth" | "z" | "z-depth" | "z_depth" => SplatSortMode::Depth,
                _ => SplatSortMode::Radial,
            };
        }
        if env_bool("BEVY_SPARK_NO_SORT", false) {
            self.sort.backend = SplatSortBackend::None;
        }
        if env_bool("BEVY_SPARK_EXPERIMENTAL_GPU_SORT", false) {
            self.sort.backend = SplatSortBackend::ExperimentalGpu;
        }
        if let Some(value) = env_lower("BEVY_SPARK_CPU_CULL_MODE") {
            self.sort.cpu_cull_mode = match value.as_str() {
                "center" | "center-only" | "center_only" | "fast" => SplatCpuCullMode::Center,
                _ => SplatCpuCullMode::Radius,
            };
        }
        if let Some(value) =
            env_lower("BEVY_SPARK_MULTI_CLOUD_MODE").or_else(|| env_lower("BEVY_SPARK_MULTI_CLOUD"))
        {
            self.multi_cloud_mode = match value.as_str() {
                "single" | "single-exact" | "single_exact" | "exact" => {
                    SplatMultiCloudMode::SingleCloudExact
                }
                _ => SplatMultiCloudMode::MultiCloudApproximate,
            };
        }
        self.upload_chunk_bytes =
            env_usize("BEVY_SPARK_UPLOAD_CHUNK_BYTES", self.upload_chunk_bytes).max(4096);
    }
}

#[derive(Clone, Copy, Debug)]
pub struct SparkDiagnosticsSnapshot {
    pub total_splats: u64,
    pub visible_splats: u64,
    pub lod_selected_splats: u64,
    pub last_sort_time_secs: f32,
    pub last_upload_time_secs: f32,
    pub upload_progress: f32,
    pub quality_preset: SparkQualityPreset,
}

impl Default for SparkDiagnosticsSnapshot {
    fn default() -> Self {
        Self {
            total_splats: 0,
            visible_splats: 0,
            lod_selected_splats: 0,
            last_sort_time_secs: 0.0,
            last_upload_time_secs: 0.0,
            upload_progress: 1.0,
            quality_preset: SparkQualityPreset::default(),
        }
    }
}

#[derive(Resource, Clone, Default, ExtractResource)]
pub struct SparkDiagnostics {
    inner: Arc<Mutex<SparkDiagnosticsSnapshot>>,
}

impl SparkDiagnostics {
    pub fn snapshot(&self) -> SparkDiagnosticsSnapshot {
        self.inner
            .lock()
            .map(|snapshot| *snapshot)
            .unwrap_or_default()
    }

    fn update(&self, update: impl FnOnce(&mut SparkDiagnosticsSnapshot)) {
        if let Ok(mut snapshot) = self.inner.lock() {
            update(&mut snapshot);
        }
    }
}

#[derive(Component, Clone, Copy, Debug, Default)]
pub struct SplatCloudSettings {
    pub quality: Option<SplatQualitySettings>,
    pub sort: Option<SplatSortSettings>,
    pub lod: Option<SplatLodSettings>,
}

#[derive(Component, Default)]
pub struct SplatSortState {
    pub last_view_pos: Vec3,
    pub last_view_dir: Vec3,
    pub initialized: bool,
    pub last_sort_duration_secs: f32,
    pub frames_since_sort: u32,
    pub sort_started_at: Option<Instant>,
}

/// Per-cloud cache of splat centres in main-world memory, shared by `Arc`
/// with background sort tasks so they don't have to clone 24 bytes × N.
#[derive(Component, Clone)]
pub struct SortCenters(pub Arc<Vec<[f32; 3]>>);

/// Per-splat local-space radius estimate used for LOD screen-coverage tests.
#[derive(Component, Clone)]
pub struct SortRadii(pub Arc<Vec<f32>>);

#[derive(Component, Clone, Copy, Debug)]
pub struct SplatCloudBounds {
    pub min: Vec3,
    pub max: Vec3,
}

#[derive(Component, Clone)]
pub struct SplatLodTree {
    pub child_counts: Arc<Vec<u16>>,
    pub child_starts: Arc<Vec<u32>>,
}

impl SplatLodTree {
    fn has_lod(&self) -> bool {
        !self.child_counts.is_empty() && self.child_counts.len() == self.child_starts.len()
    }
}

/// Render-world view key used by draw commands to select view-specific cloud
/// resources.
#[derive(Component)]
pub struct SplatViewEntity(pub Entity);

struct SplatPerViewEntry {
    upload_id: u64,
    indices: SplatIndexStorage,
    uniforms_buffer: Buffer,
    bind_group: Option<BindGroup>,
    bind_upload_id: u64,
    visible_count: u32,
    lod_selected_count: u32,
    lod_total_count: u32,
    state: SplatSortState,
    pending: Option<Task<CpuSortResult>>,
}

#[derive(Component, Default)]
pub struct SplatPerViewResources {
    entries: HashMap<Entity, SplatPerViewEntry>,
}

struct CpuSortResult {
    indices: Vec<u32>,
    lod_selected_count: u32,
    lod_total_count: u32,
}

#[derive(Clone, Copy, Debug, Eq, PartialEq, Hash)]
pub enum SplatStorageBackend {
    StorageBuffers,
    Textures,
}

#[derive(Clone)]
pub struct SplatTextureBinding {
    pub texture: Texture,
    pub view: TextureView,
}

#[allow(clippy::large_enum_variant)]
#[derive(Clone)]
pub enum SplatCloudGpuStorage {
    StorageBuffers {
        splats_buffer: Buffer,
        sh_buffer: Buffer,
        indices_buffer: Buffer,
    },
    Textures {
        splats: SplatTextureBinding,
        sh: SplatTextureBinding,
        indices: SplatTextureBinding,
        texture_width: u32,
    },
}

#[derive(Clone)]
enum SplatIndexStorage {
    Buffer(Buffer),
    Texture(SplatTextureBinding),
}

#[derive(Component, Clone)]
pub struct SplatCloudGpu {
    pub asset_id: AssetId<Splats>,
    pub upload_id: u64,
    pub storage: SplatCloudGpuStorage,
    pub uniforms_buffer: Buffer,
    pub num_splats: u32,
    pub anti_aliased: u32,
    pub sh_degree: u32,
    pub coordinate_convention: SplatCoordinateConvention,
    /// Visible (frustum-culled) splat count; instance-count for the draw call.
    /// Equals `num_splats` until the first sort completes.
    pub visible_count: u32,
}

#[derive(Component, Clone, Copy, Debug, Default)]
pub struct ExtractedSplatCloudSettings(pub SplatCloudSettings);

#[derive(Component)]
struct PendingSplatGpuUpload {
    asset_id: AssetId<Splats>,
    upload_id: u64,
    storage: SplatCloudGpuStorage,
    uniforms_buffer: Buffer,
    splats_bytes: Vec<u8>,
    sh_bytes: Vec<u8>,
    indices_bytes: Vec<u8>,
    splats_uploaded: usize,
    sh_uploaded: usize,
    indices_uploaded: usize,
    num_splats: u32,
    anti_aliased: u32,
    sh_degree: u32,
    coordinate_convention: SplatCoordinateConvention,
    centers: Arc<Vec<[f32; 3]>>,
    radii: Arc<Vec<f32>>,
    bounds: SplatCloudBounds,
    lod_child_counts: Arc<Vec<u16>>,
    lod_child_starts: Arc<Vec<u32>>,
    last_logged_percent: u32,
    upload_started_at: Instant,
}

impl PendingSplatGpuUpload {
    fn total_bytes(&self) -> usize {
        self.splats_bytes.len() + self.sh_bytes.len() + self.indices_bytes.len()
    }

    fn uploaded_bytes(&self) -> usize {
        self.splats_uploaded + self.sh_uploaded + self.indices_uploaded
    }

    fn is_complete(&self) -> bool {
        self.splats_uploaded == self.splats_bytes.len()
            && self.sh_uploaded == self.sh_bytes.len()
            && self.indices_uploaded == self.indices_bytes.len()
    }
}

impl ExtractComponent for SplatCloudGpu {
    type QueryData = (
        &'static SplatCloudGpu,
        &'static GlobalTransform,
        &'static SortCenters,
        &'static SortRadii,
        &'static SplatCloudBounds,
        &'static SplatLodTree,
        Option<&'static SplatCoordinateConvention>,
        Option<&'static SplatCloudSettings>,
    );
    type QueryFilter = ();
    type Out = (
        SplatCloudGpu,
        ExtractedSplatTransform,
        SortCenters,
        SortRadii,
        SplatCloudBounds,
        SplatLodTree,
        ExtractedSplatCloudSettings,
    );

    fn extract_component(item: QueryItem<'_, '_, Self::QueryData>) -> Option<Self::Out> {
        let (gpu, xf, centers, radii, bounds, lod_tree, convention, settings) = item;
        let source_xf = convention
            .copied()
            .unwrap_or(gpu.coordinate_convention)
            .local_from_splat();
        Some((
            gpu.clone(),
            ExtractedSplatTransform(xf.to_matrix() * source_xf),
            centers.clone(),
            radii.clone(),
            *bounds,
            lod_tree.clone(),
            ExtractedSplatCloudSettings(settings.copied().unwrap_or_default()),
        ))
    }
}

#[derive(Component, Clone)]
pub struct ExtractedSplatTransform(pub Mat4);

#[repr(C, align(16))]
#[derive(Copy, Clone, Pod, Zeroable, Debug)]
pub struct CloudUniforms {
    pub model: [[f32; 4]; 4],    // 0..64
    pub num_splats: u32,         // 64..68
    pub anti_aliased: u32,       // 68..72
    pub sh_degree: u32,          // 72..76
    pub _pad0: u32,              // 76..80
    pub render_size: [f32; 2],   // 80..88
    pub max_stddev: f32,         // 88..92
    pub min_alpha: f32,          // 92..96
    pub min_pixel_radius: f32,   // 96..100
    pub max_pixel_radius: f32,   // 100..104
    pub falloff_profile: u32,    // 104..108
    pub high_alpha_profile: u32, // 108..112
    pub texture_width: u32,      // 112..116
    pub _pad1: [u32; 3],         // 116..128
}

const _CHECK: [(); 128] = [(); core::mem::size_of::<CloudUniforms>()];

fn env_f32(name: &str, default: f32) -> f32 {
    std::env::var(name)
        .ok()
        .and_then(|value| value.parse().ok())
        .unwrap_or(default)
}

fn env_usize(name: &str, default: usize) -> usize {
    std::env::var(name)
        .ok()
        .and_then(|value| value.parse().ok())
        .unwrap_or(default)
}

fn env_bool(name: &str, default: bool) -> bool {
    std::env::var(name)
        .ok()
        .map(|value| {
            matches!(
                value.to_ascii_lowercase().as_str(),
                "1" | "true" | "yes" | "on"
            )
        })
        .unwrap_or(default)
}

fn env_lower(name: &str) -> Option<String> {
    std::env::var(name)
        .ok()
        .map(|value| value.to_ascii_lowercase())
}

pub struct SparkRenderPlugin;

fn apply_spark_settings_env_overrides_once(
    mut settings: ResMut<SparkSettings>,
    source: Res<SparkConfigSource>,
    mut applied: Local<bool>,
) {
    if *applied {
        return;
    }
    if *source == SparkConfigSource::ResourceWithEnvOverrides {
        settings.apply_env_overrides();
    }
    *applied = true;
}

impl Plugin for SparkRenderPlugin {
    fn build(&self, app: &mut App) {
        embedded_asset!(app, "splat.wgsl");
        sort::register_sort_assets(app);

        app.init_resource::<SparkSettings>()
            .init_resource::<SparkConfigSource>()
            .init_resource::<SparkDiagnostics>()
            .add_plugins((
                ExtractComponentPlugin::<SplatCloudGpu>::default(),
                ExtractResourcePlugin::<SparkSettings>::default(),
                ExtractResourcePlugin::<SparkDiagnostics>::default(),
            ))
            .add_systems(Startup, apply_spark_settings_env_overrides_once)
            .add_systems(
                Update,
                (
                    invalidate_splat_gpu_on_asset_events,
                    prepare_cloud_gpu,
                    progress_splat_gpu_uploads,
                )
                    .chain(),
            );

        let Some(render_app) = app.get_sub_app_mut(RenderApp) else {
            return;
        };

        render_app
            .init_resource::<SparkSettings>()
            .init_resource::<SparkDiagnostics>()
            .init_resource::<SpecializedRenderPipelines<SplatPipeline>>()
            .add_render_command::<Transparent3d, DrawSplatCloud>()
            .add_systems(RenderStartup, init_splat_pipeline)
            .add_systems(RenderStartup, sort::init_sort_pipeline)
            .add_systems(
                Render,
                (
                    queue_splat_clouds.in_set(RenderSystems::Queue),
                    sort::ensure_sort_buffers.in_set(RenderSystems::PrepareResources),
                    prepare_splat_view_resources.in_set(RenderSystems::PrepareResources),
                    prepare_splat_bind_groups.in_set(RenderSystems::PrepareBindGroups),
                    sort::run_gpu_sort.in_set(RenderSystems::PrepareBindGroups),
                ),
            );
    }
}

// ─── Main-world systems ───────────────────────────────────────────────────

fn clear_splat_gpu_components(commands: &mut Commands, entity: Entity) {
    commands.entity(entity).remove::<(
        SplatCloudGpu,
        PendingSplatGpuUpload,
        SortCenters,
        SortRadii,
        SplatCloudBounds,
        SplatLodTree,
    )>();
}

fn invalidate_splat_gpu_on_asset_events(
    mut commands: Commands,
    mut events: MessageReader<AssetEvent<Splats>>,
    clouds: Query<(Entity, &SplatCloud, Option<&SplatCloudGpu>)>,
) {
    let mut changed_assets = Vec::new();
    for event in events.read() {
        match event {
            AssetEvent::Modified { id } | AssetEvent::Removed { id } => changed_assets.push(*id),
            AssetEvent::Added { .. }
            | AssetEvent::LoadedWithDependencies { .. }
            | AssetEvent::Unused { .. } => {}
        }
    }
    if changed_assets.is_empty() {
        return;
    }

    for (entity, cloud, maybe_gpu) in &clouds {
        if maybe_gpu.is_some() && changed_assets.iter().any(|id| *id == cloud.handle.id()) {
            clear_splat_gpu_components(&mut commands, entity);
        }
    }
}

fn prepare_cloud_gpu(
    mut commands: Commands,
    splat_assets: Res<Assets<Splats>>,
    render_device: Res<RenderDevice>,
    query: Query<(
        Entity,
        &SplatCloud,
        Option<&SplatCloudGpu>,
        Option<&PendingSplatGpuUpload>,
    )>,
) {
    for (entity, cloud, maybe_gpu, maybe_pending) in &query {
        let asset_id = cloud.handle.id();
        if maybe_gpu.is_some_and(|gpu| gpu.asset_id == asset_id) {
            continue;
        }
        if maybe_pending.is_some_and(|pending| pending.asset_id == asset_id) {
            continue;
        }
        let Some(splats) = splat_assets.get(&cloud.handle) else {
            if maybe_gpu.is_some() || maybe_pending.is_some() {
                clear_splat_gpu_components(&mut commands, entity);
            }
            continue;
        };
        if splats.is_empty() {
            if maybe_gpu.is_some() || maybe_pending.is_some() {
                clear_splat_gpu_components(&mut commands, entity);
            }
            continue;
        }
        let upload_id = NEXT_SPLAT_UPLOAD_ID.fetch_add(1, Ordering::Relaxed);
        let gpu_splats: Vec<GpuSplat> = splats.to_gpu();
        let gpu_sh: Vec<GpuSplatSh> = splats.to_gpu_sh();
        let splats_bytes = bytemuck::cast_slice(&gpu_splats).to_vec();
        let sh_bytes = bytemuck::cast_slice(&gpu_sh).to_vec();
        let initial_indices: Vec<u32> = (0..splats.len() as u32).collect();
        let indices_bytes = bytemuck::cast_slice(&initial_indices).to_vec();
        let storage = match splat_storage_backend(&render_device) {
            SplatStorageBackend::StorageBuffers => {
                if !storage_buffer_fits_device(
                    &render_device,
                    "splat storage buffer",
                    splats_bytes.len() as u64,
                ) || !storage_buffer_fits_device(
                    &render_device,
                    "SH storage buffer",
                    sh_bytes.len() as u64,
                ) || !storage_buffer_fits_device(
                    &render_device,
                    "sorted-index storage buffer",
                    indices_bytes.len() as u64,
                ) {
                    continue;
                }
                let splats_buffer = render_device.create_buffer(&BufferDescriptor {
                    label: Some("spark.splats"),
                    size: splats_bytes.len() as u64,
                    usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
                    mapped_at_creation: false,
                });
                let sh_buffer = render_device.create_buffer(&BufferDescriptor {
                    label: Some("spark.splat_sh"),
                    size: sh_bytes.len() as u64,
                    usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
                    mapped_at_creation: false,
                });
                let indices_buffer = render_device.create_buffer(&BufferDescriptor {
                    label: Some("spark.sorted_indices"),
                    size: indices_bytes.len() as u64,
                    usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
                    mapped_at_creation: false,
                });
                SplatCloudGpuStorage::StorageBuffers {
                    splats_buffer,
                    sh_buffer,
                    indices_buffer,
                }
            }
            SplatStorageBackend::Textures => {
                let texture_width = texture_backend_width(&render_device);
                let Some(splats) = create_uint_texture_binding(
                    &render_device,
                    "spark.splats_texture",
                    TextureFormat::Rgba32Uint,
                    splats_bytes.len() / RGBA32UI_TEXEL_BYTES,
                    texture_width,
                ) else {
                    continue;
                };
                let Some(sh) = create_uint_texture_binding(
                    &render_device,
                    "spark.splat_sh_texture",
                    TextureFormat::Rgba32Uint,
                    sh_bytes.len() / RGBA32UI_TEXEL_BYTES,
                    texture_width,
                ) else {
                    continue;
                };
                let Some(indices) = create_uint_texture_binding(
                    &render_device,
                    "spark.sorted_indices_texture",
                    TextureFormat::R32Uint,
                    indices_bytes.len() / R32UI_TEXEL_BYTES,
                    texture_width,
                ) else {
                    continue;
                };
                SplatCloudGpuStorage::Textures {
                    splats,
                    sh,
                    indices,
                    texture_width,
                }
            }
        };
        let uniforms_buffer = render_device.create_buffer(&BufferDescriptor {
            label: Some("spark.cloud_uniforms"),
            size: core::mem::size_of::<CloudUniforms>() as u64,
            usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });
        let num_splats = splats.len() as u32;
        let mut centers = Vec::with_capacity(splats.splats.len());
        let mut radii = Vec::with_capacity(splats.splats.len());
        let mut bounds_min = Vec3::splat(f32::INFINITY);
        let mut bounds_max = Vec3::splat(f32::NEG_INFINITY);
        for splat in &splats.splats {
            let center = Vec3::from_array(splat.center);
            let radius = splat.scale[0].max(splat.scale[1]).max(splat.scale[2]);
            centers.push(splat.center);
            radii.push(radius);
            let extent = Vec3::splat(radius.max(0.0));
            bounds_min = bounds_min.min(center - extent);
            bounds_max = bounds_max.max(center + extent);
        }
        commands.entity(entity).insert(PendingSplatGpuUpload {
            asset_id,
            upload_id,
            storage,
            uniforms_buffer,
            splats_bytes,
            sh_bytes,
            indices_bytes,
            splats_uploaded: 0,
            sh_uploaded: 0,
            indices_uploaded: 0,
            num_splats,
            anti_aliased: u32::from(splats.anti_aliased),
            sh_degree: splats.sh_degree,
            coordinate_convention: splats.coordinate_convention,
            centers: Arc::new(centers),
            radii: Arc::new(radii),
            bounds: SplatCloudBounds {
                min: bounds_min,
                max: bounds_max,
            },
            lod_child_counts: Arc::new(splats.lod_child_counts.clone()),
            lod_child_starts: Arc::new(splats.lod_child_starts.clone()),
            last_logged_percent: 0,
            upload_started_at: Instant::now(),
        });
    }
}

fn write_staged_buffer(
    render_queue: &RenderQueue,
    buffer: &Buffer,
    bytes: &[u8],
    uploaded: usize,
    budget: &mut usize,
) -> usize {
    if *budget == 0 || uploaded == bytes.len() {
        return uploaded;
    }
    let remaining = bytes.len() - uploaded;
    let chunk = remaining.min(*budget);
    render_queue.write_buffer(
        buffer,
        uploaded as u64,
        &bytes[uploaded..(uploaded + chunk)],
    );
    *budget -= chunk;
    uploaded + chunk
}

fn create_index_storage(
    render_device: &RenderDevice,
    render_queue: &RenderQueue,
    backend: SplatStorageBackend,
    num_splats: u32,
) -> Option<SplatIndexStorage> {
    let initial_indices: Vec<u32> = (0..num_splats).collect();
    let indices_bytes = bytemuck::cast_slice(&initial_indices);
    match backend {
        SplatStorageBackend::StorageBuffers => {
            let indices_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
                label: Some("spark.view_sorted_indices"),
                contents: indices_bytes,
                usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
            });
            Some(SplatIndexStorage::Buffer(indices_buffer))
        }
        SplatStorageBackend::Textures => {
            let texture_width = texture_backend_width(render_device);
            let indices = create_uint_texture_binding(
                render_device,
                "spark.view_sorted_indices_texture",
                TextureFormat::R32Uint,
                indices_bytes.len() / R32UI_TEXEL_BYTES,
                texture_width,
            )?;
            write_texture_upload(
                render_queue,
                &indices.texture,
                indices_bytes,
                R32UI_TEXEL_BYTES,
                texture_width,
            );
            Some(SplatIndexStorage::Texture(indices))
        }
    }
}

fn write_index_storage(
    render_queue: &RenderQueue,
    storage: &SplatIndexStorage,
    indices: &[u32],
    render_device: &RenderDevice,
) {
    match storage {
        SplatIndexStorage::Buffer(buffer) => {
            render_queue.write_buffer(buffer, 0, bytemuck::cast_slice(indices));
        }
        SplatIndexStorage::Texture(texture) => {
            write_texture_upload(
                render_queue,
                &texture.texture,
                bytemuck::cast_slice(indices),
                R32UI_TEXEL_BYTES,
                texture_backend_width(render_device),
            );
        }
    }
}

fn progress_splat_gpu_uploads(
    mut commands: Commands,
    render_queue: Res<RenderQueue>,
    settings: Res<SparkSettings>,
    diagnostics: Res<SparkDiagnostics>,
    mut pending_uploads: Query<(Entity, &mut PendingSplatGpuUpload)>,
) {
    let chunk_budget = settings.upload_chunk_bytes.max(4096);
    let mut min_progress = 1.0_f32;
    let mut saw_upload = false;
    let mut last_completed_upload_time = None;
    for (entity, mut pending) in &mut pending_uploads {
        saw_upload = true;
        let mut budget = chunk_budget;
        let storage = pending.storage.clone();
        match &storage {
            SplatCloudGpuStorage::StorageBuffers {
                splats_buffer,
                sh_buffer,
                indices_buffer,
            } => {
                pending.splats_uploaded = write_staged_buffer(
                    &render_queue,
                    splats_buffer,
                    &pending.splats_bytes,
                    pending.splats_uploaded,
                    &mut budget,
                );
                pending.sh_uploaded = write_staged_buffer(
                    &render_queue,
                    sh_buffer,
                    &pending.sh_bytes,
                    pending.sh_uploaded,
                    &mut budget,
                );
                pending.indices_uploaded = write_staged_buffer(
                    &render_queue,
                    indices_buffer,
                    &pending.indices_bytes,
                    pending.indices_uploaded,
                    &mut budget,
                );
            }
            SplatCloudGpuStorage::Textures {
                splats,
                sh,
                indices,
                texture_width,
            } => {
                if !pending.is_complete() {
                    write_texture_upload(
                        &render_queue,
                        &splats.texture,
                        &pending.splats_bytes,
                        RGBA32UI_TEXEL_BYTES,
                        *texture_width,
                    );
                    write_texture_upload(
                        &render_queue,
                        &sh.texture,
                        &pending.sh_bytes,
                        RGBA32UI_TEXEL_BYTES,
                        *texture_width,
                    );
                    write_texture_upload(
                        &render_queue,
                        &indices.texture,
                        &pending.indices_bytes,
                        R32UI_TEXEL_BYTES,
                        *texture_width,
                    );
                    pending.splats_uploaded = pending.splats_bytes.len();
                    pending.sh_uploaded = pending.sh_bytes.len();
                    pending.indices_uploaded = pending.indices_bytes.len();
                }
            }
        }

        let total = pending.total_bytes().max(1);
        let percent = ((pending.uploaded_bytes() * 100) / total) as u32;
        min_progress = min_progress.min(pending.uploaded_bytes() as f32 / total as f32);
        if percent == 100 || percent >= pending.last_logged_percent.saturating_add(25) {
            pending.last_logged_percent = percent;
            bevy::log::info!(
                "[bevy_spark] staged GPU upload {}% for {} splats",
                percent,
                pending.num_splats
            );
        }

        if pending.is_complete() {
            last_completed_upload_time = Some(pending.upload_started_at.elapsed().as_secs_f32());
            commands
                .entity(entity)
                .insert((
                    SplatCloudGpu {
                        asset_id: pending.asset_id,
                        upload_id: pending.upload_id,
                        storage: pending.storage.clone(),
                        uniforms_buffer: pending.uniforms_buffer.clone(),
                        num_splats: pending.num_splats,
                        anti_aliased: pending.anti_aliased,
                        sh_degree: pending.sh_degree,
                        coordinate_convention: pending.coordinate_convention,
                        visible_count: pending.num_splats,
                    },
                    SortCenters(pending.centers.clone()),
                    SortRadii(pending.radii.clone()),
                    pending.bounds,
                    SplatLodTree {
                        child_counts: pending.lod_child_counts.clone(),
                        child_starts: pending.lod_child_starts.clone(),
                    },
                ))
                .remove::<PendingSplatGpuUpload>();
            bevy::log::info!("[bevy_spark] uploaded {} splats to GPU", pending.num_splats);
        }
    }
    diagnostics.update(|snapshot| {
        snapshot.upload_progress = if saw_upload { min_progress } else { 1.0 };
        if let Some(seconds) = last_completed_upload_time {
            snapshot.last_upload_time_secs = seconds;
        }
        snapshot.quality_preset = settings.quality_preset;
    });
}

pub(crate) fn storage_buffer_fits_device(
    render_device: &RenderDevice,
    label: &str,
    size: u64,
) -> bool {
    let limits = render_device.limits();
    if storage_buffer_fits_limits(&limits, size) {
        return true;
    }
    let max_storage_binding = u64::from(limits.max_storage_buffer_binding_size);
    bevy::log::error!(
        "[bevy_spark] {label} requires {size} bytes, exceeding adapter limits \
         (max_buffer_size={}, max_storage_buffer_binding_size={})",
        limits.max_buffer_size,
        max_storage_binding
    );
    false
}

fn storage_buffer_fits_limits(limits: &WgpuLimits, size: u64) -> bool {
    let max_storage_binding = u64::from(limits.max_storage_buffer_binding_size);
    size <= limits.max_buffer_size && size <= max_storage_binding
}

pub(crate) fn splat_storage_backend(render_device: &RenderDevice) -> SplatStorageBackend {
    let limits = render_device.limits();
    if cfg!(feature = "webgl2-textures") || limits.max_storage_buffers_per_shader_stage < 4 {
        SplatStorageBackend::Textures
    } else {
        SplatStorageBackend::StorageBuffers
    }
}

fn texture_backend_width(render_device: &RenderDevice) -> u32 {
    TEXTURE_BACKEND_TEXELS_PER_ROW
        .min(render_device.limits().max_texture_dimension_2d)
        .max(1)
}

fn texture_backend_extent(
    render_device: &RenderDevice,
    label: &str,
    texel_count: usize,
    texture_width: u32,
) -> Option<Extent3d> {
    let texel_count = texel_count.max(1) as u32;
    let height = texel_count.div_ceil(texture_width).max(1);
    let max_dimension = render_device.limits().max_texture_dimension_2d;
    if texture_width > max_dimension || height > max_dimension {
        bevy::log::error!(
            "[bevy_spark] {label} requires a {}x{} data texture, exceeding adapter max_texture_dimension_2d={}",
            texture_width,
            height,
            max_dimension
        );
        return None;
    }
    Some(Extent3d {
        width: texture_width,
        height,
        depth_or_array_layers: 1,
    })
}

fn create_uint_texture_binding(
    render_device: &RenderDevice,
    label: &'static str,
    format: TextureFormat,
    texel_count: usize,
    texture_width: u32,
) -> Option<SplatTextureBinding> {
    let extent = texture_backend_extent(render_device, label, texel_count, texture_width)?;
    let texture = render_device.create_texture(&TextureDescriptor {
        label: Some(label),
        size: extent,
        mip_level_count: 1,
        sample_count: 1,
        dimension: TextureDimension::D2,
        format,
        usage: TextureUsages::TEXTURE_BINDING | TextureUsages::COPY_DST,
        view_formats: &[],
    });
    let view = texture.create_view(&TextureViewDescriptor::default());
    Some(SplatTextureBinding { texture, view })
}

fn padded_texture_upload_bytes(
    bytes: &[u8],
    texel_size: usize,
    texture_width: u32,
) -> (Vec<u8>, Extent3d, u32) {
    let texel_count = (bytes.len() / texel_size).max(1) as u32;
    let height = texel_count.div_ceil(texture_width).max(1);
    let row_bytes = texture_width as usize * texel_size;
    let mut padded = vec![0u8; row_bytes * height as usize];
    for row in 0..height as usize {
        let src_start = row * row_bytes;
        if src_start >= bytes.len() {
            break;
        }
        let src_end = (src_start + row_bytes).min(bytes.len());
        let dst = &mut padded[src_start..src_start + (src_end - src_start)];
        dst.copy_from_slice(&bytes[src_start..src_end]);
    }
    (
        padded,
        Extent3d {
            width: texture_width,
            height,
            depth_or_array_layers: 1,
        },
        row_bytes as u32,
    )
}

fn write_texture_upload(
    render_queue: &RenderQueue,
    texture: &Texture,
    bytes: &[u8],
    texel_size: usize,
    texture_width: u32,
) {
    let (upload_bytes, extent, row_bytes) =
        padded_texture_upload_bytes(bytes, texel_size, texture_width);
    render_queue.write_texture(
        texture.as_image_copy(),
        &upload_bytes,
        TexelCopyBufferLayout {
            offset: 0,
            bytes_per_row: Some(row_bytes),
            rows_per_image: Some(extent.height),
        },
        extent,
    );
}

fn gpu_compute_sort_supported_limits(limits: &WgpuLimits) -> bool {
    limits.max_compute_workgroup_size_x >= 256
        && limits.max_compute_invocations_per_workgroup >= 256
        && limits.max_storage_buffers_per_shader_stage >= 7
}

pub(crate) fn gpu_compute_sort_supported(render_device: &RenderDevice) -> bool {
    let limits = render_device.limits();
    let texture_backend = splat_storage_backend(render_device) == SplatStorageBackend::Textures;
    let supported = !texture_backend && gpu_compute_sort_supported_limits(&limits);
    if !supported {
        static WARNED: Once = Once::new();
        WARNED.call_once(|| {
            if texture_backend {
                bevy::log::warn!(
                    "[bevy_spark] experimental GPU sort is disabled on the texture data backend; \
                     falling back to CPU sort when enabled"
                );
            } else {
                bevy::log::warn!(
                    "[bevy_spark] experimental GPU sort is unavailable on this adapter \
                     (max_compute_workgroup_size_x={}, max_compute_invocations_per_workgroup={}, \
                      max_storage_buffers_per_shader_stage={}); falling back to CPU sort when enabled",
                    limits.max_compute_workgroup_size_x,
                    limits.max_compute_invocations_per_workgroup,
                    limits.max_storage_buffers_per_shader_stage
                );
            }
        });
    }
    supported
}

pub(crate) fn effective_sort_settings(
    settings: &SparkSettings,
    cloud_settings: Option<&ExtractedSplatCloudSettings>,
) -> SplatSortSettings {
    cloud_settings
        .and_then(|cloud| cloud.0.sort)
        .unwrap_or(settings.sort)
}

fn effective_quality_settings(
    settings: &SparkSettings,
    cloud_settings: Option<&ExtractedSplatCloudSettings>,
) -> SplatQualitySettings {
    cloud_settings
        .and_then(|cloud| cloud.0.quality)
        .unwrap_or(settings.quality)
}

fn effective_lod_settings(
    settings: &SparkSettings,
    cloud_settings: Option<&ExtractedSplatCloudSettings>,
) -> SplatLodSettings {
    cloud_settings
        .and_then(|cloud| cloud.0.lod)
        .unwrap_or(settings.lod)
}

pub(crate) fn experimental_gpu_sort_enabled(sort: SplatSortSettings) -> bool {
    let enabled = sort.backend == SplatSortBackend::ExperimentalGpu;
    if enabled {
        static WARNED: Once = Once::new();
        WARNED.call_once(|| {
            bevy::log::warn!(
                "[bevy_spark] the experimental GPU sorter is enabled. \
                 It is still under verification and is not a supported shipping mode."
            );
        });
    }
    enabled
}

fn cpu_sort_enabled_for_device(render_device: &RenderDevice, sort: SplatSortSettings) -> bool {
    match sort.backend {
        SplatSortBackend::Cpu => true,
        SplatSortBackend::ExperimentalGpu => !gpu_compute_sort_supported(render_device),
        SplatSortBackend::None => false,
    }
}

fn adaptive_sort_interval_frames(last_sort_duration_secs: f32) -> u32 {
    if last_sort_duration_secs >= 0.050 {
        6
    } else if last_sort_duration_secs >= 0.020 {
        3
    } else if last_sort_duration_secs >= 0.008 {
        1
    } else {
        0
    }
}

fn row_dot(row: [f32; 4], c: [f32; 3]) -> f32 {
    row[0] * c[0] + row[1] * c[1] + row[2] * c[2] + row[3]
}

fn row_xyz_len(row: [f32; 4]) -> f32 {
    (row[0] * row[0] + row[1] * row[1] + row[2] * row[2]).sqrt()
}

fn outside_cpu_clip_bounds(
    center: [f32; 3],
    radius: f32,
    r0: [f32; 4],
    r1: [f32; 4],
    r3: [f32; 4],
    limit: f32,
    cull_mode: SplatCpuCullMode,
) -> bool {
    let w = row_dot(r3, center);
    if cull_mode == SplatCpuCullMode::Center {
        if w <= 0.0 {
            return true;
        }
        let lim = limit * w;
        return row_dot(r0, center).abs() > lim || row_dot(r1, center).abs() > lim;
    }

    let radius = radius.max(0.0);
    let w_margin = radius * row_xyz_len(r3);
    if w + w_margin <= 0.0 {
        return true;
    }
    let lim = limit * w;
    let x_margin = radius * (row_xyz_len(r0) + limit * row_xyz_len(r3));
    if row_dot(r0, center).abs() > lim + x_margin {
        return true;
    }
    let y_margin = radius * (row_xyz_len(r1) + limit * row_xyz_len(r3));
    row_dot(r1, center).abs() > lim + y_margin
}

fn projected_splat_pixel_radius(
    radius: f32,
    radius_scale: f32,
    focal_y_pixels: f32,
    orthographic: bool,
    depth: f32,
    extent_scale: f32,
) -> f32 {
    let world_radius = radius.max(0.0) * radius_scale.max(0.0);
    let sigma_pixels = if orthographic {
        world_radius * focal_y_pixels
    } else {
        world_radius * focal_y_pixels / depth.max(1e-3)
    };
    sigma_pixels * extent_scale.max(0.0)
}

fn cloud_bounds_visible_in_clip(
    bounds: SplatCloudBounds,
    model: Mat4,
    clip_from_world: Mat4,
    limit: f32,
) -> bool {
    let clip_from_local = clip_from_world * model;
    let corners = [
        Vec3::new(bounds.min.x, bounds.min.y, bounds.min.z),
        Vec3::new(bounds.max.x, bounds.min.y, bounds.min.z),
        Vec3::new(bounds.min.x, bounds.max.y, bounds.min.z),
        Vec3::new(bounds.max.x, bounds.max.y, bounds.min.z),
        Vec3::new(bounds.min.x, bounds.min.y, bounds.max.z),
        Vec3::new(bounds.max.x, bounds.min.y, bounds.max.z),
        Vec3::new(bounds.min.x, bounds.max.y, bounds.max.z),
        Vec3::new(bounds.max.x, bounds.max.y, bounds.max.z),
    ];

    let mut any_in_front = false;
    let mut any_behind_near = false;
    let mut all_left = true;
    let mut all_right = true;
    let mut all_bottom = true;
    let mut all_top = true;

    for corner in corners {
        let clip = clip_from_local * corner.extend(1.0);
        if clip.w <= 0.0 {
            any_behind_near = true;
            continue;
        }
        any_in_front = true;
        let lim = limit * clip.w;
        all_left &= clip.x < -lim;
        all_right &= clip.x > lim;
        all_bottom &= clip.y < -lim;
        all_top &= clip.y > lim;
    }

    if !any_in_front {
        return false;
    }
    if any_behind_near {
        return true;
    }
    !(all_left || all_right || all_bottom || all_top)
}

fn select_lod_candidates(
    n: usize,
    centers: &[[f32; 3]],
    radii: &[f32],
    child_counts: &[u16],
    child_starts: &[u32],
    view_pos: [f32; 3],
    view_dir: [f32; 3],
    model_rows: [[f32; 4]; 3],
    radius_scale: f32,
    focal_y_pixels: f32,
    orthographic: bool,
    pixel_threshold: f32,
    center_density_scale: f32,
    peripheral_density_scale: f32,
    behind_density_scale: f32,
) -> Vec<u32> {
    debug_assert_eq!(
        child_counts.len(),
        n,
        "LOD child-count table must match splat count"
    );
    debug_assert_eq!(
        child_starts.len(),
        n,
        "LOD child-start table must match splat count"
    );
    debug_assert_eq!(radii.len(), n, "LOD radii table must match splat count");
    debug_assert_eq!(centers.len(), n, "LOD center table must match splat count");
    if child_counts.len() != n || child_starts.len() != n || radii.len() != n || centers.len() != n
    {
        return (0..n as u32).collect();
    }

    let mut selected = Vec::new();
    let mut stack = vec![0usize];
    let mut visited = vec![false; n];

    while let Some(i) = stack.pop() {
        if i >= n || visited[i] {
            continue;
        }
        visited[i] = true;

        let count = child_counts[i] as usize;
        if count == 0 {
            selected.push(i as u32);
            continue;
        }

        let c = centers[i];
        let wx = model_rows[0][0] * c[0]
            + model_rows[0][1] * c[1]
            + model_rows[0][2] * c[2]
            + model_rows[0][3];
        let wy = model_rows[1][0] * c[0]
            + model_rows[1][1] * c[1]
            + model_rows[1][2] * c[2]
            + model_rows[1][3];
        let wz = model_rows[2][0] * c[0]
            + model_rows[2][1] * c[1]
            + model_rows[2][2] * c[2]
            + model_rows[2][3];
        let dx = wx - view_pos[0];
        let dy = wy - view_pos[1];
        let dz = wz - view_pos[2];
        let raw_depth = dx * view_dir[0] + dy * view_dir[1] + dz * view_dir[2];
        let depth = if raw_depth > 0.0 {
            raw_depth.max(1e-3)
        } else {
            raw_depth.abs().max(1.0)
        };
        let distance = (dx * dx + dy * dy + dz * dz).sqrt().max(1e-3);
        let center_alignment = (raw_depth / distance).clamp(0.0, 1.0);
        let foveal_t = center_alignment * center_alignment;
        let density_scale = if raw_depth <= 0.0 {
            behind_density_scale
        } else {
            peripheral_density_scale + (center_density_scale - peripheral_density_scale) * foveal_t
        }
        .max(0.0);
        let radius = radii[i].max(0.0) * radius_scale.max(0.0);
        let pixel_radius = if orthographic {
            radius * focal_y_pixels
        } else {
            radius * focal_y_pixels / depth
        } * density_scale;

        if pixel_radius <= pixel_threshold {
            selected.push(i as u32);
            continue;
        }

        let start = child_starts[i] as usize;
        let end = start.saturating_add(count).min(n);
        if start >= end {
            selected.push(i as u32);
            continue;
        }
        for child in (start..end).rev() {
            if child != i {
                stack.push(child);
            }
        }
    }

    if selected.is_empty() {
        (0..n as u32).collect()
    } else {
        selected
    }
}

fn prepare_splat_view_resources(
    mut commands: Commands,
    render_device: Res<RenderDevice>,
    render_queue: Res<RenderQueue>,
    settings: Res<SparkSettings>,
    diagnostics: Res<SparkDiagnostics>,
    transparent_phases: Res<ViewSortedRenderPhases<Transparent3d>>,
    views: Query<(Entity, &ExtractedView)>,
    mut clouds: Query<(
        Entity,
        &SplatCloudGpu,
        &SortCenters,
        &SortRadii,
        &SplatCloudBounds,
        &SplatLodTree,
        &ExtractedSplatTransform,
        Option<&ExtractedSplatCloudSettings>,
        Option<&mut SplatPerViewResources>,
    )>,
) {
    let active_views: Vec<(Entity, Vec3, Vec3, Mat4, f32, bool)> = views
        .iter()
        .filter(|(_, view)| transparent_phases.get(&view.retained_view_entity).is_some())
        .map(|(view_entity, view)| {
            let view_pos = view.world_from_view.translation();
            let view_dir = view.world_from_view.forward().as_vec3();
            let clip_from_world = view.clip_from_world.unwrap_or_else(|| {
                view.clip_from_view * view.world_from_view.to_matrix().inverse()
            });
            let viewport = view.viewport.zw();
            let focal_y_pixels = view.clip_from_view.y_axis.y.abs() * viewport.y as f32 * 0.5;
            let orthographic = (view.clip_from_view.w_axis.w - 1.0).abs() < 1e-6;
            (
                view_entity,
                view_pos,
                view_dir,
                clip_from_world,
                focal_y_pixels,
                orthographic,
            )
        })
        .collect();
    if active_views.is_empty() {
        return;
    }

    let pool = AsyncComputeTaskPool::get();
    let mut total_splats = 0_u64;
    let mut visible_splats = 0_u64;
    let mut lod_selected_splats = 0_u64;
    let mut last_sort_time_secs = 0.0_f32;

    for (
        entity,
        gpu,
        centers_arc,
        radii_arc,
        bounds,
        lod_tree,
        model_xf,
        cloud_settings,
        maybe_resources,
    ) in &mut clouds
    {
        total_splats = total_splats.saturating_add(u64::from(gpu.num_splats));
        let sort_settings = effective_sort_settings(&settings, cloud_settings);
        let quality = effective_quality_settings(&settings, cloud_settings);
        let lod_settings = effective_lod_settings(&settings, cloud_settings);
        let enable_cpu_sort = cpu_sort_enabled_for_device(&render_device, sort_settings);
        let sort_mode = sort_settings.mode;
        let cull_mode = sort_settings.cpu_cull_mode;
        let min_pixel_radius = quality.min_pixel_radius.max(0.0);
        let extent_scale = quality.max_stddev.max(0.0);

        let Some(mut resources) = maybe_resources else {
            commands
                .entity(entity)
                .insert(SplatPerViewResources::default());
            continue;
        };
        resources.entries.retain(|view_entity, _| {
            active_views
                .iter()
                .any(|(e, _, _, _, _, _)| e == view_entity)
        });

        for (view_entity, view_pos, view_dir, clip_from_world, focal_y_pixels, orthographic) in
            &active_views
        {
            let needs_entry = resources
                .entries
                .get(view_entity)
                .map(|entry| entry.upload_id != gpu.upload_id)
                .unwrap_or(true);
            if needs_entry {
                let Some(indices) = create_index_storage(
                    &render_device,
                    &render_queue,
                    splat_storage_backend(&render_device),
                    gpu.num_splats,
                ) else {
                    continue;
                };
                let uniforms_buffer = render_device.create_buffer(&BufferDescriptor {
                    label: Some("spark.view_cloud_uniforms"),
                    size: core::mem::size_of::<CloudUniforms>() as u64,
                    usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST,
                    mapped_at_creation: false,
                });
                resources.entries.insert(
                    *view_entity,
                    SplatPerViewEntry {
                        upload_id: gpu.upload_id,
                        indices,
                        uniforms_buffer,
                        bind_group: None,
                        bind_upload_id: 0,
                        visible_count: gpu.num_splats,
                        lod_selected_count: gpu.num_splats,
                        lod_total_count: gpu.num_splats,
                        state: SplatSortState::default(),
                        pending: None,
                    },
                );
            }
            let Some(entry) = resources.entries.get_mut(view_entity) else {
                continue;
            };

            if !cloud_bounds_visible_in_clip(*bounds, model_xf.0, *clip_from_world, 1.4) {
                entry.pending = None;
                entry.state.initialized = false;
                entry.state.sort_started_at = None;
                entry.state.frames_since_sort = 0;
                entry.visible_count = 0;
                entry.lod_selected_count = 0;
                entry.lod_total_count = gpu.num_splats;
                continue;
            }

            if !enable_cpu_sort {
                entry.pending = None;
                entry.state.initialized = false;
                entry.state.sort_started_at = None;
                entry.state.frames_since_sort = 0;
                entry.visible_count = gpu.visible_count;
                entry.lod_selected_count = gpu.visible_count;
                entry.lod_total_count = gpu.num_splats;
                continue;
            }

            entry.state.frames_since_sort = entry.state.frames_since_sort.saturating_add(1);

            // First: poll any in-flight sort. If it finished, upload its result
            // and free the slot. If still busy, render with the previous indices
            // and try again next frame — never block the render app.
            if let Some(mut pending) = entry.pending.take() {
                match block_on(future::poll_once(&mut pending)) {
                    Some(result) => {
                        if let Some(started_at) = entry.state.sort_started_at.take() {
                            entry.state.last_sort_duration_secs =
                                started_at.elapsed().as_secs_f32();
                        }
                        last_sort_time_secs =
                            last_sort_time_secs.max(entry.state.last_sort_duration_secs);
                        let visible_count = result.indices.len() as u32;
                        if visible_count > 0 {
                            write_index_storage(
                                &render_queue,
                                &entry.indices,
                                &result.indices,
                                &render_device,
                            );
                        }
                        entry.visible_count = visible_count;
                        entry.lod_selected_count = result.lod_selected_count;
                        entry.lod_total_count = result.lod_total_count;
                        entry.state.initialized = true;
                        if lod_settings.debug_counters && result.lod_total_count > 0 {
                            bevy::log::info!(
                                "[bevy_spark] LOD selected {} / {} splats before frustum/sort; drew {}",
                                result.lod_selected_count,
                                result.lod_total_count,
                                visible_count
                            );
                        }
                    }
                    None => {
                        entry.pending = Some(pending);
                        continue;
                    }
                }
            }

            let state = &mut entry.state;
            let model = model_xf.0;
            let model_inv = model.inverse();
            let local_view = (model_inv * view_pos.extend(1.0)).truncate();
            let local_dir = (model_inv * view_dir.extend(0.0))
                .truncate()
                .normalize_or_zero();

            let drift = state.last_view_pos.distance(local_view);
            let dot = state.last_view_dir.dot(local_dir);
            // Sort when the view jumps enough to make stale back-to-front order obvious.
            let forced_jump = drift >= SORT_FORCED_DRIFT_METERS || dot < SORT_FORCED_DIR_DOT;
            let interval_frames = adaptive_sort_interval_frames(state.last_sort_duration_secs);
            let waiting_for_interval = state.frames_since_sort < interval_frames;
            // Reuse previous indices for tiny camera drift to avoid CPU sort churn.
            let below_motion_threshold =
                drift < SORT_STABLE_DRIFT_METERS && dot > SORT_STABLE_DIR_DOT;
            if state.initialized && !forced_jump && (below_motion_threshold || waiting_for_interval)
            {
                continue;
            }

            // Spawn a background sort for this cloud/view pair. The task carries
            // a cheap Arc<Vec<...>> of centres and computes distance in world
            // space so scaled clouds keep correct radial ordering.
            let centers = centers_arc.0.clone();
            let radii = radii_arc.0.clone();
            let child_counts = lod_tree.child_counts.clone();
            let child_starts = lod_tree.child_starts.clone();
            let use_lod = lod_settings.enabled && lod_tree.has_lod();
            let n = centers.len();
            let wvp = [view_pos.x, view_pos.y, view_pos.z];
            let wvd = [view_dir.x, view_dir.y, view_dir.z];
            let m0 = model.row(0).to_array();
            let m1 = model.row(1).to_array();
            let m2 = model.row(2).to_array();
            let radius_scale = model
                .x_axis
                .truncate()
                .length()
                .max(model.y_axis.truncate().length())
                .max(model.z_axis.truncate().length());
            let focal_y_pixels = *focal_y_pixels;
            let orthographic = *orthographic;
            let lod_threshold = lod_settings.pixel_radius.max(0.0);
            let lod_center_density = lod_settings.center_density_scale;
            let lod_peripheral_density = lod_settings.peripheral_density_scale;
            let lod_behind_density = lod_settings.behind_density_scale;
            // Frustum culling matrix in the cloud's local source space.
            // Radius mode pads the center test by each splat's local radius
            // through the clip rows so large splats near an edge are retained.
            // The z-bounds test in the shader is omitted here — it culls
            // almost nothing and adds a row dot.
            let clip_from_local = *clip_from_world * model;
            let r0 = clip_from_local.row(0).to_array();
            let r1 = clip_from_local.row(1).to_array();
            let r3 = clip_from_local.row(3).to_array();
            let task = pool.spawn(async move {
                use rayon::prelude::*;
                const LIM: f32 = 1.4;
                let candidates = if use_lod && n > 0 {
                    select_lod_candidates(
                        n,
                        &centers,
                        &radii,
                        &child_counts,
                        &child_starts,
                        wvp,
                        wvd,
                        [m0, m1, m2],
                        radius_scale,
                        focal_y_pixels,
                        orthographic,
                        lod_threshold,
                        lod_center_density,
                        lod_peripheral_density,
                        lod_behind_density,
                    )
                } else {
                    (0..n as u32).collect()
                };
                let lod_selected_count = candidates.len() as u32;
                let lod_total_count = n as u32;
                let mut visible: Vec<(u32, u32)> = candidates
                    .into_par_iter()
                    .filter_map(|splat_index| {
                        let i = splat_index as usize;
                        let c = centers[i];
                        let (c0, c1, c2) = (c[0], c[1], c[2]);
                        let radius = radii.get(i).copied().unwrap_or(0.0);
                        if outside_cpu_clip_bounds(c, radius, r0, r1, r3, LIM, cull_mode) {
                            return None;
                        }
                        let wx = m0[0] * c0 + m0[1] * c1 + m0[2] * c2 + m0[3];
                        let wy = m1[0] * c0 + m1[1] * c1 + m1[2] * c2 + m1[3];
                        let wz = m2[0] * c0 + m2[1] * c1 + m2[2] * c2 + m2[3];
                        let dx = wx - wvp[0];
                        let dy = wy - wvp[1];
                        let dz = wz - wvp[2];
                        let d2 = dx * dx + dy * dy + dz * dz;
                        if min_pixel_radius > 0.0 {
                            let raw_depth = dx * wvd[0] + dy * wvd[1] + dz * wvd[2];
                            if !orthographic && raw_depth <= 0.0 {
                                return None;
                            }
                            let pixel_radius = projected_splat_pixel_radius(
                                radius,
                                radius_scale,
                                focal_y_pixels,
                                orthographic,
                                raw_depth,
                                extent_scale,
                            );
                            if pixel_radius < min_pixel_radius {
                                return None;
                            }
                        }
                        let metric = match sort_mode {
                            SplatSortMode::Radial => d2,
                            SplatSortMode::Depth => {
                                (dx * wvd[0] + dy * wvd[1] + dz * wvd[2]).max(0.0)
                            }
                        };
                        // Invert so larger distance/depth sorts first
                        // (back-to-front).
                        let key = u32::MAX - metric.to_bits();
                        Some((key, splat_index))
                    })
                    .collect();
                radsort::sort_by_key(&mut visible, |&(k, _)| k);
                CpuSortResult {
                    indices: visible.into_iter().map(|(_, i)| i).collect::<Vec<u32>>(),
                    lod_selected_count,
                    lod_total_count,
                }
            });
            entry.pending = Some(task);

            state.last_view_pos = local_view;
            state.last_view_dir = local_dir;
            state.frames_since_sort = 0;
            state.sort_started_at = Some(Instant::now());
        }

        for entry in resources.entries.values() {
            visible_splats = visible_splats.saturating_add(u64::from(entry.visible_count));
            lod_selected_splats =
                lod_selected_splats.saturating_add(u64::from(entry.lod_selected_count));
            last_sort_time_secs = last_sort_time_secs.max(entry.state.last_sort_duration_secs);
        }
    }

    diagnostics.update(|snapshot| {
        snapshot.total_splats = total_splats;
        snapshot.visible_splats = visible_splats;
        snapshot.lod_selected_splats = lod_selected_splats;
        snapshot.last_sort_time_secs = last_sort_time_secs;
        snapshot.quality_preset = settings.quality_preset;
    });
}

// ─── Render-world ──────────────────────────────────────────────────────────

#[derive(Resource)]
pub struct SplatPipeline {
    pub view_layout: BindGroupLayoutDescriptor,
    pub cloud_layout: BindGroupLayoutDescriptor,
    pub shader: Handle<Shader>,
    pub storage_backend: SplatStorageBackend,
}

fn init_splat_pipeline(
    mut commands: Commands,
    asset_server: Res<AssetServer>,
    render_device: Res<RenderDevice>,
) {
    let storage_backend = splat_storage_backend(&render_device);
    let view_layout = BindGroupLayoutDescriptor::new(
        "spark.view_layout",
        &[BindGroupLayoutEntry {
            binding: 0,
            visibility: ShaderStages::VERTEX_FRAGMENT,
            ty: BindingType::Buffer {
                ty: BufferBindingType::Uniform,
                has_dynamic_offset: true,
                min_binding_size: None,
            },
            count: None,
        }],
    );
    let cloud_entries = match storage_backend {
        SplatStorageBackend::StorageBuffers => vec![
            BindGroupLayoutEntry {
                binding: 0,
                visibility: ShaderStages::VERTEX,
                ty: BindingType::Buffer {
                    ty: BufferBindingType::Storage { read_only: true },
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
            BindGroupLayoutEntry {
                binding: 1,
                visibility: ShaderStages::VERTEX,
                ty: BindingType::Buffer {
                    ty: BufferBindingType::Storage { read_only: true },
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
            BindGroupLayoutEntry {
                binding: 2,
                visibility: ShaderStages::VERTEX_FRAGMENT,
                ty: BindingType::Buffer {
                    ty: BufferBindingType::Uniform,
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
            BindGroupLayoutEntry {
                binding: 3,
                visibility: ShaderStages::VERTEX,
                ty: BindingType::Buffer {
                    ty: BufferBindingType::Storage { read_only: true },
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
        ],
        SplatStorageBackend::Textures => vec![
            BindGroupLayoutEntry {
                binding: 0,
                visibility: ShaderStages::VERTEX,
                ty: BindingType::Texture {
                    sample_type: TextureSampleType::Uint,
                    view_dimension: TextureViewDimension::D2,
                    multisampled: false,
                },
                count: None,
            },
            BindGroupLayoutEntry {
                binding: 1,
                visibility: ShaderStages::VERTEX,
                ty: BindingType::Texture {
                    sample_type: TextureSampleType::Uint,
                    view_dimension: TextureViewDimension::D2,
                    multisampled: false,
                },
                count: None,
            },
            BindGroupLayoutEntry {
                binding: 2,
                visibility: ShaderStages::VERTEX_FRAGMENT,
                ty: BindingType::Buffer {
                    ty: BufferBindingType::Uniform,
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
            BindGroupLayoutEntry {
                binding: 3,
                visibility: ShaderStages::VERTEX,
                ty: BindingType::Texture {
                    sample_type: TextureSampleType::Uint,
                    view_dimension: TextureViewDimension::D2,
                    multisampled: false,
                },
                count: None,
            },
        ],
    };
    let cloud_layout = BindGroupLayoutDescriptor::new("spark.cloud_layout", &cloud_entries);
    let shader = asset_server.load(SHADER_PATH);
    commands.insert_resource(SplatPipeline {
        view_layout,
        cloud_layout,
        shader,
        storage_backend,
    });
}

#[derive(Eq, PartialEq, Hash, Clone)]
pub struct SplatPipelineKey {
    pub hdr: bool,
    pub msaa: u32,
    pub storage_backend: SplatStorageBackend,
}

impl SpecializedRenderPipeline for SplatPipeline {
    type Key = SplatPipelineKey;

    fn specialize(&self, key: Self::Key) -> RenderPipelineDescriptor {
        let format = if key.hdr {
            bevy::render::view::ViewTarget::TEXTURE_FORMAT_HDR
        } else {
            TextureFormat::bevy_default()
        };
        RenderPipelineDescriptor {
            label: Some("spark.pipeline".into()),
            layout: vec![self.view_layout.clone(), self.cloud_layout.clone()],
            vertex: VertexState {
                shader: self.shader.clone(),
                entry_point: Some("vs_main".into()),
                buffers: vec![],
                shader_defs: match key.storage_backend {
                    SplatStorageBackend::StorageBuffers => vec![],
                    SplatStorageBackend::Textures => vec!["SPARK_TEXTURE_BACKEND".into()],
                },
            },
            fragment: Some(FragmentState {
                shader: self.shader.clone(),
                entry_point: Some("fs_main".into()),
                targets: vec![Some(ColorTargetState {
                    format,
                    blend: Some(BlendState {
                        color: BlendComponent {
                            src_factor: BlendFactor::One,
                            dst_factor: BlendFactor::OneMinusSrcAlpha,
                            operation: BlendOperation::Add,
                        },
                        alpha: BlendComponent {
                            src_factor: BlendFactor::One,
                            dst_factor: BlendFactor::OneMinusSrcAlpha,
                            operation: BlendOperation::Add,
                        },
                    }),
                    write_mask: ColorWrites::ALL,
                })],
                shader_defs: match key.storage_backend {
                    SplatStorageBackend::StorageBuffers => vec![],
                    SplatStorageBackend::Textures => vec!["SPARK_TEXTURE_BACKEND".into()],
                },
            }),
            primitive: PrimitiveState {
                topology: PrimitiveTopology::TriangleStrip,
                front_face: FrontFace::Ccw,
                cull_mode: None,
                polygon_mode: PolygonMode::Fill,
                ..default()
            },
            depth_stencil: Some(DepthStencilState {
                format: bevy::core_pipeline::core_3d::CORE_3D_DEPTH_FORMAT,
                depth_write_enabled: false,
                depth_compare: CompareFunction::Greater,
                stencil: default(),
                bias: default(),
            }),
            multisample: MultisampleState {
                count: key.msaa.max(1),
                ..default()
            },
            push_constant_ranges: vec![],
            zero_initialize_workgroup_memory: false,
        }
    }
}

#[derive(Component, Clone)]
pub struct SplatCloudBindGroup {
    pub bind_group: BindGroup,
    pub upload_id: u64,
}

#[derive(Component, Clone)]
pub struct SplatViewBindGroup(pub BindGroup);

fn prepare_splat_bind_groups(
    mut commands: Commands,
    pipeline: Res<SplatPipeline>,
    pipeline_cache: Res<PipelineCache>,
    render_device: Res<RenderDevice>,
    render_queue: Res<RenderQueue>,
    settings: Res<SparkSettings>,
    view_uniforms: Res<ViewUniforms>,
    mut clouds: Query<(
        Entity,
        &SplatCloudGpu,
        &ExtractedSplatTransform,
        Option<&ExtractedSplatCloudSettings>,
        Option<&SplatCloudBindGroup>,
        Option<&mut SplatPerViewResources>,
    )>,
    views: Query<(Entity, &ExtractedView, Option<&SplatViewBindGroup>)>,
) {
    let view_layout = pipeline_cache.get_bind_group_layout(&pipeline.view_layout);
    let cloud_layout = pipeline_cache.get_bind_group_layout(&pipeline.cloud_layout);

    // View bind group (shared across all clouds for a given view).
    if let Some(view_binding) = view_uniforms.uniforms.binding() {
        for (view_entity, _, maybe_view_bg) in &views {
            commands
                .entity(view_entity)
                .insert(SplatViewEntity(view_entity));
            if maybe_view_bg.is_some() {
                continue;
            }
            let bg = render_device.create_bind_group(
                "spark.view_bg",
                &view_layout,
                &[BindGroupEntry {
                    binding: 0,
                    resource: view_binding.clone(),
                }],
            );
            commands.entity(view_entity).insert(SplatViewBindGroup(bg));
        }
    }

    // Global fallback bind group + uniforms upload. Normal draws use the
    // per-view bind groups below once the per-view entries have been prepared.
    let fallback_render_size = views
        .iter()
        .next()
        .map(|(_, v, _)| {
            let r = v.viewport.zw();
            [r.x as f32, r.y as f32]
        })
        .unwrap_or([1280.0, 720.0]);
    let render_sizes: HashMap<Entity, [f32; 2]> = views
        .iter()
        .map(|(view_entity, v, _)| {
            let r = v.viewport.zw();
            (view_entity, [r.x as f32, r.y as f32])
        })
        .collect();

    for (entity, gpu, xf, cloud_settings, maybe_global_bg, maybe_resources) in &mut clouds {
        let q = effective_quality_settings(&settings, cloud_settings);
        let texture_width = match &gpu.storage {
            SplatCloudGpuStorage::Textures { texture_width, .. } => *texture_width,
            SplatCloudGpuStorage::StorageBuffers { .. } => 1,
        };
        let make_uniforms = |render_size| CloudUniforms {
            model: xf.0.to_cols_array_2d(),
            num_splats: gpu.num_splats,
            anti_aliased: gpu.anti_aliased,
            sh_degree: gpu.sh_degree,
            _pad0: 0,
            render_size,
            max_stddev: q.max_stddev,
            min_alpha: q.min_alpha,
            min_pixel_radius: q.min_pixel_radius,
            max_pixel_radius: q.max_pixel_radius,
            falloff_profile: q.falloff_profile.shader_code(),
            high_alpha_profile: q.high_alpha_profile.shader_code(),
            texture_width,
            _pad1: [0; 3],
        };
        let uniforms = make_uniforms(fallback_render_size);
        render_queue.write_buffer(&gpu.uniforms_buffer, 0, bytemuck::cast_slice(&[uniforms]));

        let global_bg_stale = maybe_global_bg
            .map(|bind_group| bind_group.upload_id != gpu.upload_id)
            .unwrap_or(true);
        if global_bg_stale {
            let bg = match &gpu.storage {
                SplatCloudGpuStorage::StorageBuffers {
                    splats_buffer,
                    sh_buffer,
                    indices_buffer,
                } => render_device.create_bind_group(
                    "spark.cloud_bg",
                    &cloud_layout,
                    &[
                        BindGroupEntry {
                            binding: 0,
                            resource: BindingResource::Buffer(BufferBinding {
                                buffer: splats_buffer,
                                offset: 0,
                                size: None,
                            }),
                        },
                        BindGroupEntry {
                            binding: 1,
                            resource: BindingResource::Buffer(BufferBinding {
                                buffer: indices_buffer,
                                offset: 0,
                                size: None,
                            }),
                        },
                        BindGroupEntry {
                            binding: 2,
                            resource: BindingResource::Buffer(BufferBinding {
                                buffer: &gpu.uniforms_buffer,
                                offset: 0,
                                size: None,
                            }),
                        },
                        BindGroupEntry {
                            binding: 3,
                            resource: BindingResource::Buffer(BufferBinding {
                                buffer: sh_buffer,
                                offset: 0,
                                size: None,
                            }),
                        },
                    ],
                ),
                SplatCloudGpuStorage::Textures {
                    splats,
                    sh,
                    indices,
                    ..
                } => render_device.create_bind_group(
                    "spark.cloud_bg",
                    &cloud_layout,
                    &[
                        BindGroupEntry {
                            binding: 0,
                            resource: BindingResource::TextureView(&splats.view),
                        },
                        BindGroupEntry {
                            binding: 1,
                            resource: BindingResource::TextureView(&indices.view),
                        },
                        BindGroupEntry {
                            binding: 2,
                            resource: BindingResource::Buffer(BufferBinding {
                                buffer: &gpu.uniforms_buffer,
                                offset: 0,
                                size: None,
                            }),
                        },
                        BindGroupEntry {
                            binding: 3,
                            resource: BindingResource::TextureView(&sh.view),
                        },
                    ],
                ),
            };
            commands.entity(entity).insert(SplatCloudBindGroup {
                bind_group: bg,
                upload_id: gpu.upload_id,
            });
        }

        if let Some(mut resources) = maybe_resources {
            for (view_entity, entry) in resources.entries.iter_mut() {
                let render_size = render_sizes
                    .get(view_entity)
                    .copied()
                    .unwrap_or(fallback_render_size);
                let uniforms = make_uniforms(render_size);
                render_queue.write_buffer(
                    &entry.uniforms_buffer,
                    0,
                    bytemuck::cast_slice(&[uniforms]),
                );
                if entry.bind_group.is_none() || entry.bind_upload_id != gpu.upload_id {
                    entry.bind_group = match (&gpu.storage, &entry.indices) {
                        (
                            SplatCloudGpuStorage::StorageBuffers {
                                splats_buffer,
                                sh_buffer,
                                ..
                            },
                            SplatIndexStorage::Buffer(indices_buffer),
                        ) => Some(render_device.create_bind_group(
                            "spark.view_cloud_bg",
                            &cloud_layout,
                            &[
                                BindGroupEntry {
                                    binding: 0,
                                    resource: BindingResource::Buffer(BufferBinding {
                                        buffer: splats_buffer,
                                        offset: 0,
                                        size: None,
                                    }),
                                },
                                BindGroupEntry {
                                    binding: 1,
                                    resource: BindingResource::Buffer(BufferBinding {
                                        buffer: indices_buffer,
                                        offset: 0,
                                        size: None,
                                    }),
                                },
                                BindGroupEntry {
                                    binding: 2,
                                    resource: BindingResource::Buffer(BufferBinding {
                                        buffer: &entry.uniforms_buffer,
                                        offset: 0,
                                        size: None,
                                    }),
                                },
                                BindGroupEntry {
                                    binding: 3,
                                    resource: BindingResource::Buffer(BufferBinding {
                                        buffer: sh_buffer,
                                        offset: 0,
                                        size: None,
                                    }),
                                },
                            ],
                        )),
                        (
                            SplatCloudGpuStorage::Textures { splats, sh, .. },
                            SplatIndexStorage::Texture(indices),
                        ) => Some(render_device.create_bind_group(
                            "spark.view_cloud_bg",
                            &cloud_layout,
                            &[
                                BindGroupEntry {
                                    binding: 0,
                                    resource: BindingResource::TextureView(&splats.view),
                                },
                                BindGroupEntry {
                                    binding: 1,
                                    resource: BindingResource::TextureView(&indices.view),
                                },
                                BindGroupEntry {
                                    binding: 2,
                                    resource: BindingResource::Buffer(BufferBinding {
                                        buffer: &entry.uniforms_buffer,
                                        offset: 0,
                                        size: None,
                                    }),
                                },
                                BindGroupEntry {
                                    binding: 3,
                                    resource: BindingResource::TextureView(&sh.view),
                                },
                            ],
                        )),
                        _ => None,
                    };
                    entry.bind_upload_id = gpu.upload_id;
                }
            }
        }
    }
}

// ─── Queueing ─────────────────────────────────────────────────────────────

fn queue_splat_clouds(
    transparent_draw: Res<DrawFunctions<Transparent3d>>,
    pipeline: Res<SplatPipeline>,
    pipeline_cache: Res<PipelineCache>,
    mut pipelines: ResMut<SpecializedRenderPipelines<SplatPipeline>>,
    mut transparent_phases: ResMut<ViewSortedRenderPhases<Transparent3d>>,
    settings: Res<SparkSettings>,
    views: Query<(&ExtractedView, &Msaa)>,
    clouds: Query<
        (
            Entity,
            &MainEntity,
            &ExtractedSplatTransform,
            &SplatCloudBounds,
        ),
        With<SplatCloudGpu>,
    >,
    mut warned_approximate: Local<bool>,
    mut warned_single_exact: Local<bool>,
) {
    let draw_fn = transparent_draw.read().id::<DrawSplatCloud>();
    let multi_cloud_mode = settings.multi_cloud_mode;
    for (view, msaa) in &views {
        let Some(phase) = transparent_phases.get_mut(&view.retained_view_entity) else {
            continue;
        };
        let clip_from_world = view
            .clip_from_world
            .unwrap_or_else(|| view.clip_from_view * view.world_from_view.to_matrix().inverse());
        let key = SplatPipelineKey {
            hdr: view.hdr,
            msaa: msaa.samples(),
            storage_backend: pipeline.storage_backend,
        };
        let pipeline_id = pipelines.specialize(&pipeline_cache, &pipeline, key);
        let mut queued_clouds: Vec<(Entity, MainEntity, f32)> = clouds
            .iter()
            .filter_map(|(entity, main_entity, xf, bounds)| {
                if !cloud_bounds_visible_in_clip(*bounds, xf.0, clip_from_world, 1.4) {
                    return None;
                }
                Some((entity, *main_entity, f32::NEG_INFINITY))
            })
            .collect();

        if queued_clouds.len() > 1 {
            match multi_cloud_mode {
                SplatMultiCloudMode::MultiCloudApproximate => {
                    if !*warned_approximate {
                        bevy::log::warn!(
                            "[bevy_spark] multiple SplatClouds are visible; \
                             BEVY_SPARK_MULTI_CLOUD_MODE=approximate sorts splats within each \
                             cloud, then blends whole clouds by entity distance. Overlapping \
                             clouds can show transparency artifacts."
                        );
                        *warned_approximate = true;
                    }
                }
                SplatMultiCloudMode::SingleCloudExact => {
                    if !*warned_single_exact {
                        bevy::log::warn!(
                            "[bevy_spark] BEVY_SPARK_MULTI_CLOUD_MODE=single-exact only draws \
                             one SplatCloud per view; skipping additional visible clouds to avoid \
                             approximate cross-cloud blending."
                        );
                        *warned_single_exact = true;
                    }
                    queued_clouds.sort_by(|a, b| a.2.total_cmp(&b.2));
                    queued_clouds.truncate(1);
                }
            }
        }

        for (entity, main_entity, distance) in queued_clouds {
            phase.add(Transparent3d {
                entity: (entity, main_entity),
                pipeline: pipeline_id,
                draw_function: draw_fn,
                distance,
                batch_range: 0..1,
                extra_index: PhaseItemExtraIndex::None,
                indexed: false,
            });
        }
    }
}

// ─── Draw command ─────────────────────────────────────────────────────────

type DrawSplatCloud = (
    SetItemPipeline,
    SetSplatViewBindGroup<0>,
    SetSplatCloudBindGroup<1>,
    DrawSplatQuads,
);

struct SetSplatViewBindGroup<const I: usize>;
impl<P: bevy::render::render_phase::PhaseItem, const I: usize> RenderCommand<P>
    for SetSplatViewBindGroup<I>
{
    type Param = ();
    type ViewQuery = (Read<ViewUniformOffset>, Read<SplatViewBindGroup>);
    type ItemQuery = ();

    fn render<'w>(
        _item: &P,
        (view_uniform_offset, view_bg): (&'w ViewUniformOffset, &'w SplatViewBindGroup),
        _: Option<()>,
        _: SystemParamItem<'w, '_, Self::Param>,
        pass: &mut TrackedRenderPass<'w>,
    ) -> RenderCommandResult {
        pass.set_bind_group(I, &view_bg.0, &[view_uniform_offset.offset]);
        RenderCommandResult::Success
    }
}

struct SetSplatCloudBindGroup<const I: usize>;
impl<P: bevy::render::render_phase::PhaseItem, const I: usize> RenderCommand<P>
    for SetSplatCloudBindGroup<I>
{
    type Param = ();
    type ViewQuery = Read<SplatViewEntity>;
    type ItemQuery = (
        Read<SplatCloudBindGroup>,
        Option<Read<SplatPerViewResources>>,
    );

    fn render<'w>(
        _item: &P,
        view_entity: &'w SplatViewEntity,
        item: Option<(&'w SplatCloudBindGroup, Option<&'w SplatPerViewResources>)>,
        _: SystemParamItem<'w, '_, Self::Param>,
        pass: &mut TrackedRenderPass<'w>,
    ) -> RenderCommandResult {
        let Some((global_bg, per_view_sorts)) = item else {
            return RenderCommandResult::Skip;
        };
        if let Some(view_bg) = per_view_sorts
            .and_then(|sorts| sorts.entries.get(&view_entity.0))
            .and_then(|entry| entry.bind_group.as_ref())
        {
            pass.set_bind_group(I, view_bg, &[]);
        } else {
            pass.set_bind_group(I, &global_bg.bind_group, &[]);
        }
        RenderCommandResult::Success
    }
}

struct DrawSplatQuads;
impl<P: bevy::render::render_phase::PhaseItem> RenderCommand<P> for DrawSplatQuads {
    type Param = ();
    type ViewQuery = Read<SplatViewEntity>;
    type ItemQuery = (Read<SplatCloudGpu>, Option<Read<SplatPerViewResources>>);

    fn render<'w>(
        _item: &P,
        view_entity: &'w SplatViewEntity,
        item: Option<(&'w SplatCloudGpu, Option<&'w SplatPerViewResources>)>,
        _: SystemParamItem<'w, '_, Self::Param>,
        pass: &mut TrackedRenderPass<'w>,
    ) -> RenderCommandResult {
        let Some((gpu, per_view_sorts)) = item else {
            return RenderCommandResult::Skip;
        };
        let visible_count = per_view_sorts
            .and_then(|sorts| sorts.entries.get(&view_entity.0))
            .map(|entry| entry.visible_count)
            .unwrap_or(gpu.visible_count);
        pass.draw(0..4, 0..visible_count);
        RenderCommandResult::Success
    }
}

#[cfg(test)]
mod tests {
    use bevy::math::{Mat4, Vec3};
    use bevy::render::settings::WgpuLimits;

    use super::{
        ExtractedSplatCloudSettings, SparkConfigSource, SparkDiagnostics, SparkQualityPreset,
        SparkSettings, SplatCloudBounds, SplatCloudSettings, SplatCpuCullMode, SplatFalloffProfile,
        SplatHighAlphaProfile, SplatLodSettings, SplatMultiCloudMode, SplatQualitySettings,
        SplatSortBackend, SplatSortMode, SplatSortSettings, adaptive_sort_interval_frames,
        cloud_bounds_visible_in_clip, effective_lod_settings, effective_quality_settings,
        effective_sort_settings, gpu_compute_sort_supported_limits, outside_cpu_clip_bounds,
        projected_splat_pixel_radius, select_lod_candidates, storage_buffer_fits_limits,
    };

    fn identity_rows() -> [[f32; 4]; 3] {
        [
            [1.0, 0.0, 0.0, 0.0],
            [0.0, 1.0, 0.0, 0.0],
            [0.0, 0.0, 1.0, 0.0],
        ]
    }

    #[test]
    fn adaptive_sort_interval_increases_with_sort_cost() {
        assert_eq!(adaptive_sort_interval_frames(0.001), 0);
        assert_eq!(adaptive_sort_interval_frames(0.010), 1);
        assert_eq!(adaptive_sort_interval_frames(0.025), 3);
        assert_eq!(adaptive_sort_interval_frames(0.075), 6);
    }

    #[test]
    fn storage_buffer_limits_check_binding_and_buffer_caps() {
        let mut limits = WgpuLimits {
            max_storage_buffer_binding_size: 16,
            max_buffer_size: 32,
            ..Default::default()
        };

        assert!(storage_buffer_fits_limits(&limits, 16));
        assert!(!storage_buffer_fits_limits(&limits, 17));

        limits.max_storage_buffer_binding_size = 64;
        assert!(storage_buffer_fits_limits(&limits, 32));
        assert!(!storage_buffer_fits_limits(&limits, 33));
    }

    #[test]
    fn per_cloud_settings_override_global_defaults() {
        let global = SparkSettings {
            quality_preset: SparkQualityPreset::Reference,
            quality: SplatQualitySettings::default(),
            sort: SplatSortSettings::default(),
            lod: SplatLodSettings::default(),
            multi_cloud_mode: SplatMultiCloudMode::MultiCloudApproximate,
            upload_chunk_bytes: 4096,
        };
        let quality = SplatQualitySettings {
            max_stddev: 2.0,
            min_alpha: 0.01,
            min_pixel_radius: 0.5,
            max_pixel_radius: 128.0,
            falloff_profile: SplatFalloffProfile::EdgeNormalized,
            high_alpha_profile: SplatHighAlphaProfile::Bounded,
        };
        let sort = SplatSortSettings {
            backend: SplatSortBackend::None,
            mode: SplatSortMode::Depth,
            cpu_cull_mode: SplatCpuCullMode::Center,
        };
        let lod = SplatLodSettings {
            enabled: false,
            pixel_radius: 3.0,
            center_density_scale: 0.75,
            peripheral_density_scale: 0.25,
            behind_density_scale: 0.0,
            debug_counters: true,
        };
        let cloud = ExtractedSplatCloudSettings(SplatCloudSettings {
            quality: Some(quality),
            sort: Some(sort),
            lod: Some(lod),
        });

        assert_eq!(effective_quality_settings(&global, Some(&cloud)), quality);
        assert_eq!(effective_sort_settings(&global, Some(&cloud)), sort);
        assert_eq!(effective_lod_settings(&global, Some(&cloud)), lod);
    }

    #[test]
    fn diagnostics_snapshot_reflects_shared_updates() {
        let diagnostics = SparkDiagnostics::default();
        diagnostics.update(|snapshot| {
            snapshot.total_splats = 10;
            snapshot.visible_splats = 7;
            snapshot.lod_selected_splats = 8;
            snapshot.last_sort_time_secs = 0.002;
            snapshot.last_upload_time_secs = 0.125;
            snapshot.upload_progress = 0.5;
            snapshot.quality_preset = SparkQualityPreset::Performance;
        });

        let snapshot = diagnostics.snapshot();
        assert_eq!(snapshot.total_splats, 10);
        assert_eq!(snapshot.visible_splats, 7);
        assert_eq!(snapshot.lod_selected_splats, 8);
        assert_eq!(snapshot.quality_preset, SparkQualityPreset::Performance);
        assert!((snapshot.last_sort_time_secs - 0.002).abs() < f32::EPSILON);
        assert!((snapshot.last_upload_time_secs - 0.125).abs() < f32::EPSILON);
        assert!((snapshot.upload_progress - 0.5).abs() < f32::EPSILON);
    }

    #[test]
    fn quality_presets_cover_reference_performance_and_capture_modes() {
        assert_eq!(SparkQualityPreset::default(), SparkQualityPreset::Balanced);
        assert_eq!(
            SparkSettings::default().quality_preset,
            SparkQualityPreset::Balanced
        );

        let reference = SparkSettings::from_preset(SparkQualityPreset::Reference);
        assert_eq!(reference.quality, SplatQualitySettings::default());
        assert_eq!(reference.sort.backend, SplatSortBackend::Cpu);
        assert_eq!(
            SparkConfigSource::default(),
            SparkConfigSource::ResourceWithEnvOverrides
        );

        let balanced = SparkSettings::from_preset(SparkQualityPreset::Balanced);
        assert_eq!(balanced.quality.max_stddev, 2.0);
        assert_eq!(balanced.quality.min_alpha, reference.quality.min_alpha);
        assert_eq!(balanced.quality.min_pixel_radius, 0.25);
        assert_eq!(
            balanced.quality.falloff_profile,
            SplatFalloffProfile::EdgeNormalized
        );
        assert_eq!(
            balanced.quality.high_alpha_profile,
            SplatHighAlphaProfile::Bounded
        );
        assert_eq!(balanced.lod.pixel_radius, reference.lod.pixel_radius);

        let performance = SparkSettings::from_preset(SparkQualityPreset::Performance);
        assert_eq!(performance.quality.max_stddev, 2.0);
        assert_eq!(
            performance.quality.falloff_profile,
            SplatFalloffProfile::EdgeNormalized
        );
        assert!(performance.lod.enabled);
        assert!(performance.lod.pixel_radius > reference.lod.pixel_radius);

        let capture = SparkSettings::from_preset(SparkQualityPreset::Capture);
        assert!(!capture.lod.enabled);
        assert_eq!(capture.sort.backend, SplatSortBackend::Cpu);
        assert_eq!(
            capture.multi_cloud_mode,
            SplatMultiCloudMode::SingleCloudExact
        );
    }

    #[test]
    fn gpu_compute_sort_limits_require_shader_shape_and_storage_bindings() {
        let limits = WgpuLimits::default();
        assert!(gpu_compute_sort_supported_limits(&limits));

        let mut low_workgroup_size = limits.clone();
        low_workgroup_size.max_compute_workgroup_size_x = 255;
        assert!(!gpu_compute_sort_supported_limits(&low_workgroup_size));

        let mut low_invocations = limits.clone();
        low_invocations.max_compute_invocations_per_workgroup = 255;
        assert!(!gpu_compute_sort_supported_limits(&low_invocations));

        let mut low_storage_bindings = limits;
        low_storage_bindings.max_storage_buffers_per_shader_stage = 6;
        assert!(!gpu_compute_sort_supported_limits(&low_storage_bindings));
    }

    #[test]
    fn radius_cull_keeps_large_splats_crossing_clip_edge() {
        let r0 = [1.0, 0.0, 0.0, 0.0];
        let r1 = [0.0, 1.0, 0.0, 0.0];
        let r3 = [0.0, 0.0, 0.0, 1.0];

        assert!(outside_cpu_clip_bounds(
            [1.5, 0.0, 0.0],
            0.2,
            r0,
            r1,
            r3,
            1.4,
            SplatCpuCullMode::Center,
        ));
        assert!(!outside_cpu_clip_bounds(
            [1.5, 0.0, 0.0],
            0.2,
            r0,
            r1,
            r3,
            1.4,
            SplatCpuCullMode::Radius,
        ));
        assert!(outside_cpu_clip_bounds(
            [2.0, 0.0, 0.0],
            0.2,
            r0,
            r1,
            r3,
            1.4,
            SplatCpuCullMode::Radius,
        ));
    }

    #[test]
    fn projected_pixel_radius_uses_depth_for_perspective_only() {
        let perspective = projected_splat_pixel_radius(0.5, 2.0, 100.0, false, 10.0, 2.0);
        assert!((perspective - 20.0).abs() < 1e-5);

        let orthographic = projected_splat_pixel_radius(0.5, 2.0, 100.0, true, 10.0, 2.0);
        assert!((orthographic - 200.0).abs() < 1e-5);
    }

    #[test]
    fn cloud_bounds_visibility_rejects_offscreen_aabb() {
        let visible = SplatCloudBounds {
            min: Vec3::new(-0.5, -0.5, 0.0),
            max: Vec3::new(0.5, 0.5, 0.0),
        };
        assert!(cloud_bounds_visible_in_clip(
            visible,
            Mat4::IDENTITY,
            Mat4::IDENTITY,
            1.4,
        ));

        let offscreen = SplatCloudBounds {
            min: Vec3::new(3.0, -0.5, 0.0),
            max: Vec3::new(4.0, 0.5, 0.0),
        };
        assert!(!cloud_bounds_visible_in_clip(
            offscreen,
            Mat4::IDENTITY,
            Mat4::IDENTITY,
            1.4,
        ));
    }

    #[test]
    fn lod_keeps_parent_below_pixel_threshold() {
        let selected = select_lod_candidates(
            3,
            &[[0.0, 0.0, -3.0], [0.0, 0.0, -3.0], [0.0, 0.0, -3.0]],
            &[1.0, 0.1, 0.1],
            &[2, 0, 0],
            &[1, 0, 0],
            [0.0, 0.0, 0.0],
            [0.0, 0.0, -1.0],
            identity_rows(),
            1.0,
            100.0,
            false,
            100.0,
            1.0,
            1.0,
            1.0,
        );

        assert_eq!(selected, vec![0]);
    }

    #[test]
    fn lod_expands_parent_above_pixel_threshold() {
        let selected = select_lod_candidates(
            3,
            &[[0.0, 0.0, -3.0], [0.0, 0.0, -3.0], [0.0, 0.0, -3.0]],
            &[1.0, 0.1, 0.1],
            &[2, 0, 0],
            &[1, 0, 0],
            [0.0, 0.0, 0.0],
            [0.0, 0.0, -1.0],
            identity_rows(),
            1.0,
            100.0,
            false,
            1.0,
            1.0,
            1.0,
            1.0,
        );

        assert_eq!(selected, vec![1, 2]);
    }

    #[test]
    fn lod_uses_lower_density_for_peripheral_splats() {
        let centers = [[10.0, 0.0, -3.0], [10.0, 0.0, -3.0], [10.0, 0.0, -3.0]];
        let without_foveation = select_lod_candidates(
            3,
            &centers,
            &[1.0, 0.1, 0.1],
            &[2, 0, 0],
            &[1, 0, 0],
            [0.0, 0.0, 0.0],
            [0.0, 0.0, -1.0],
            identity_rows(),
            1.0,
            100.0,
            false,
            10.0,
            1.0,
            1.0,
            1.0,
        );
        let with_foveation = select_lod_candidates(
            3,
            &centers,
            &[1.0, 0.1, 0.1],
            &[2, 0, 0],
            &[1, 0, 0],
            [0.0, 0.0, 0.0],
            [0.0, 0.0, -1.0],
            identity_rows(),
            1.0,
            100.0,
            false,
            10.0,
            1.0,
            0.1,
            1.0,
        );

        assert_eq!(without_foveation, vec![1, 2]);
        assert_eq!(with_foveation, vec![0]);
    }
}