Skip to main content

bevy_pbr/render/
gpu_preprocess.rs

1//! GPU mesh preprocessing.
2//!
3//! This is an optional pass that uses a compute shader to reduce the amount of
4//! data that has to be transferred from the CPU to the GPU. When enabled,
5//! instead of transferring [`MeshUniform`]s to the GPU, we transfer the smaller
6//! [`MeshInputUniform`]s instead and use the GPU to calculate the remaining
7//! derived fields in [`MeshUniform`].
8
9use core::num::{NonZero, NonZeroU64};
10
11use bevy_app::{App, Plugin};
12use bevy_asset::{embedded_asset, load_embedded_asset, Handle};
13use bevy_core_pipeline::{
14    deferred::node::late_deferred_prepass,
15    mip_generation::experimental::depth::{early_downsample_depth, ViewDepthPyramid},
16    prepass::{
17        node::{early_prepass, late_prepass},
18        DeferredPrepass, DepthPrepass, MotionVectorPrepass, NormalPrepass, PreviousViewData,
19        PreviousViewUniformOffset, PreviousViewUniforms,
20    },
21    schedule::{Core3d, Core3dSystems},
22};
23use bevy_derive::{Deref, DerefMut};
24use bevy_ecs::{
25    component::Component,
26    entity::Entity,
27    prelude::resource_exists,
28    query::{Has, Or, With, Without},
29    resource::Resource,
30    schedule::{common_conditions::any_match_filter, IntoScheduleConfigs as _},
31    system::{Commands, Query, Res, ResMut},
32    world::{FromWorld, World},
33};
34use bevy_log::warn_once;
35use bevy_math::Vec4;
36use bevy_platform::collections::HashMap;
37use bevy_render::{
38    batching::gpu_preprocessing::{
39        clear_scene_unpacking_buffers, BatchedInstanceBuffers, BinUnpackingMetadataIndex,
40        BuildIndirectParametersMetadata, GpuBinMetadata, GpuBinUnpackingMetadata,
41        GpuOcclusionCullingWorkItemBuffers, GpuPreprocessingMode, GpuPreprocessingSupport,
42        GpuUniformAllocationMetadata, IndirectBatchSet, IndirectParametersBuffers,
43        IndirectParametersBuildJob, IndirectParametersBuildJobs, IndirectParametersIndexed,
44        IndirectParametersMetadata, IndirectParametersNonIndexed,
45        LatePreprocessWorkItemIndirectParameters, PreprocessWorkItem, PreprocessWorkItemBuffers,
46        SceneUnpackingBuffers, SceneUnpackingBuffersKey, SceneUnpackingJob,
47        UniformAllocationMetadataIndex, UntypedPhaseBatchedInstanceBuffers,
48        UntypedPhaseIndirectParametersBuffers,
49    },
50    diagnostic::RecordDiagnostics as _,
51    occlusion_culling::OcclusionCulling,
52    render_phase::{GpuRenderBinnedMeshInstance, UNIFORM_ALLOCATION_WORKGROUP_SIZE},
53    render_resource::{
54        binding_types::{storage_buffer, storage_buffer_read_only, texture_2d, uniform_buffer},
55        BindGroup, BindGroupEntries, BindGroupLayoutDescriptor, BindGroupLayoutEntries,
56        BindingResource, Buffer, BufferBinding, BufferVec, CachedComputePipelineId,
57        ComputePassDescriptor, ComputePipelineDescriptor, DynamicBindGroupLayoutEntries,
58        PartialBufferVec, PipelineCache, RawBufferVec, ShaderStages, ShaderType,
59        SparseBufferUpdateBindGroups, SparseBufferUpdateJobs, SparseBufferUpdatePipelines,
60        SpecializedComputePipeline, SpecializedComputePipelines, TextureSampleType,
61        UninitBufferVec,
62    },
63    renderer::{RenderContext, RenderDevice, RenderQueue, ViewQuery},
64    settings::WgpuFeatures,
65    view::{
66        ExtractedView, NoIndirectDrawing, RenderVisibilityRanges, RetainedViewEntity, ViewUniform,
67        ViewUniformOffset, ViewUniforms,
68    },
69    GpuResourceAppExt, Render, RenderApp, RenderSystems,
70};
71use bevy_shader::Shader;
72use bevy_utils::{default, TypeIdHashMap};
73use bitflags::bitflags;
74use smallvec::{smallvec, SmallVec};
75use tracing::warn;
76
77use crate::{
78    LightEntity, MeshCullingData, MeshCullingDataBuffer, MeshInputUniform, MeshUniform,
79    PreviousMeshInputUniform,
80};
81
82use super::{ShadowView, ViewLightEntities};
83
84/// The GPU workgroup size.
85const WORKGROUP_SIZE: usize = 64;
86
87/// A plugin that builds mesh uniforms on GPU.
88///
89/// This will only be added if the platform supports compute shaders (e.g. not
90/// on WebGL 2).
91pub struct GpuMeshPreprocessPlugin {
92    /// Whether we're building [`MeshUniform`]s on GPU.
93    ///
94    /// This requires compute shader support and so will be forcibly disabled if
95    /// the platform doesn't support those.
96    pub use_gpu_instance_buffer_builder: bool,
97}
98
99/// The compute shader pipelines for the GPU mesh preprocessing and indirect
100/// parameter building passes.
101#[derive(impl bevy_ecs::component::Component for PreprocessPipelines where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {
    const STORAGE_TYPE: bevy_ecs::component::StorageType =
        bevy_ecs::component::StorageType::SparseSet;
    type Mutability = bevy_ecs::component::Mutable;
    fn register_required_components(_requiree:
            bevy_ecs::component::ComponentId,
        required_components:
            &mut bevy_ecs::component::RequiredComponentsRegistrator) {
        let resource_component_id =
            if let ::core::option::Option::Some(id) =
                    required_components.components_registrator().component_id::<PreprocessPipelines>()
                {
                id
            } else {
                required_components.components_registrator().register_component::<PreprocessPipelines>()
            };
        required_components.register_required::<bevy_ecs::resource::IsResource>(move
                ||
                bevy_ecs::resource::IsResource::new(resource_component_id));
    }
    fn clone_behavior() -> bevy_ecs::component::ComponentCloneBehavior {
        use bevy_ecs::component::{
            DefaultCloneBehaviorBase, DefaultCloneBehaviorViaClone,
        };
        (&&&bevy_ecs::component::DefaultCloneBehaviorSpecialization::<Self>::default()).default_clone_behavior()
    }
    fn relationship_accessor()
        ->
            ::core::option::Option<bevy_ecs::relationship::ComponentRelationshipAccessor<Self>> {
        ::core::option::Option::None
    }
}
impl bevy_ecs::resource::Resource for PreprocessPipelines where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {}Resource)]
102pub struct PreprocessPipelines {
103    /// The pipeline used for CPU culling. This pipeline doesn't populate
104    /// indirect parameter metadata.
105    pub direct_preprocess: PreprocessPipeline,
106    /// The pipeline used for mesh preprocessing when GPU frustum culling is in
107    /// use, but occlusion culling isn't.
108    ///
109    /// This pipeline populates indirect parameter metadata.
110    pub gpu_frustum_culling_preprocess: PreprocessPipeline,
111    /// The pipeline used for the first phase of occlusion culling.
112    ///
113    /// This pipeline culls, transforms meshes, and populates indirect parameter
114    /// metadata.
115    pub early_gpu_occlusion_culling_preprocess: PreprocessPipeline,
116    /// The pipeline used for the second phase of occlusion culling.
117    ///
118    /// This pipeline culls, transforms meshes, and populates indirect parameter
119    /// metadata.
120    pub late_gpu_occlusion_culling_preprocess: PreprocessPipeline,
121    /// The pipeline that builds indirect draw parameters for indexed meshes,
122    /// when frustum culling is enabled but occlusion culling *isn't* enabled.
123    pub gpu_frustum_culling_build_indexed_indirect_params: BuildIndirectParametersPipeline,
124    /// The pipeline that builds indirect draw parameters for non-indexed
125    /// meshes, when frustum culling is enabled but occlusion culling *isn't*
126    /// enabled.
127    pub gpu_frustum_culling_build_non_indexed_indirect_params: BuildIndirectParametersPipeline,
128    /// Compute shader pipelines for the early prepass phase that draws meshes
129    /// visible in the previous frame.
130    pub early_phase: PreprocessPhasePipelines,
131    /// Compute shader pipelines for the late prepass phase that draws meshes
132    /// that weren't visible in the previous frame, but became visible this
133    /// frame.
134    pub late_phase: PreprocessPhasePipelines,
135    /// Compute shader pipelines for the main color phase.
136    pub main_phase: PreprocessPhasePipelines,
137    /// Compute shader pipelines for the bin unpacking step.
138    pub bin_unpacking: BinUnpackingPipeline,
139    /// Compute shader pipelines for the uniform allocation step.
140    pub uniform_allocation: UniformAllocationPipelines,
141}
142
143/// Compute shader pipelines for a specific phase: early, late, or main.
144///
145/// The distinction between these phases is relevant for occlusion culling.
146#[derive(#[automatically_derived]
impl ::core::clone::Clone for PreprocessPhasePipelines {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            reset_indirect_batch_sets: ::core::clone::Clone::clone(&self.reset_indirect_batch_sets),
            gpu_occlusion_culling_build_indexed_indirect_params: ::core::clone::Clone::clone(&self.gpu_occlusion_culling_build_indexed_indirect_params),
            gpu_occlusion_culling_build_non_indexed_indirect_params: ::core::clone::Clone::clone(&self.gpu_occlusion_culling_build_non_indexed_indirect_params),
        }
    }
}Clone)]
147pub struct PreprocessPhasePipelines {
148    /// The pipeline that resets the indirect draw counts used in
149    /// `multi_draw_indirect_count` to 0 in preparation for a new pass.
150    pub reset_indirect_batch_sets: ResetIndirectBatchSetsPipeline,
151    /// The pipeline used for indexed indirect parameter building.
152    ///
153    /// This pipeline converts indirect parameter metadata into indexed indirect
154    /// parameters.
155    pub gpu_occlusion_culling_build_indexed_indirect_params: BuildIndirectParametersPipeline,
156    /// The pipeline used for non-indexed indirect parameter building.
157    ///
158    /// This pipeline converts indirect parameter metadata into non-indexed
159    /// indirect parameters.
160    pub gpu_occlusion_culling_build_non_indexed_indirect_params: BuildIndirectParametersPipeline,
161}
162
163/// The pipeline for the GPU mesh preprocessing shader.
164pub struct PreprocessPipeline {
165    /// The bind group layout for the compute shader.
166    pub bind_group_layout: BindGroupLayoutDescriptor,
167    /// The shader asset handle.
168    pub shader: Handle<Shader>,
169    /// The pipeline ID for the compute shader.
170    ///
171    /// This gets filled in `prepare_preprocess_pipelines`.
172    pub pipeline_id: Option<CachedComputePipelineId>,
173}
174
175/// The pipeline for the batch set count reset shader.
176///
177/// This shader resets the indirect batch set count to 0 for each view. It runs
178/// in between every phase (early, late, and main).
179#[derive(#[automatically_derived]
impl ::core::clone::Clone for ResetIndirectBatchSetsPipeline {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            bind_group_layout: ::core::clone::Clone::clone(&self.bind_group_layout),
            shader: ::core::clone::Clone::clone(&self.shader),
            pipeline_id: ::core::clone::Clone::clone(&self.pipeline_id),
        }
    }
}Clone)]
180pub struct ResetIndirectBatchSetsPipeline {
181    /// The bind group layout for the compute shader.
182    pub bind_group_layout: BindGroupLayoutDescriptor,
183    /// The shader asset handle.
184    pub shader: Handle<Shader>,
185    /// The pipeline ID for the compute shader.
186    ///
187    /// This gets filled in `prepare_preprocess_pipelines`.
188    pub pipeline_id: Option<CachedComputePipelineId>,
189}
190
191/// The pipeline for the indirect parameter building shader.
192#[derive(#[automatically_derived]
impl ::core::clone::Clone for BuildIndirectParametersPipeline {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            bind_group_layout: ::core::clone::Clone::clone(&self.bind_group_layout),
            shader: ::core::clone::Clone::clone(&self.shader),
            pipeline_id: ::core::clone::Clone::clone(&self.pipeline_id),
        }
    }
}Clone)]
193pub struct BuildIndirectParametersPipeline {
194    /// The bind group layout for the compute shader.
195    pub bind_group_layout: BindGroupLayoutDescriptor,
196    /// The shader asset handle.
197    pub shader: Handle<Shader>,
198    /// The pipeline ID for the compute shader.
199    ///
200    /// This gets filled in `prepare_preprocess_pipelines`.
201    pub pipeline_id: Option<CachedComputePipelineId>,
202}
203
204/// The pipeline for the `unpack_bins` compute shader.
205#[derive(#[automatically_derived]
impl ::core::clone::Clone for BinUnpackingPipeline {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            bind_group_layout: ::core::clone::Clone::clone(&self.bind_group_layout),
            shader: ::core::clone::Clone::clone(&self.shader),
            pipeline_id: ::core::clone::Clone::clone(&self.pipeline_id),
        }
    }
}Clone)]
206pub struct BinUnpackingPipeline {
207    /// The layout of the single bind group for that shader.
208    pub bind_group_layout: BindGroupLayoutDescriptor,
209    /// The shader asset handle.
210    pub shader: Handle<Shader>,
211    /// The pipeline ID for the compute shader.
212    ///
213    /// This gets filled in in the [`prepare_preprocess_pipelines`] system.
214    pub pipeline_id: Option<CachedComputePipelineId>,
215}
216
217/// Pipelines for the `allocate_uniforms` compute shader.
218///
219/// This shader has three steps, so we have three pipelines.
220///
221/// Although the `Handle<Shader>` is the same among these three pipelines, they
222/// have to be separate so that the `SpecializedComputePipeline` implementation
223/// on each sub-pipeline can access it.
224#[derive(#[automatically_derived]
impl ::core::clone::Clone for UniformAllocationPipelines {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            local_scan: ::core::clone::Clone::clone(&self.local_scan),
            global_scan: ::core::clone::Clone::clone(&self.global_scan),
            fan: ::core::clone::Clone::clone(&self.fan),
        }
    }
}Clone)]
225pub struct UniformAllocationPipelines {
226    /// The pipeline for step 1: local scan.
227    pub local_scan: UniformAllocationLocalScanPipeline,
228    /// The pipeline for step 2: global scan.
229    pub global_scan: UniformAllocationGlobalScanPipeline,
230    /// The pipeline for step 3: fan.
231    pub fan: UniformAllocationFanPipeline,
232}
233
234/// The pipeline for the first step of the `allocate_uniforms` shader.
235#[derive(#[automatically_derived]
impl ::core::clone::Clone for UniformAllocationLocalScanPipeline {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            bind_group_layout: ::core::clone::Clone::clone(&self.bind_group_layout),
            shader: ::core::clone::Clone::clone(&self.shader),
            pipeline_id_local_scan: ::core::clone::Clone::clone(&self.pipeline_id_local_scan),
        }
    }
}Clone)]
236pub struct UniformAllocationLocalScanPipeline {
237    /// The bind group layout, shared among all the uniform allocation
238    /// pipelines.
239    pub bind_group_layout: BindGroupLayoutDescriptor,
240    /// The shader, also shared among all uniform allocation pipelines.
241    pub shader: Handle<Shader>,
242    /// The pipeline ID for the first step of the `allocate_uniforms` shader.
243    pub pipeline_id_local_scan: Option<CachedComputePipelineId>,
244}
245
246/// The pipeline for the second step of the `allocate_uniforms` shader.
247///
248/// This step is skipped if the number of bins in the batch set is 256 or fewer.
249#[derive(#[automatically_derived]
impl ::core::clone::Clone for UniformAllocationGlobalScanPipeline {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            bind_group_layout: ::core::clone::Clone::clone(&self.bind_group_layout),
            shader: ::core::clone::Clone::clone(&self.shader),
            pipeline_id_global_scan: ::core::clone::Clone::clone(&self.pipeline_id_global_scan),
        }
    }
}Clone)]
250pub struct UniformAllocationGlobalScanPipeline {
251    /// The bind group layout, shared among all the uniform allocation
252    /// pipelines.
253    pub bind_group_layout: BindGroupLayoutDescriptor,
254    /// The shader, also shared among all uniform allocation pipelines.
255    pub shader: Handle<Shader>,
256    /// The pipeline ID for the second step of the `allocate_uniforms` shader.
257    pub pipeline_id_global_scan: Option<CachedComputePipelineId>,
258}
259
260/// The pipeline for the third step of the `allocate_uniforms` shader.
261///
262/// This step is skipped if the number of bins in the batch set is 256 or fewer.
263#[derive(#[automatically_derived]
impl ::core::clone::Clone for UniformAllocationFanPipeline {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            bind_group_layout: ::core::clone::Clone::clone(&self.bind_group_layout),
            shader: ::core::clone::Clone::clone(&self.shader),
            pipeline_id_fan: ::core::clone::Clone::clone(&self.pipeline_id_fan),
        }
    }
}Clone)]
264pub struct UniformAllocationFanPipeline {
265    /// The bind group layout, shared among all the uniform allocation
266    /// pipelines.
267    pub bind_group_layout: BindGroupLayoutDescriptor,
268    /// The shader, also shared among all uniform allocation pipelines.
269    pub shader: Handle<Shader>,
270    /// The pipeline ID for the third step of the `allocate_uniforms` shader.
271    pub pipeline_id_fan: Option<CachedComputePipelineId>,
272}
273
274#[doc = r" Specifies variants of the mesh preprocessing shader."]
pub struct PreprocessPipelineKey(<PreprocessPipelineKey as
    ::bitflags::__private::PublicFlags>::Internal);
#[automatically_derived]
#[doc(hidden)]
unsafe impl ::core::clone::TrivialClone for PreprocessPipelineKey { }
#[automatically_derived]
impl ::core::clone::Clone for PreprocessPipelineKey {
    #[inline]
    fn clone(&self) -> Self {
        let _:
                ::core::clone::AssertParamIsClone<<PreprocessPipelineKey as
                ::bitflags::__private::PublicFlags>::Internal>;
        *self
    }
}
#[automatically_derived]
impl ::core::marker::Copy for PreprocessPipelineKey { }
#[automatically_derived]
impl ::core::marker::StructuralPartialEq for PreprocessPipelineKey { }
#[automatically_derived]
impl ::core::cmp::PartialEq for PreprocessPipelineKey {
    #[inline]
    fn eq(&self, other: &Self) -> bool { self.0 == other.0 }
}
#[automatically_derived]
impl ::core::cmp::Eq for PreprocessPipelineKey {
    #[inline]
    #[doc(hidden)]
    #[coverage(off)]
    fn assert_fields_are_eq(&self) {
        let _:
                ::core::cmp::AssertParamIsEq<<PreprocessPipelineKey as
                ::bitflags::__private::PublicFlags>::Internal>;
    }
}
#[automatically_derived]
impl ::core::hash::Hash for PreprocessPipelineKey {
    #[inline]
    fn hash<__H: ::core::hash::Hasher>(&self, state: &mut __H) {
        ::core::hash::Hash::hash(&self.0, state)
    }
}
#[allow(dead_code, deprecated, unused_doc_comments, unused_attributes,
unused_mut, unused_imports, non_upper_case_globals, clippy :: min_ident_chars,
clippy :: assign_op_pattern, clippy :: indexing_slicing, clippy ::
same_name_method, clippy :: iter_without_into_iter,)]
impl PreprocessPipelineKey {
    #[doc = r" Whether GPU frustum culling is in use."]
    #[doc = r""]
    #[doc = r" This `#define`'s `FRUSTUM_CULLING` in the shader."]
    pub const FRUSTUM_CULLING: Self = Self::from_bits_retain(1);
    #[doc = r" Whether GPU two-phase occlusion culling is in use."]
    #[doc = r""]
    #[doc = r" This `#define`'s `OCCLUSION_CULLING` in the shader."]
    pub const OCCLUSION_CULLING: Self = Self::from_bits_retain(2);
    #[doc =
    r" Whether this is the early phase of GPU two-phase occlusion culling."]
    #[doc = r""]
    #[doc = r" This `#define`'s `EARLY_PHASE` in the shader."]
    pub const EARLY_PHASE: Self = Self::from_bits_retain(4);
}
#[allow(dead_code, deprecated, unused_doc_comments, unused_attributes,
unused_mut, unused_imports, non_upper_case_globals, clippy :: min_ident_chars,
clippy :: assign_op_pattern, clippy :: indexing_slicing, clippy ::
same_name_method, clippy :: iter_without_into_iter,)]
impl ::bitflags::Flags for PreprocessPipelineKey {
    const FLAGS: &'static [::bitflags::Flag<PreprocessPipelineKey>] =
        {
            mod __bitflags_flag_names {
                #[allow(unused_imports)]
                use super::*;
                pub(super) const FRUSTUM_CULLING: &'static str =
                    "FRUSTUM_CULLING";
                pub(super) const OCCLUSION_CULLING: &'static str =
                    "OCCLUSION_CULLING";
                pub(super) const EARLY_PHASE: &'static str = "EARLY_PHASE";
            }
            &[{
                            ::bitflags::Flag::new(__bitflags_flag_names::FRUSTUM_CULLING,
                                PreprocessPipelineKey::FRUSTUM_CULLING)
                        },
                        {
                            ::bitflags::Flag::new(__bitflags_flag_names::OCCLUSION_CULLING,
                                PreprocessPipelineKey::OCCLUSION_CULLING)
                        },
                        {
                            ::bitflags::Flag::new(__bitflags_flag_names::EARLY_PHASE,
                                PreprocessPipelineKey::EARLY_PHASE)
                        }]
        };
    type Bits = u8;
    fn bits(&self) -> u8 { PreprocessPipelineKey::bits(self) }
    fn from_bits_retain(bits: u8) -> PreprocessPipelineKey {
        PreprocessPipelineKey::from_bits_retain(bits)
    }
    fn all_named() -> PreprocessPipelineKey {
        const ALL_NAMED: u8 =
            {
                let mut truncated = <u8 as ::bitflags::Bits>::EMPTY;
                let mut i = 0;
                {
                    {
                        let flag =
                            &<PreprocessPipelineKey as ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                {
                    {
                        let flag =
                            &<PreprocessPipelineKey as ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                {
                    {
                        let flag =
                            &<PreprocessPipelineKey as ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                let _ = i;
                truncated
            };
        PreprocessPipelineKey::from_bits_retain(ALL_NAMED)
    }
}
#[allow(dead_code, deprecated, unused_doc_comments, unused_attributes,
unused_mut, unused_imports, non_upper_case_globals, clippy :: min_ident_chars,
clippy :: assign_op_pattern, clippy :: indexing_slicing, clippy ::
same_name_method, clippy :: iter_without_into_iter,)]
const _: () =
    {
        #[repr(transparent)]
        pub struct InternalBitFlags(u8);
        #[automatically_derived]
        #[doc(hidden)]
        unsafe impl ::core::clone::TrivialClone for InternalBitFlags { }
        #[automatically_derived]
        impl ::core::clone::Clone for InternalBitFlags {
            #[inline]
            fn clone(&self) -> Self {
                let _: ::core::clone::AssertParamIsClone<u8>;
                *self
            }
        }
        #[automatically_derived]
        impl ::core::marker::Copy for InternalBitFlags { }
        #[automatically_derived]
        impl ::core::marker::StructuralPartialEq for InternalBitFlags { }
        #[automatically_derived]
        impl ::core::cmp::PartialEq for InternalBitFlags {
            #[inline]
            fn eq(&self, other: &Self) -> bool { self.0 == other.0 }
        }
        #[automatically_derived]
        impl ::core::cmp::Eq for InternalBitFlags {
            #[inline]
            #[doc(hidden)]
            #[coverage(off)]
            fn assert_fields_are_eq(&self) {
                let _: ::core::cmp::AssertParamIsEq<u8>;
            }
        }
        #[automatically_derived]
        impl ::core::cmp::PartialOrd for InternalBitFlags {
            #[inline]
            fn partial_cmp(&self, other: &Self)
                -> ::core::option::Option<::core::cmp::Ordering> {
                ::core::option::Option::Some(::core::cmp::Ord::cmp(self,
                        other))
            }
        }
        #[automatically_derived]
        impl ::core::cmp::Ord for InternalBitFlags {
            #[inline]
            fn cmp(&self, other: &Self) -> ::core::cmp::Ordering {
                ::core::cmp::Ord::cmp(&self.0, &other.0)
            }
        }
        #[automatically_derived]
        impl ::core::hash::Hash for InternalBitFlags {
            #[inline]
            fn hash<__H: ::core::hash::Hasher>(&self, state: &mut __H) {
                ::core::hash::Hash::hash(&self.0, state)
            }
        }
        impl ::bitflags::__private::PublicFlags for PreprocessPipelineKey {
            type Primitive = u8;
            type Internal = InternalBitFlags;
        }
        impl ::bitflags::__private::core::default::Default for
            InternalBitFlags {
            #[inline]
            fn default() -> Self { InternalBitFlags::empty() }
        }
        impl ::bitflags::__private::core::fmt::Debug for InternalBitFlags {
            fn fmt(&self,
                f: &mut ::bitflags::__private::core::fmt::Formatter<'_>)
                -> ::bitflags::__private::core::fmt::Result {
                if self.is_empty() {
                    f.write_fmt(format_args!("{0:#x}",
                            <u8 as ::bitflags::Bits>::EMPTY))
                } else {
                    ::bitflags::__private::core::fmt::Display::fmt(self, f)
                }
            }
        }
        impl ::bitflags::__private::core::fmt::Display for InternalBitFlags {
            fn fmt(&self,
                f: &mut ::bitflags::__private::core::fmt::Formatter<'_>)
                -> ::bitflags::__private::core::fmt::Result {
                ::bitflags::parser::to_writer(&PreprocessPipelineKey(*self),
                    f)
            }
        }
        impl ::bitflags::__private::core::str::FromStr for InternalBitFlags {
            type Err = ::bitflags::parser::ParseError;
            fn from_str(s: &str)
                ->
                    ::bitflags::__private::core::result::Result<Self,
                    Self::Err> {
                ::bitflags::parser::from_str::<PreprocessPipelineKey>(s).map(|flags|
                        flags.0)
            }
        }
        impl ::bitflags::__private::core::convert::AsRef<u8> for
            InternalBitFlags {
            fn as_ref(&self) -> &u8 { &self.0 }
        }
        impl ::bitflags::__private::core::convert::From<u8> for
            InternalBitFlags {
            fn from(bits: u8) -> Self { Self::from_bits_retain(bits) }
        }
        impl InternalBitFlags {
            /// Get a flags value with all bits unset.
            #[inline]
            pub const fn empty() -> Self {
                Self(<u8 as ::bitflags::Bits>::EMPTY)
            }
            /// Get a flags value with all known bits set.
            #[inline]
            pub const fn all() -> Self {
                const ALL: InternalBitFlags =
                    {
                        let mut truncated = <u8 as ::bitflags::Bits>::EMPTY;
                        let mut _i = 0;
                        {
                            {
                                truncated |=
                                    <PreprocessPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        {
                            {
                                truncated |=
                                    <PreprocessPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        {
                            {
                                truncated |=
                                    <PreprocessPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        InternalBitFlags(truncated)
                    };
                ALL
            }
            /// Get the underlying bits value.
            ///
            /// The returned value is exactly the bits set in this flags value.
            #[inline]
            pub const fn bits(&self) -> u8 { self.0 }
            /// Convert from a bits value.
            ///
            /// This method will return `None` if any unknown bits are set.
            #[inline]
            pub const fn from_bits(bits: u8)
                -> ::bitflags::__private::core::option::Option<Self> {
                let truncated = Self::from_bits_truncate(bits).0;
                if truncated == bits {
                    ::bitflags::__private::core::option::Option::Some(Self(bits))
                } else { ::bitflags::__private::core::option::Option::None }
            }
            /// Convert from a bits value, unsetting any unknown bits.
            #[inline]
            pub const fn from_bits_truncate(bits: u8) -> Self {
                Self(bits & Self::all().0)
            }
            /// Convert from a bits value exactly.
            #[inline]
            pub const fn from_bits_retain(bits: u8) -> Self { Self(bits) }
            /// Get a flags value with the bits of a flag with the given name set.
            ///
            /// This method will return `None` if `name` is empty or doesn't
            /// correspond to any named flag.
            #[inline]
            pub fn from_name(name: &str)
                -> ::bitflags::__private::core::option::Option<Self> {
                mod __bitflags_flag_names {
                    #[allow(unused_imports)]
                    use super::*;
                    pub(super) const FRUSTUM_CULLING: &'static str =
                        "FRUSTUM_CULLING";
                    pub(super) const OCCLUSION_CULLING: &'static str =
                        "OCCLUSION_CULLING";
                    pub(super) const EARLY_PHASE: &'static str = "EARLY_PHASE";
                }
                {
                    {
                        if name == __bitflags_flag_names::FRUSTUM_CULLING {
                            return ::bitflags::__private::core::option::Option::Some(Self(PreprocessPipelineKey::FRUSTUM_CULLING.bits()));
                        }
                    };
                };
                {
                    {
                        if name == __bitflags_flag_names::OCCLUSION_CULLING {
                            return ::bitflags::__private::core::option::Option::Some(Self(PreprocessPipelineKey::OCCLUSION_CULLING.bits()));
                        }
                    };
                };
                {
                    {
                        if name == __bitflags_flag_names::EARLY_PHASE {
                            return ::bitflags::__private::core::option::Option::Some(Self(PreprocessPipelineKey::EARLY_PHASE.bits()));
                        }
                    };
                };
                let _ = name;
                ::bitflags::__private::core::option::Option::None
            }
            /// Whether all bits in `self` are unset.
            #[inline]
            pub const fn is_empty(&self) -> bool {
                self.0 == <u8 as ::bitflags::Bits>::EMPTY
            }
            /// Whether all known bits in this flags value are set.
            #[inline]
            pub const fn is_all(&self) -> bool {
                Self::all().0 | self.0 == self.0
            }
            /// Whether any set bits in `other` are also set in `self`.
            #[inline]
            pub const fn intersects(&self, other: Self) -> bool {
                self.0 & other.0 != <u8 as ::bitflags::Bits>::EMPTY
            }
            /// Whether all set bits in `other` are also set in `self`.
            #[inline]
            pub const fn contains(&self, other: Self) -> bool {
                self.0 & other.0 == other.0
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            pub fn insert(&mut self, other: Self) {
                *self = Self(self.0).union(other);
            }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `remove` won't truncate `other`, but the `!` operator will.
            #[inline]
            pub fn remove(&mut self, other: Self) {
                *self = Self(self.0).difference(other);
            }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            pub fn toggle(&mut self, other: Self) {
                *self = Self(self.0).symmetric_difference(other);
            }
            /// Call `insert` when `value` is `true` or `remove` when `value` is `false`.
            #[inline]
            pub fn set(&mut self, other: Self, value: bool) {
                if value { self.insert(other); } else { self.remove(other); }
            }
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn intersection(self, other: Self) -> Self {
                Self(self.0 & other.0)
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn union(self, other: Self) -> Self {
                Self(self.0 | other.0)
            }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            #[must_use]
            pub const fn difference(self, other: Self) -> Self {
                Self(self.0 & !other.0)
            }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn symmetric_difference(self, other: Self) -> Self {
                Self(self.0 ^ other.0)
            }
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            #[must_use]
            pub const fn complement(self) -> Self {
                Self::from_bits_truncate(!self.0)
            }
        }
        impl ::bitflags::__private::core::fmt::Binary for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Binary::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::Octal for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Octal::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::LowerHex for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::LowerHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::UpperHex for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::UpperHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::ops::BitOr for InternalBitFlags {
            type Output = Self;
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor(self, other: InternalBitFlags) -> Self {
                self.union(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitOrAssign for
            InternalBitFlags {
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor_assign(&mut self, other: Self) { self.insert(other); }
        }
        impl ::bitflags::__private::core::ops::BitXor for InternalBitFlags {
            type Output = Self;
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor(self, other: Self) -> Self {
                self.symmetric_difference(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitXorAssign for
            InternalBitFlags {
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor_assign(&mut self, other: Self) { self.toggle(other); }
        }
        impl ::bitflags::__private::core::ops::BitAnd for InternalBitFlags {
            type Output = Self;
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand(self, other: Self) -> Self { self.intersection(other) }
        }
        impl ::bitflags::__private::core::ops::BitAndAssign for
            InternalBitFlags {
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand_assign(&mut self, other: Self) {
                *self =
                    Self::from_bits_retain(self.bits()).intersection(other);
            }
        }
        impl ::bitflags::__private::core::ops::Sub for InternalBitFlags {
            type Output = Self;
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub(self, other: Self) -> Self { self.difference(other) }
        }
        impl ::bitflags::__private::core::ops::SubAssign for InternalBitFlags
            {
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub_assign(&mut self, other: Self) { self.remove(other); }
        }
        impl ::bitflags::__private::core::ops::Not for InternalBitFlags {
            type Output = Self;
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            fn not(self) -> Self { self.complement() }
        }
        impl ::bitflags::__private::core::iter::Extend<InternalBitFlags> for
            InternalBitFlags {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn extend<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(&mut self, iterator: T) {
                for item in iterator { self.insert(item) }
            }
        }
        impl ::bitflags::__private::core::iter::FromIterator<InternalBitFlags>
            for InternalBitFlags {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn from_iter<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(iterator: T) -> Self {
                use ::bitflags::__private::core::iter::Extend;
                let mut result = Self::empty();
                result.extend(iterator);
                result
            }
        }
        impl InternalBitFlags {
            /// Yield a set of contained flags values.
            ///
            /// Each yielded flags value will correspond to a defined named flag. Any unknown bits
            /// will be yielded together as a final flags value.
            #[inline]
            pub const fn iter(&self)
                -> ::bitflags::iter::Iter<PreprocessPipelineKey> {
                ::bitflags::iter::Iter::__private_const_new(<PreprocessPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    PreprocessPipelineKey::from_bits_retain(self.bits()),
                    PreprocessPipelineKey::from_bits_retain(self.bits()))
            }
            /// Yield a set of contained named flags values.
            ///
            /// This method is like [`iter`](#method.iter), except only yields bits in contained named flags.
            /// Any unknown bits, or bits not corresponding to a contained flag will not be yielded.
            #[inline]
            pub const fn iter_names(&self)
                -> ::bitflags::iter::IterNames<PreprocessPipelineKey> {
                ::bitflags::iter::IterNames::__private_const_new(<PreprocessPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    PreprocessPipelineKey::from_bits_retain(self.bits()),
                    PreprocessPipelineKey::from_bits_retain(self.bits()))
            }
        }
        impl ::bitflags::__private::core::iter::IntoIterator for
            InternalBitFlags {
            type Item = PreprocessPipelineKey;
            type IntoIter = ::bitflags::iter::Iter<PreprocessPipelineKey>;
            fn into_iter(self) -> Self::IntoIter { self.iter() }
        }
        impl InternalBitFlags {
            /// Returns a mutable reference to the raw value of the flags currently stored.
            #[inline]
            pub fn bits_mut(&mut self) -> &mut u8 { &mut self.0 }
        }
        impl ::bitflags::__private::serde::Serialize for InternalBitFlags {
            fn serialize<S: ::bitflags::__private::serde::Serializer>(&self,
                serializer: S)
                ->
                    ::bitflags::__private::core::result::Result<S::Ok,
                    S::Error> {
                ::bitflags::serde::serialize(&PreprocessPipelineKey::from_bits_retain(self.bits()),
                    serializer)
            }
        }
        impl<'de> ::bitflags::__private::serde::Deserialize<'de> for
            InternalBitFlags {
            fn deserialize<D: ::bitflags::__private::serde::Deserializer<'de>>(deserializer:
                    D)
                ->
                    ::bitflags::__private::core::result::Result<Self,
                    D::Error> {
                let flags: PreprocessPipelineKey =
                    ::bitflags::serde::deserialize(deserializer)?;
                ::bitflags::__private::core::result::Result::Ok(flags.0)
            }
        }
        unsafe impl ::bitflags::__private::bytemuck::Pod for InternalBitFlags
            where u8: ::bitflags::__private::bytemuck::Pod {}
        unsafe impl ::bitflags::__private::bytemuck::Zeroable for
            InternalBitFlags where
            u8: ::bitflags::__private::bytemuck::Zeroable {}
        impl PreprocessPipelineKey {
            /// Get a flags value with all bits unset.
            #[inline]
            pub const fn empty() -> Self { Self(InternalBitFlags::empty()) }
            /// Get a flags value with all known bits set.
            #[inline]
            pub const fn all() -> Self { Self(InternalBitFlags::all()) }
            /// Get the underlying bits value.
            ///
            /// The returned value is exactly the bits set in this flags value.
            #[inline]
            pub const fn bits(&self) -> u8 { self.0.bits() }
            /// Convert from a bits value.
            ///
            /// This method will return `None` if any unknown bits are set.
            #[inline]
            pub const fn from_bits(bits: u8)
                -> ::bitflags::__private::core::option::Option<Self> {
                match InternalBitFlags::from_bits(bits) {
                    ::bitflags::__private::core::option::Option::Some(bits) =>
                        ::bitflags::__private::core::option::Option::Some(Self(bits)),
                    ::bitflags::__private::core::option::Option::None =>
                        ::bitflags::__private::core::option::Option::None,
                }
            }
            /// Convert from a bits value, unsetting any unknown bits.
            #[inline]
            pub const fn from_bits_truncate(bits: u8) -> Self {
                Self(InternalBitFlags::from_bits_truncate(bits))
            }
            /// Convert from a bits value exactly.
            #[inline]
            pub const fn from_bits_retain(bits: u8) -> Self {
                Self(InternalBitFlags::from_bits_retain(bits))
            }
            /// Get a flags value with the bits of a flag with the given name set.
            ///
            /// This method will return `None` if `name` is empty or doesn't
            /// correspond to any named flag.
            #[inline]
            pub fn from_name(name: &str)
                -> ::bitflags::__private::core::option::Option<Self> {
                match InternalBitFlags::from_name(name) {
                    ::bitflags::__private::core::option::Option::Some(bits) =>
                        ::bitflags::__private::core::option::Option::Some(Self(bits)),
                    ::bitflags::__private::core::option::Option::None =>
                        ::bitflags::__private::core::option::Option::None,
                }
            }
            /// Whether all bits in `self` are unset.
            #[inline]
            pub const fn is_empty(&self) -> bool { self.0.is_empty() }
            /// Whether all known bits in this flags value are set.
            #[inline]
            pub const fn is_all(&self) -> bool { self.0.is_all() }
            /// Whether any set bits in `other` are also set in `self`.
            #[inline]
            pub const fn intersects(&self, other: Self) -> bool {
                self.0.intersects(other.0)
            }
            /// Whether all set bits in `other` are also set in `self`.
            #[inline]
            pub const fn contains(&self, other: Self) -> bool {
                self.0.contains(other.0)
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            pub fn insert(&mut self, other: Self) { self.0.insert(other.0) }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `remove` won't truncate `other`, but the `!` operator will.
            #[inline]
            pub fn remove(&mut self, other: Self) { self.0.remove(other.0) }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            pub fn toggle(&mut self, other: Self) { self.0.toggle(other.0) }
            /// Call `insert` when `value` is `true` or `remove` when `value` is `false`.
            #[inline]
            pub fn set(&mut self, other: Self, value: bool) {
                self.0.set(other.0, value)
            }
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn intersection(self, other: Self) -> Self {
                Self(self.0.intersection(other.0))
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn union(self, other: Self) -> Self {
                Self(self.0.union(other.0))
            }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            #[must_use]
            pub const fn difference(self, other: Self) -> Self {
                Self(self.0.difference(other.0))
            }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn symmetric_difference(self, other: Self) -> Self {
                Self(self.0.symmetric_difference(other.0))
            }
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            #[must_use]
            pub const fn complement(self) -> Self {
                Self(self.0.complement())
            }
        }
        impl ::bitflags::__private::core::fmt::Binary for
            PreprocessPipelineKey {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Binary::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::Octal for PreprocessPipelineKey
            {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Octal::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::LowerHex for
            PreprocessPipelineKey {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::LowerHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::UpperHex for
            PreprocessPipelineKey {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::UpperHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::ops::BitOr for PreprocessPipelineKey
            {
            type Output = Self;
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor(self, other: PreprocessPipelineKey) -> Self {
                self.union(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitOrAssign for
            PreprocessPipelineKey {
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor_assign(&mut self, other: Self) { self.insert(other); }
        }
        impl ::bitflags::__private::core::ops::BitXor for
            PreprocessPipelineKey {
            type Output = Self;
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor(self, other: Self) -> Self {
                self.symmetric_difference(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitXorAssign for
            PreprocessPipelineKey {
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor_assign(&mut self, other: Self) { self.toggle(other); }
        }
        impl ::bitflags::__private::core::ops::BitAnd for
            PreprocessPipelineKey {
            type Output = Self;
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand(self, other: Self) -> Self { self.intersection(other) }
        }
        impl ::bitflags::__private::core::ops::BitAndAssign for
            PreprocessPipelineKey {
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand_assign(&mut self, other: Self) {
                *self =
                    Self::from_bits_retain(self.bits()).intersection(other);
            }
        }
        impl ::bitflags::__private::core::ops::Sub for PreprocessPipelineKey {
            type Output = Self;
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub(self, other: Self) -> Self { self.difference(other) }
        }
        impl ::bitflags::__private::core::ops::SubAssign for
            PreprocessPipelineKey {
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub_assign(&mut self, other: Self) { self.remove(other); }
        }
        impl ::bitflags::__private::core::ops::Not for PreprocessPipelineKey {
            type Output = Self;
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            fn not(self) -> Self { self.complement() }
        }
        impl ::bitflags::__private::core::iter::Extend<PreprocessPipelineKey>
            for PreprocessPipelineKey {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn extend<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(&mut self, iterator: T) {
                for item in iterator { self.insert(item) }
            }
        }
        impl ::bitflags::__private::core::iter::FromIterator<PreprocessPipelineKey>
            for PreprocessPipelineKey {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn from_iter<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(iterator: T) -> Self {
                use ::bitflags::__private::core::iter::Extend;
                let mut result = Self::empty();
                result.extend(iterator);
                result
            }
        }
        impl PreprocessPipelineKey {
            /// Yield a set of contained flags values.
            ///
            /// Each yielded flags value will correspond to a defined named flag. Any unknown bits
            /// will be yielded together as a final flags value.
            #[inline]
            pub const fn iter(&self)
                -> ::bitflags::iter::Iter<PreprocessPipelineKey> {
                ::bitflags::iter::Iter::__private_const_new(<PreprocessPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    PreprocessPipelineKey::from_bits_retain(self.bits()),
                    PreprocessPipelineKey::from_bits_retain(self.bits()))
            }
            /// Yield a set of contained named flags values.
            ///
            /// This method is like [`iter`](#method.iter), except only yields bits in contained named flags.
            /// Any unknown bits, or bits not corresponding to a contained flag will not be yielded.
            #[inline]
            pub const fn iter_names(&self)
                -> ::bitflags::iter::IterNames<PreprocessPipelineKey> {
                ::bitflags::iter::IterNames::__private_const_new(<PreprocessPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    PreprocessPipelineKey::from_bits_retain(self.bits()),
                    PreprocessPipelineKey::from_bits_retain(self.bits()))
            }
        }
        impl ::bitflags::__private::core::iter::IntoIterator for
            PreprocessPipelineKey {
            type Item = PreprocessPipelineKey;
            type IntoIter = ::bitflags::iter::Iter<PreprocessPipelineKey>;
            fn into_iter(self) -> Self::IntoIter { self.iter() }
        }
    };
#[doc = r" Specifies variants of the indirect parameter building shader."]
pub struct BuildIndirectParametersPipelineKey(<BuildIndirectParametersPipelineKey
    as ::bitflags::__private::PublicFlags>::Internal);
#[automatically_derived]
#[doc(hidden)]
unsafe impl ::core::clone::TrivialClone for BuildIndirectParametersPipelineKey
    {
}
#[automatically_derived]
impl ::core::clone::Clone for BuildIndirectParametersPipelineKey {
    #[inline]
    fn clone(&self) -> Self {
        let _:
                ::core::clone::AssertParamIsClone<<BuildIndirectParametersPipelineKey
                as ::bitflags::__private::PublicFlags>::Internal>;
        *self
    }
}
#[automatically_derived]
impl ::core::marker::Copy for BuildIndirectParametersPipelineKey { }
#[automatically_derived]
impl ::core::marker::StructuralPartialEq for
    BuildIndirectParametersPipelineKey {
}
#[automatically_derived]
impl ::core::cmp::PartialEq for BuildIndirectParametersPipelineKey {
    #[inline]
    fn eq(&self, other: &Self) -> bool { self.0 == other.0 }
}
#[automatically_derived]
impl ::core::cmp::Eq for BuildIndirectParametersPipelineKey {
    #[inline]
    #[doc(hidden)]
    #[coverage(off)]
    fn assert_fields_are_eq(&self) {
        let _:
                ::core::cmp::AssertParamIsEq<<BuildIndirectParametersPipelineKey
                as ::bitflags::__private::PublicFlags>::Internal>;
    }
}
#[automatically_derived]
impl ::core::hash::Hash for BuildIndirectParametersPipelineKey {
    #[inline]
    fn hash<__H: ::core::hash::Hasher>(&self, state: &mut __H) {
        ::core::hash::Hash::hash(&self.0, state)
    }
}
#[allow(dead_code, deprecated, unused_doc_comments, unused_attributes,
unused_mut, unused_imports, non_upper_case_globals, clippy :: min_ident_chars,
clippy :: assign_op_pattern, clippy :: indexing_slicing, clippy ::
same_name_method, clippy :: iter_without_into_iter,)]
impl BuildIndirectParametersPipelineKey {
    #[doc =
    r" Whether the indirect parameter building shader is processing indexed"]
    #[doc = r" meshes (those that have index buffers)."]
    #[doc = r""]
    #[doc = r" This defines `INDEXED` in the shader."]
    pub const INDEXED: Self = Self::from_bits_retain(1);
    #[doc =
    r" Whether the GPU and driver supports `multi_draw_indirect_count`."]
    #[doc = r""]
    #[doc =
    r" This defines `MULTI_DRAW_INDIRECT_COUNT_SUPPORTED` in the shader."]
    pub const MULTI_DRAW_INDIRECT_COUNT_SUPPORTED: Self =
        Self::from_bits_retain(2);
    #[doc = r" Whether GPU two-phase occlusion culling is in use."]
    #[doc = r""]
    #[doc = r" This `#define`'s `OCCLUSION_CULLING` in the shader."]
    pub const OCCLUSION_CULLING: Self = Self::from_bits_retain(4);
    #[doc =
    r" Whether this is the early phase of GPU two-phase occlusion culling."]
    #[doc = r""]
    #[doc = r" This `#define`'s `EARLY_PHASE` in the shader."]
    pub const EARLY_PHASE: Self = Self::from_bits_retain(8);
    #[doc =
    r" Whether this is the late phase of GPU two-phase occlusion culling."]
    #[doc = r""]
    #[doc = r" This `#define`'s `LATE_PHASE` in the shader."]
    pub const LATE_PHASE: Self = Self::from_bits_retain(16);
    #[doc =
    r" Whether this is the phase that runs after the early and late phases,"]
    #[doc = r" and right before the main drawing logic, when GPU two-phase"]
    #[doc = r" occlusion culling is in use."]
    #[doc = r""]
    #[doc = r" This `#define`'s `MAIN_PHASE` in the shader."]
    pub const MAIN_PHASE: Self = Self::from_bits_retain(32);
}
#[allow(dead_code, deprecated, unused_doc_comments, unused_attributes,
unused_mut, unused_imports, non_upper_case_globals, clippy :: min_ident_chars,
clippy :: assign_op_pattern, clippy :: indexing_slicing, clippy ::
same_name_method, clippy :: iter_without_into_iter,)]
impl ::bitflags::Flags for BuildIndirectParametersPipelineKey {
    const FLAGS:
        &'static [::bitflags::Flag<BuildIndirectParametersPipelineKey>] =
        {
            mod __bitflags_flag_names {
                #[allow(unused_imports)]
                use super::*;
                pub(super) const INDEXED: &'static str = "INDEXED";
                pub(super) const MULTI_DRAW_INDIRECT_COUNT_SUPPORTED:
                    &'static str =
                    "MULTI_DRAW_INDIRECT_COUNT_SUPPORTED";
                pub(super) const OCCLUSION_CULLING: &'static str =
                    "OCCLUSION_CULLING";
                pub(super) const EARLY_PHASE: &'static str = "EARLY_PHASE";
                pub(super) const LATE_PHASE: &'static str = "LATE_PHASE";
                pub(super) const MAIN_PHASE: &'static str = "MAIN_PHASE";
            }
            &[{
                            ::bitflags::Flag::new(__bitflags_flag_names::INDEXED,
                                BuildIndirectParametersPipelineKey::INDEXED)
                        },
                        {
                            ::bitflags::Flag::new(__bitflags_flag_names::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED,
                                BuildIndirectParametersPipelineKey::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED)
                        },
                        {
                            ::bitflags::Flag::new(__bitflags_flag_names::OCCLUSION_CULLING,
                                BuildIndirectParametersPipelineKey::OCCLUSION_CULLING)
                        },
                        {
                            ::bitflags::Flag::new(__bitflags_flag_names::EARLY_PHASE,
                                BuildIndirectParametersPipelineKey::EARLY_PHASE)
                        },
                        {
                            ::bitflags::Flag::new(__bitflags_flag_names::LATE_PHASE,
                                BuildIndirectParametersPipelineKey::LATE_PHASE)
                        },
                        {
                            ::bitflags::Flag::new(__bitflags_flag_names::MAIN_PHASE,
                                BuildIndirectParametersPipelineKey::MAIN_PHASE)
                        }]
        };
    type Bits = u8;
    fn bits(&self) -> u8 { BuildIndirectParametersPipelineKey::bits(self) }
    fn from_bits_retain(bits: u8) -> BuildIndirectParametersPipelineKey {
        BuildIndirectParametersPipelineKey::from_bits_retain(bits)
    }
    fn all_named() -> BuildIndirectParametersPipelineKey {
        const ALL_NAMED: u8 =
            {
                let mut truncated = <u8 as ::bitflags::Bits>::EMPTY;
                let mut i = 0;
                {
                    {
                        let flag =
                            &<BuildIndirectParametersPipelineKey as
                                        ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                {
                    {
                        let flag =
                            &<BuildIndirectParametersPipelineKey as
                                        ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                {
                    {
                        let flag =
                            &<BuildIndirectParametersPipelineKey as
                                        ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                {
                    {
                        let flag =
                            &<BuildIndirectParametersPipelineKey as
                                        ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                {
                    {
                        let flag =
                            &<BuildIndirectParametersPipelineKey as
                                        ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                {
                    {
                        let flag =
                            &<BuildIndirectParametersPipelineKey as
                                        ::bitflags::Flags>::FLAGS[i];
                        if flag.is_named() {
                            truncated = truncated | flag.value().bits();
                        }
                        i += 1;
                    }
                };
                let _ = i;
                truncated
            };
        BuildIndirectParametersPipelineKey::from_bits_retain(ALL_NAMED)
    }
}
#[allow(dead_code, deprecated, unused_doc_comments, unused_attributes,
unused_mut, unused_imports, non_upper_case_globals, clippy :: min_ident_chars,
clippy :: assign_op_pattern, clippy :: indexing_slicing, clippy ::
same_name_method, clippy :: iter_without_into_iter,)]
const _: () =
    {
        #[repr(transparent)]
        pub struct InternalBitFlags(u8);
        #[automatically_derived]
        #[doc(hidden)]
        unsafe impl ::core::clone::TrivialClone for InternalBitFlags { }
        #[automatically_derived]
        impl ::core::clone::Clone for InternalBitFlags {
            #[inline]
            fn clone(&self) -> Self {
                let _: ::core::clone::AssertParamIsClone<u8>;
                *self
            }
        }
        #[automatically_derived]
        impl ::core::marker::Copy for InternalBitFlags { }
        #[automatically_derived]
        impl ::core::marker::StructuralPartialEq for InternalBitFlags { }
        #[automatically_derived]
        impl ::core::cmp::PartialEq for InternalBitFlags {
            #[inline]
            fn eq(&self, other: &Self) -> bool { self.0 == other.0 }
        }
        #[automatically_derived]
        impl ::core::cmp::Eq for InternalBitFlags {
            #[inline]
            #[doc(hidden)]
            #[coverage(off)]
            fn assert_fields_are_eq(&self) {
                let _: ::core::cmp::AssertParamIsEq<u8>;
            }
        }
        #[automatically_derived]
        impl ::core::cmp::PartialOrd for InternalBitFlags {
            #[inline]
            fn partial_cmp(&self, other: &Self)
                -> ::core::option::Option<::core::cmp::Ordering> {
                ::core::option::Option::Some(::core::cmp::Ord::cmp(self,
                        other))
            }
        }
        #[automatically_derived]
        impl ::core::cmp::Ord for InternalBitFlags {
            #[inline]
            fn cmp(&self, other: &Self) -> ::core::cmp::Ordering {
                ::core::cmp::Ord::cmp(&self.0, &other.0)
            }
        }
        #[automatically_derived]
        impl ::core::hash::Hash for InternalBitFlags {
            #[inline]
            fn hash<__H: ::core::hash::Hasher>(&self, state: &mut __H) {
                ::core::hash::Hash::hash(&self.0, state)
            }
        }
        impl ::bitflags::__private::PublicFlags for
            BuildIndirectParametersPipelineKey {
            type Primitive = u8;
            type Internal = InternalBitFlags;
        }
        impl ::bitflags::__private::core::default::Default for
            InternalBitFlags {
            #[inline]
            fn default() -> Self { InternalBitFlags::empty() }
        }
        impl ::bitflags::__private::core::fmt::Debug for InternalBitFlags {
            fn fmt(&self,
                f: &mut ::bitflags::__private::core::fmt::Formatter<'_>)
                -> ::bitflags::__private::core::fmt::Result {
                if self.is_empty() {
                    f.write_fmt(format_args!("{0:#x}",
                            <u8 as ::bitflags::Bits>::EMPTY))
                } else {
                    ::bitflags::__private::core::fmt::Display::fmt(self, f)
                }
            }
        }
        impl ::bitflags::__private::core::fmt::Display for InternalBitFlags {
            fn fmt(&self,
                f: &mut ::bitflags::__private::core::fmt::Formatter<'_>)
                -> ::bitflags::__private::core::fmt::Result {
                ::bitflags::parser::to_writer(&BuildIndirectParametersPipelineKey(*self),
                    f)
            }
        }
        impl ::bitflags::__private::core::str::FromStr for InternalBitFlags {
            type Err = ::bitflags::parser::ParseError;
            fn from_str(s: &str)
                ->
                    ::bitflags::__private::core::result::Result<Self,
                    Self::Err> {
                ::bitflags::parser::from_str::<BuildIndirectParametersPipelineKey>(s).map(|flags|
                        flags.0)
            }
        }
        impl ::bitflags::__private::core::convert::AsRef<u8> for
            InternalBitFlags {
            fn as_ref(&self) -> &u8 { &self.0 }
        }
        impl ::bitflags::__private::core::convert::From<u8> for
            InternalBitFlags {
            fn from(bits: u8) -> Self { Self::from_bits_retain(bits) }
        }
        impl InternalBitFlags {
            /// Get a flags value with all bits unset.
            #[inline]
            pub const fn empty() -> Self {
                Self(<u8 as ::bitflags::Bits>::EMPTY)
            }
            /// Get a flags value with all known bits set.
            #[inline]
            pub const fn all() -> Self {
                const ALL: InternalBitFlags =
                    {
                        let mut truncated = <u8 as ::bitflags::Bits>::EMPTY;
                        let mut _i = 0;
                        {
                            {
                                truncated |=
                                    <BuildIndirectParametersPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        {
                            {
                                truncated |=
                                    <BuildIndirectParametersPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        {
                            {
                                truncated |=
                                    <BuildIndirectParametersPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        {
                            {
                                truncated |=
                                    <BuildIndirectParametersPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        {
                            {
                                truncated |=
                                    <BuildIndirectParametersPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        {
                            {
                                truncated |=
                                    <BuildIndirectParametersPipelineKey as
                                                    ::bitflags::Flags>::FLAGS[_i].value().bits();
                                _i += 1;
                            }
                        };
                        InternalBitFlags(truncated)
                    };
                ALL
            }
            /// Get the underlying bits value.
            ///
            /// The returned value is exactly the bits set in this flags value.
            #[inline]
            pub const fn bits(&self) -> u8 { self.0 }
            /// Convert from a bits value.
            ///
            /// This method will return `None` if any unknown bits are set.
            #[inline]
            pub const fn from_bits(bits: u8)
                -> ::bitflags::__private::core::option::Option<Self> {
                let truncated = Self::from_bits_truncate(bits).0;
                if truncated == bits {
                    ::bitflags::__private::core::option::Option::Some(Self(bits))
                } else { ::bitflags::__private::core::option::Option::None }
            }
            /// Convert from a bits value, unsetting any unknown bits.
            #[inline]
            pub const fn from_bits_truncate(bits: u8) -> Self {
                Self(bits & Self::all().0)
            }
            /// Convert from a bits value exactly.
            #[inline]
            pub const fn from_bits_retain(bits: u8) -> Self { Self(bits) }
            /// Get a flags value with the bits of a flag with the given name set.
            ///
            /// This method will return `None` if `name` is empty or doesn't
            /// correspond to any named flag.
            #[inline]
            pub fn from_name(name: &str)
                -> ::bitflags::__private::core::option::Option<Self> {
                mod __bitflags_flag_names {
                    #[allow(unused_imports)]
                    use super::*;
                    pub(super) const INDEXED: &'static str = "INDEXED";
                    pub(super) const MULTI_DRAW_INDIRECT_COUNT_SUPPORTED:
                        &'static str =
                        "MULTI_DRAW_INDIRECT_COUNT_SUPPORTED";
                    pub(super) const OCCLUSION_CULLING: &'static str =
                        "OCCLUSION_CULLING";
                    pub(super) const EARLY_PHASE: &'static str = "EARLY_PHASE";
                    pub(super) const LATE_PHASE: &'static str = "LATE_PHASE";
                    pub(super) const MAIN_PHASE: &'static str = "MAIN_PHASE";
                }
                {
                    {
                        if name == __bitflags_flag_names::INDEXED {
                            return ::bitflags::__private::core::option::Option::Some(Self(BuildIndirectParametersPipelineKey::INDEXED.bits()));
                        }
                    };
                };
                {
                    {
                        if name ==
                                __bitflags_flag_names::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED {
                            return ::bitflags::__private::core::option::Option::Some(Self(BuildIndirectParametersPipelineKey::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED.bits()));
                        }
                    };
                };
                {
                    {
                        if name == __bitflags_flag_names::OCCLUSION_CULLING {
                            return ::bitflags::__private::core::option::Option::Some(Self(BuildIndirectParametersPipelineKey::OCCLUSION_CULLING.bits()));
                        }
                    };
                };
                {
                    {
                        if name == __bitflags_flag_names::EARLY_PHASE {
                            return ::bitflags::__private::core::option::Option::Some(Self(BuildIndirectParametersPipelineKey::EARLY_PHASE.bits()));
                        }
                    };
                };
                {
                    {
                        if name == __bitflags_flag_names::LATE_PHASE {
                            return ::bitflags::__private::core::option::Option::Some(Self(BuildIndirectParametersPipelineKey::LATE_PHASE.bits()));
                        }
                    };
                };
                {
                    {
                        if name == __bitflags_flag_names::MAIN_PHASE {
                            return ::bitflags::__private::core::option::Option::Some(Self(BuildIndirectParametersPipelineKey::MAIN_PHASE.bits()));
                        }
                    };
                };
                let _ = name;
                ::bitflags::__private::core::option::Option::None
            }
            /// Whether all bits in `self` are unset.
            #[inline]
            pub const fn is_empty(&self) -> bool {
                self.0 == <u8 as ::bitflags::Bits>::EMPTY
            }
            /// Whether all known bits in this flags value are set.
            #[inline]
            pub const fn is_all(&self) -> bool {
                Self::all().0 | self.0 == self.0
            }
            /// Whether any set bits in `other` are also set in `self`.
            #[inline]
            pub const fn intersects(&self, other: Self) -> bool {
                self.0 & other.0 != <u8 as ::bitflags::Bits>::EMPTY
            }
            /// Whether all set bits in `other` are also set in `self`.
            #[inline]
            pub const fn contains(&self, other: Self) -> bool {
                self.0 & other.0 == other.0
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            pub fn insert(&mut self, other: Self) {
                *self = Self(self.0).union(other);
            }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `remove` won't truncate `other`, but the `!` operator will.
            #[inline]
            pub fn remove(&mut self, other: Self) {
                *self = Self(self.0).difference(other);
            }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            pub fn toggle(&mut self, other: Self) {
                *self = Self(self.0).symmetric_difference(other);
            }
            /// Call `insert` when `value` is `true` or `remove` when `value` is `false`.
            #[inline]
            pub fn set(&mut self, other: Self, value: bool) {
                if value { self.insert(other); } else { self.remove(other); }
            }
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn intersection(self, other: Self) -> Self {
                Self(self.0 & other.0)
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn union(self, other: Self) -> Self {
                Self(self.0 | other.0)
            }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            #[must_use]
            pub const fn difference(self, other: Self) -> Self {
                Self(self.0 & !other.0)
            }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn symmetric_difference(self, other: Self) -> Self {
                Self(self.0 ^ other.0)
            }
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            #[must_use]
            pub const fn complement(self) -> Self {
                Self::from_bits_truncate(!self.0)
            }
        }
        impl ::bitflags::__private::core::fmt::Binary for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Binary::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::Octal for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Octal::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::LowerHex for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::LowerHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::UpperHex for InternalBitFlags {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::UpperHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::ops::BitOr for InternalBitFlags {
            type Output = Self;
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor(self, other: InternalBitFlags) -> Self {
                self.union(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitOrAssign for
            InternalBitFlags {
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor_assign(&mut self, other: Self) { self.insert(other); }
        }
        impl ::bitflags::__private::core::ops::BitXor for InternalBitFlags {
            type Output = Self;
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor(self, other: Self) -> Self {
                self.symmetric_difference(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitXorAssign for
            InternalBitFlags {
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor_assign(&mut self, other: Self) { self.toggle(other); }
        }
        impl ::bitflags::__private::core::ops::BitAnd for InternalBitFlags {
            type Output = Self;
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand(self, other: Self) -> Self { self.intersection(other) }
        }
        impl ::bitflags::__private::core::ops::BitAndAssign for
            InternalBitFlags {
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand_assign(&mut self, other: Self) {
                *self =
                    Self::from_bits_retain(self.bits()).intersection(other);
            }
        }
        impl ::bitflags::__private::core::ops::Sub for InternalBitFlags {
            type Output = Self;
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub(self, other: Self) -> Self { self.difference(other) }
        }
        impl ::bitflags::__private::core::ops::SubAssign for InternalBitFlags
            {
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub_assign(&mut self, other: Self) { self.remove(other); }
        }
        impl ::bitflags::__private::core::ops::Not for InternalBitFlags {
            type Output = Self;
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            fn not(self) -> Self { self.complement() }
        }
        impl ::bitflags::__private::core::iter::Extend<InternalBitFlags> for
            InternalBitFlags {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn extend<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(&mut self, iterator: T) {
                for item in iterator { self.insert(item) }
            }
        }
        impl ::bitflags::__private::core::iter::FromIterator<InternalBitFlags>
            for InternalBitFlags {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn from_iter<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(iterator: T) -> Self {
                use ::bitflags::__private::core::iter::Extend;
                let mut result = Self::empty();
                result.extend(iterator);
                result
            }
        }
        impl InternalBitFlags {
            /// Yield a set of contained flags values.
            ///
            /// Each yielded flags value will correspond to a defined named flag. Any unknown bits
            /// will be yielded together as a final flags value.
            #[inline]
            pub const fn iter(&self)
                ->
                    ::bitflags::iter::Iter<BuildIndirectParametersPipelineKey> {
                ::bitflags::iter::Iter::__private_const_new(<BuildIndirectParametersPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()),
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()))
            }
            /// Yield a set of contained named flags values.
            ///
            /// This method is like [`iter`](#method.iter), except only yields bits in contained named flags.
            /// Any unknown bits, or bits not corresponding to a contained flag will not be yielded.
            #[inline]
            pub const fn iter_names(&self)
                ->
                    ::bitflags::iter::IterNames<BuildIndirectParametersPipelineKey> {
                ::bitflags::iter::IterNames::__private_const_new(<BuildIndirectParametersPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()),
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()))
            }
        }
        impl ::bitflags::__private::core::iter::IntoIterator for
            InternalBitFlags {
            type Item = BuildIndirectParametersPipelineKey;
            type IntoIter =
                ::bitflags::iter::Iter<BuildIndirectParametersPipelineKey>;
            fn into_iter(self) -> Self::IntoIter { self.iter() }
        }
        impl InternalBitFlags {
            /// Returns a mutable reference to the raw value of the flags currently stored.
            #[inline]
            pub fn bits_mut(&mut self) -> &mut u8 { &mut self.0 }
        }
        impl ::bitflags::__private::serde::Serialize for InternalBitFlags {
            fn serialize<S: ::bitflags::__private::serde::Serializer>(&self,
                serializer: S)
                ->
                    ::bitflags::__private::core::result::Result<S::Ok,
                    S::Error> {
                ::bitflags::serde::serialize(&BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()),
                    serializer)
            }
        }
        impl<'de> ::bitflags::__private::serde::Deserialize<'de> for
            InternalBitFlags {
            fn deserialize<D: ::bitflags::__private::serde::Deserializer<'de>>(deserializer:
                    D)
                ->
                    ::bitflags::__private::core::result::Result<Self,
                    D::Error> {
                let flags: BuildIndirectParametersPipelineKey =
                    ::bitflags::serde::deserialize(deserializer)?;
                ::bitflags::__private::core::result::Result::Ok(flags.0)
            }
        }
        unsafe impl ::bitflags::__private::bytemuck::Pod for InternalBitFlags
            where u8: ::bitflags::__private::bytemuck::Pod {}
        unsafe impl ::bitflags::__private::bytemuck::Zeroable for
            InternalBitFlags where
            u8: ::bitflags::__private::bytemuck::Zeroable {}
        impl BuildIndirectParametersPipelineKey {
            /// Get a flags value with all bits unset.
            #[inline]
            pub const fn empty() -> Self { Self(InternalBitFlags::empty()) }
            /// Get a flags value with all known bits set.
            #[inline]
            pub const fn all() -> Self { Self(InternalBitFlags::all()) }
            /// Get the underlying bits value.
            ///
            /// The returned value is exactly the bits set in this flags value.
            #[inline]
            pub const fn bits(&self) -> u8 { self.0.bits() }
            /// Convert from a bits value.
            ///
            /// This method will return `None` if any unknown bits are set.
            #[inline]
            pub const fn from_bits(bits: u8)
                -> ::bitflags::__private::core::option::Option<Self> {
                match InternalBitFlags::from_bits(bits) {
                    ::bitflags::__private::core::option::Option::Some(bits) =>
                        ::bitflags::__private::core::option::Option::Some(Self(bits)),
                    ::bitflags::__private::core::option::Option::None =>
                        ::bitflags::__private::core::option::Option::None,
                }
            }
            /// Convert from a bits value, unsetting any unknown bits.
            #[inline]
            pub const fn from_bits_truncate(bits: u8) -> Self {
                Self(InternalBitFlags::from_bits_truncate(bits))
            }
            /// Convert from a bits value exactly.
            #[inline]
            pub const fn from_bits_retain(bits: u8) -> Self {
                Self(InternalBitFlags::from_bits_retain(bits))
            }
            /// Get a flags value with the bits of a flag with the given name set.
            ///
            /// This method will return `None` if `name` is empty or doesn't
            /// correspond to any named flag.
            #[inline]
            pub fn from_name(name: &str)
                -> ::bitflags::__private::core::option::Option<Self> {
                match InternalBitFlags::from_name(name) {
                    ::bitflags::__private::core::option::Option::Some(bits) =>
                        ::bitflags::__private::core::option::Option::Some(Self(bits)),
                    ::bitflags::__private::core::option::Option::None =>
                        ::bitflags::__private::core::option::Option::None,
                }
            }
            /// Whether all bits in `self` are unset.
            #[inline]
            pub const fn is_empty(&self) -> bool { self.0.is_empty() }
            /// Whether all known bits in this flags value are set.
            #[inline]
            pub const fn is_all(&self) -> bool { self.0.is_all() }
            /// Whether any set bits in `other` are also set in `self`.
            #[inline]
            pub const fn intersects(&self, other: Self) -> bool {
                self.0.intersects(other.0)
            }
            /// Whether all set bits in `other` are also set in `self`.
            #[inline]
            pub const fn contains(&self, other: Self) -> bool {
                self.0.contains(other.0)
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            pub fn insert(&mut self, other: Self) { self.0.insert(other.0) }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `remove` won't truncate `other`, but the `!` operator will.
            #[inline]
            pub fn remove(&mut self, other: Self) { self.0.remove(other.0) }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            pub fn toggle(&mut self, other: Self) { self.0.toggle(other.0) }
            /// Call `insert` when `value` is `true` or `remove` when `value` is `false`.
            #[inline]
            pub fn set(&mut self, other: Self, value: bool) {
                self.0.set(other.0, value)
            }
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn intersection(self, other: Self) -> Self {
                Self(self.0.intersection(other.0))
            }
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn union(self, other: Self) -> Self {
                Self(self.0.union(other.0))
            }
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            #[must_use]
            pub const fn difference(self, other: Self) -> Self {
                Self(self.0.difference(other.0))
            }
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            #[must_use]
            pub const fn symmetric_difference(self, other: Self) -> Self {
                Self(self.0.symmetric_difference(other.0))
            }
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            #[must_use]
            pub const fn complement(self) -> Self {
                Self(self.0.complement())
            }
        }
        impl ::bitflags::__private::core::fmt::Binary for
            BuildIndirectParametersPipelineKey {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Binary::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::Octal for
            BuildIndirectParametersPipelineKey {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::Octal::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::LowerHex for
            BuildIndirectParametersPipelineKey {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::LowerHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::fmt::UpperHex for
            BuildIndirectParametersPipelineKey {
            fn fmt(&self, f: &mut ::bitflags::__private::core::fmt::Formatter)
                -> ::bitflags::__private::core::fmt::Result {
                let inner = self.0;
                ::bitflags::__private::core::fmt::UpperHex::fmt(&inner, f)
            }
        }
        impl ::bitflags::__private::core::ops::BitOr for
            BuildIndirectParametersPipelineKey {
            type Output = Self;
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor(self, other: BuildIndirectParametersPipelineKey)
                -> Self {
                self.union(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitOrAssign for
            BuildIndirectParametersPipelineKey {
            /// The bitwise or (`|`) of the bits in `self` and `other`.
            #[inline]
            fn bitor_assign(&mut self, other: Self) { self.insert(other); }
        }
        impl ::bitflags::__private::core::ops::BitXor for
            BuildIndirectParametersPipelineKey {
            type Output = Self;
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor(self, other: Self) -> Self {
                self.symmetric_difference(other)
            }
        }
        impl ::bitflags::__private::core::ops::BitXorAssign for
            BuildIndirectParametersPipelineKey {
            /// The bitwise exclusive-or (`^`) of the bits in `self` and `other`.
            #[inline]
            fn bitxor_assign(&mut self, other: Self) { self.toggle(other); }
        }
        impl ::bitflags::__private::core::ops::BitAnd for
            BuildIndirectParametersPipelineKey {
            type Output = Self;
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand(self, other: Self) -> Self { self.intersection(other) }
        }
        impl ::bitflags::__private::core::ops::BitAndAssign for
            BuildIndirectParametersPipelineKey {
            /// The bitwise and (`&`) of the bits in `self` and `other`.
            #[inline]
            fn bitand_assign(&mut self, other: Self) {
                *self =
                    Self::from_bits_retain(self.bits()).intersection(other);
            }
        }
        impl ::bitflags::__private::core::ops::Sub for
            BuildIndirectParametersPipelineKey {
            type Output = Self;
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub(self, other: Self) -> Self { self.difference(other) }
        }
        impl ::bitflags::__private::core::ops::SubAssign for
            BuildIndirectParametersPipelineKey {
            /// The intersection of `self` with the complement of `other` (`&!`).
            ///
            /// This method is not equivalent to `self & !other` when `other` has unknown bits set.
            /// `difference` won't truncate `other`, but the `!` operator will.
            #[inline]
            fn sub_assign(&mut self, other: Self) { self.remove(other); }
        }
        impl ::bitflags::__private::core::ops::Not for
            BuildIndirectParametersPipelineKey {
            type Output = Self;
            /// The bitwise negation (`!`) of the bits in `self`, truncating the result.
            #[inline]
            fn not(self) -> Self { self.complement() }
        }
        impl ::bitflags::__private::core::iter::Extend<BuildIndirectParametersPipelineKey>
            for BuildIndirectParametersPipelineKey {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn extend<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(&mut self, iterator: T) {
                for item in iterator { self.insert(item) }
            }
        }
        impl ::bitflags::__private::core::iter::FromIterator<BuildIndirectParametersPipelineKey>
            for BuildIndirectParametersPipelineKey {
            /// The bitwise or (`|`) of the bits in each flags value.
            fn from_iter<T: ::bitflags::__private::core::iter::IntoIterator<Item
                = Self>>(iterator: T) -> Self {
                use ::bitflags::__private::core::iter::Extend;
                let mut result = Self::empty();
                result.extend(iterator);
                result
            }
        }
        impl BuildIndirectParametersPipelineKey {
            /// Yield a set of contained flags values.
            ///
            /// Each yielded flags value will correspond to a defined named flag. Any unknown bits
            /// will be yielded together as a final flags value.
            #[inline]
            pub const fn iter(&self)
                ->
                    ::bitflags::iter::Iter<BuildIndirectParametersPipelineKey> {
                ::bitflags::iter::Iter::__private_const_new(<BuildIndirectParametersPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()),
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()))
            }
            /// Yield a set of contained named flags values.
            ///
            /// This method is like [`iter`](#method.iter), except only yields bits in contained named flags.
            /// Any unknown bits, or bits not corresponding to a contained flag will not be yielded.
            #[inline]
            pub const fn iter_names(&self)
                ->
                    ::bitflags::iter::IterNames<BuildIndirectParametersPipelineKey> {
                ::bitflags::iter::IterNames::__private_const_new(<BuildIndirectParametersPipelineKey
                        as ::bitflags::Flags>::FLAGS,
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()),
                    BuildIndirectParametersPipelineKey::from_bits_retain(self.bits()))
            }
        }
        impl ::bitflags::__private::core::iter::IntoIterator for
            BuildIndirectParametersPipelineKey {
            type Item = BuildIndirectParametersPipelineKey;
            type IntoIter =
                ::bitflags::iter::Iter<BuildIndirectParametersPipelineKey>;
            fn into_iter(self) -> Self::IntoIter { self.iter() }
        }
    };bitflags! {
275    /// Specifies variants of the mesh preprocessing shader.
276    #[derive(Clone, Copy, PartialEq, Eq, Hash)]
277    pub struct PreprocessPipelineKey: u8 {
278        /// Whether GPU frustum culling is in use.
279        ///
280        /// This `#define`'s `FRUSTUM_CULLING` in the shader.
281        const FRUSTUM_CULLING = 1;
282        /// Whether GPU two-phase occlusion culling is in use.
283        ///
284        /// This `#define`'s `OCCLUSION_CULLING` in the shader.
285        const OCCLUSION_CULLING = 2;
286        /// Whether this is the early phase of GPU two-phase occlusion culling.
287        ///
288        /// This `#define`'s `EARLY_PHASE` in the shader.
289        const EARLY_PHASE = 4;
290    }
291
292    /// Specifies variants of the indirect parameter building shader.
293    #[derive(Clone, Copy, PartialEq, Eq, Hash)]
294    pub struct BuildIndirectParametersPipelineKey: u8 {
295        /// Whether the indirect parameter building shader is processing indexed
296        /// meshes (those that have index buffers).
297        ///
298        /// This defines `INDEXED` in the shader.
299        const INDEXED = 1;
300        /// Whether the GPU and driver supports `multi_draw_indirect_count`.
301        ///
302        /// This defines `MULTI_DRAW_INDIRECT_COUNT_SUPPORTED` in the shader.
303        const MULTI_DRAW_INDIRECT_COUNT_SUPPORTED = 2;
304        /// Whether GPU two-phase occlusion culling is in use.
305        ///
306        /// This `#define`'s `OCCLUSION_CULLING` in the shader.
307        const OCCLUSION_CULLING = 4;
308        /// Whether this is the early phase of GPU two-phase occlusion culling.
309        ///
310        /// This `#define`'s `EARLY_PHASE` in the shader.
311        const EARLY_PHASE = 8;
312        /// Whether this is the late phase of GPU two-phase occlusion culling.
313        ///
314        /// This `#define`'s `LATE_PHASE` in the shader.
315        const LATE_PHASE = 16;
316        /// Whether this is the phase that runs after the early and late phases,
317        /// and right before the main drawing logic, when GPU two-phase
318        /// occlusion culling is in use.
319        ///
320        /// This `#define`'s `MAIN_PHASE` in the shader.
321        const MAIN_PHASE = 32;
322    }
323}
324
325/// The compute shader bind group for the mesh preprocessing pass for each
326/// render phase.
327///
328/// This goes on the view. It maps the [`core::any::TypeId`] of a render phase
329/// (e.g.  [`bevy_core_pipeline::core_3d::Opaque3d`]) to the
330/// [`PhasePreprocessBindGroups`] for that phase.
331#[derive(impl bevy_ecs::component::Component for PreprocessBindGroups where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {
    const STORAGE_TYPE: bevy_ecs::component::StorageType =
        bevy_ecs::component::StorageType::Table;
    type Mutability = bevy_ecs::component::Mutable;
    fn register_required_components(_requiree:
            bevy_ecs::component::ComponentId,
        required_components:
            &mut bevy_ecs::component::RequiredComponentsRegistrator) {}
    fn clone_behavior() -> bevy_ecs::component::ComponentCloneBehavior {
        use bevy_ecs::component::{
            DefaultCloneBehaviorBase, DefaultCloneBehaviorViaClone,
        };
        (&&&bevy_ecs::component::DefaultCloneBehaviorSpecialization::<Self>::default()).default_clone_behavior()
    }
    fn relationship_accessor()
        ->
            ::core::option::Option<bevy_ecs::relationship::ComponentRelationshipAccessor<Self>> {
        ::core::option::Option::None
    }
}Component, #[automatically_derived]
impl ::core::clone::Clone for PreprocessBindGroups {
    #[inline]
    fn clone(&self) -> Self { Self(::core::clone::Clone::clone(&self.0)) }
}Clone, impl ::core::ops::Deref for PreprocessBindGroups {
    type Target = TypeIdHashMap<PhasePreprocessBindGroups>;
    fn deref(&self) -> &Self::Target { &self.0 }
}Deref, impl ::core::ops::DerefMut for PreprocessBindGroups {
    fn deref_mut(&mut self) -> &mut Self::Target { &mut self.0 }
}DerefMut)]
332pub struct PreprocessBindGroups(pub TypeIdHashMap<PhasePreprocessBindGroups>);
333
334/// The compute shader bind group for the mesh preprocessing step for a single
335/// render phase on a single view.
336#[derive(#[automatically_derived]
impl ::core::clone::Clone for PhasePreprocessBindGroups {
    #[inline]
    fn clone(&self) -> Self {
        match self {
            Self::Direct(__self_0) =>
                Self::Direct(::core::clone::Clone::clone(__self_0)),
            Self::IndirectFrustumCulling {
                indexed: __self_0, non_indexed: __self_1 } =>
                Self::IndirectFrustumCulling {
                    indexed: ::core::clone::Clone::clone(__self_0),
                    non_indexed: ::core::clone::Clone::clone(__self_1),
                },
            Self::IndirectOcclusionCulling {
                early_indexed: __self_0,
                early_non_indexed: __self_1,
                late_indexed: __self_2,
                late_non_indexed: __self_3 } =>
                Self::IndirectOcclusionCulling {
                    early_indexed: ::core::clone::Clone::clone(__self_0),
                    early_non_indexed: ::core::clone::Clone::clone(__self_1),
                    late_indexed: ::core::clone::Clone::clone(__self_2),
                    late_non_indexed: ::core::clone::Clone::clone(__self_3),
                },
        }
    }
}Clone)]
337pub enum PhasePreprocessBindGroups {
338    /// The bind group used for the single invocation of the compute shader when
339    /// indirect drawing is *not* being used.
340    ///
341    /// Because direct drawing doesn't require splitting the meshes into indexed
342    /// and non-indexed meshes, there's only one bind group in this case.
343    Direct(BindGroup),
344
345    /// The bind groups used for the compute shader when indirect drawing is
346    /// being used, but occlusion culling isn't being used.
347    ///
348    /// Because indirect drawing requires splitting the meshes into indexed and
349    /// non-indexed meshes, there are two bind groups here.
350    IndirectFrustumCulling {
351        /// The bind group for indexed meshes.
352        indexed: Option<BindGroup>,
353        /// The bind group for non-indexed meshes.
354        non_indexed: Option<BindGroup>,
355    },
356
357    /// The bind groups used for the compute shader when indirect drawing is
358    /// being used, but occlusion culling isn't being used.
359    ///
360    /// Because indirect drawing requires splitting the meshes into indexed and
361    /// non-indexed meshes, and because occlusion culling requires splitting
362    /// this phase into early and late versions, there are four bind groups
363    /// here.
364    IndirectOcclusionCulling {
365        /// The bind group for indexed meshes during the early mesh
366        /// preprocessing phase.
367        early_indexed: Option<BindGroup>,
368        /// The bind group for non-indexed meshes during the early mesh
369        /// preprocessing phase.
370        early_non_indexed: Option<BindGroup>,
371        /// The bind group for indexed meshes during the late mesh preprocessing
372        /// phase.
373        late_indexed: Option<BindGroup>,
374        /// The bind group for non-indexed meshes during the late mesh
375        /// preprocessing phase.
376        late_non_indexed: Option<BindGroup>,
377    },
378}
379
380/// The bind groups for the compute shaders that reset indirect draw counts and
381/// build indirect parameters.
382///
383/// There's one set of bind group for each phase. Phases are keyed off their
384/// [`core::any::TypeId`].
385#[derive(impl bevy_ecs::component::Component for BuildIndirectParametersBindGroups
    where Self: ::core::marker::Send + ::core::marker::Sync + 'static {
    const STORAGE_TYPE: bevy_ecs::component::StorageType =
        bevy_ecs::component::StorageType::SparseSet;
    type Mutability = bevy_ecs::component::Mutable;
    fn register_required_components(_requiree:
            bevy_ecs::component::ComponentId,
        required_components:
            &mut bevy_ecs::component::RequiredComponentsRegistrator) {
        let resource_component_id =
            if let ::core::option::Option::Some(id) =
                    required_components.components_registrator().component_id::<BuildIndirectParametersBindGroups>()
                {
                id
            } else {
                required_components.components_registrator().register_component::<BuildIndirectParametersBindGroups>()
            };
        required_components.register_required::<bevy_ecs::resource::IsResource>(move
                ||
                bevy_ecs::resource::IsResource::new(resource_component_id));
    }
    fn clone_behavior() -> bevy_ecs::component::ComponentCloneBehavior {
        use bevy_ecs::component::{
            DefaultCloneBehaviorBase, DefaultCloneBehaviorViaClone,
        };
        (&&&bevy_ecs::component::DefaultCloneBehaviorSpecialization::<Self>::default()).default_clone_behavior()
    }
    fn relationship_accessor()
        ->
            ::core::option::Option<bevy_ecs::relationship::ComponentRelationshipAccessor<Self>> {
        ::core::option::Option::None
    }
}
impl bevy_ecs::resource::Resource for BuildIndirectParametersBindGroups where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {}Resource, #[automatically_derived]
impl ::core::default::Default for BuildIndirectParametersBindGroups {
    #[inline]
    fn default() -> Self { Self(::core::default::Default::default()) }
}Default, impl ::core::ops::Deref for BuildIndirectParametersBindGroups {
    type Target = TypeIdHashMap<PhaseBuildIndirectParametersBindGroups>;
    fn deref(&self) -> &Self::Target { &self.0 }
}Deref, impl ::core::ops::DerefMut for BuildIndirectParametersBindGroups {
    fn deref_mut(&mut self) -> &mut Self::Target { &mut self.0 }
}DerefMut)]
386pub struct BuildIndirectParametersBindGroups(
387    pub TypeIdHashMap<PhaseBuildIndirectParametersBindGroups>,
388);
389
390impl BuildIndirectParametersBindGroups {
391    /// Creates a new, empty [`BuildIndirectParametersBindGroups`] table.
392    pub fn new() -> BuildIndirectParametersBindGroups {
393        Self::default()
394    }
395}
396
397/// The per-phase set of bind groups for the compute shaders that reset indirect
398/// draw counts and build indirect parameters.
399pub struct PhaseBuildIndirectParametersBindGroups {
400    /// The bind group for the `reset_indirect_batch_sets.wesl` shader, for
401    /// indexed meshes.
402    reset_indexed_indirect_batch_sets: Option<BindGroup>,
403    /// The bind group for the `reset_indirect_batch_sets.wesl` shader, for
404    /// non-indexed meshes.
405    reset_non_indexed_indirect_batch_sets: Option<BindGroup>,
406    /// The bind group for the `build_indirect_params.wesl` shader, for indexed
407    /// meshes.
408    build_indexed_indirect: Option<BindGroup>,
409    /// The bind group for the `build_indirect_params.wesl` shader, for
410    /// non-indexed meshes.
411    build_non_indexed_indirect: Option<BindGroup>,
412}
413
414/// A resource, part of the render world, that stores all the bind groups for
415/// the bin unpacking shader.
416///
417/// There will be one such bind group for each combination of view, phase, and
418/// mesh indexed-ness.
419#[derive(#[automatically_derived]
impl ::core::clone::Clone for BinUnpackingBindGroups {
    #[inline]
    fn clone(&self) -> Self { Self(::core::clone::Clone::clone(&self.0)) }
}Clone, impl bevy_ecs::component::Component for BinUnpackingBindGroups where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {
    const STORAGE_TYPE: bevy_ecs::component::StorageType =
        bevy_ecs::component::StorageType::SparseSet;
    type Mutability = bevy_ecs::component::Mutable;
    fn register_required_components(_requiree:
            bevy_ecs::component::ComponentId,
        required_components:
            &mut bevy_ecs::component::RequiredComponentsRegistrator) {
        let resource_component_id =
            if let ::core::option::Option::Some(id) =
                    required_components.components_registrator().component_id::<BinUnpackingBindGroups>()
                {
                id
            } else {
                required_components.components_registrator().register_component::<BinUnpackingBindGroups>()
            };
        required_components.register_required::<bevy_ecs::resource::IsResource>(move
                ||
                bevy_ecs::resource::IsResource::new(resource_component_id));
    }
    fn clone_behavior() -> bevy_ecs::component::ComponentCloneBehavior {
        use bevy_ecs::component::{
            DefaultCloneBehaviorBase, DefaultCloneBehaviorViaClone,
        };
        (&&&bevy_ecs::component::DefaultCloneBehaviorSpecialization::<Self>::default()).default_clone_behavior()
    }
    fn relationship_accessor()
        ->
            ::core::option::Option<bevy_ecs::relationship::ComponentRelationshipAccessor<Self>> {
        ::core::option::Option::None
    }
}
impl bevy_ecs::resource::Resource for BinUnpackingBindGroups where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {}Resource, #[automatically_derived]
impl ::core::default::Default for BinUnpackingBindGroups {
    #[inline]
    fn default() -> Self { Self(::core::default::Default::default()) }
}Default, impl ::core::ops::Deref for BinUnpackingBindGroups {
    type Target =
        HashMap<SceneUnpackingBuffersKey, ViewPhaseBinUnpackingBindGroups>;
    fn deref(&self) -> &Self::Target { &self.0 }
}Deref, impl ::core::ops::DerefMut for BinUnpackingBindGroups {
    fn deref_mut(&mut self) -> &mut Self::Target { &mut self.0 }
}DerefMut)]
420pub struct BinUnpackingBindGroups(
421    pub HashMap<SceneUnpackingBuffersKey, ViewPhaseBinUnpackingBindGroups>,
422);
423
424/// The bind groups for the `unpack_bins` shader for a single (view, phase)
425/// combination.
426#[derive(#[automatically_derived]
impl ::core::clone::Clone for ViewPhaseBinUnpackingBindGroups {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            indexed: ::core::clone::Clone::clone(&self.indexed),
            non_indexed: ::core::clone::Clone::clone(&self.non_indexed),
        }
    }
}Clone)]
427pub struct ViewPhaseBinUnpackingBindGroups {
428    /// The bind groups for the indexed meshes, one for each batch set.
429    indexed: Vec<ViewPhaseBinUnpackingBindGroup>,
430    /// The bind groups for the non-indexed meshes, one for each batch set.
431    non_indexed: Vec<ViewPhaseBinUnpackingBindGroup>,
432}
433
434/// The bind group for the `unpack_bins` shader for a single combination of
435/// view, phase, and mesh indexed-ness.
436#[derive(#[automatically_derived]
impl ::core::clone::Clone for ViewPhaseBinUnpackingBindGroup {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            metadata_index: ::core::clone::Clone::clone(&self.metadata_index),
            bind_group: ::core::clone::Clone::clone(&self.bind_group),
            mesh_instance_count: ::core::clone::Clone::clone(&self.mesh_instance_count),
        }
    }
}Clone)]
437pub struct ViewPhaseBinUnpackingBindGroup {
438    /// The index of the metadata in the
439    /// [`SceneUnpackingBuffers::bin_unpacking_metadata`] buffer.
440    pub metadata_index: BinUnpackingMetadataIndex,
441    /// The actual shader bind group.
442    pub bind_group: BindGroup,
443    /// The number of mesh instances of the appropriate type (indexed or
444    /// non-indexed) for this batch set.
445    pub mesh_instance_count: u32,
446}
447
448/// A resource, part of the render world, that stores all the bind groups for
449/// the uniform allocation shader.
450///
451/// There will be one such bind group for each combination of view, phase, and
452/// mesh indexed-ness.
453#[derive(#[automatically_derived]
impl ::core::clone::Clone for UniformAllocationBindGroups {
    #[inline]
    fn clone(&self) -> Self { Self(::core::clone::Clone::clone(&self.0)) }
}Clone, impl bevy_ecs::component::Component for UniformAllocationBindGroups where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {
    const STORAGE_TYPE: bevy_ecs::component::StorageType =
        bevy_ecs::component::StorageType::SparseSet;
    type Mutability = bevy_ecs::component::Mutable;
    fn register_required_components(_requiree:
            bevy_ecs::component::ComponentId,
        required_components:
            &mut bevy_ecs::component::RequiredComponentsRegistrator) {
        let resource_component_id =
            if let ::core::option::Option::Some(id) =
                    required_components.components_registrator().component_id::<UniformAllocationBindGroups>()
                {
                id
            } else {
                required_components.components_registrator().register_component::<UniformAllocationBindGroups>()
            };
        required_components.register_required::<bevy_ecs::resource::IsResource>(move
                ||
                bevy_ecs::resource::IsResource::new(resource_component_id));
    }
    fn clone_behavior() -> bevy_ecs::component::ComponentCloneBehavior {
        use bevy_ecs::component::{
            DefaultCloneBehaviorBase, DefaultCloneBehaviorViaClone,
        };
        (&&&bevy_ecs::component::DefaultCloneBehaviorSpecialization::<Self>::default()).default_clone_behavior()
    }
    fn relationship_accessor()
        ->
            ::core::option::Option<bevy_ecs::relationship::ComponentRelationshipAccessor<Self>> {
        ::core::option::Option::None
    }
}
impl bevy_ecs::resource::Resource for UniformAllocationBindGroups where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {}Resource, #[automatically_derived]
impl ::core::default::Default for UniformAllocationBindGroups {
    #[inline]
    fn default() -> Self { Self(::core::default::Default::default()) }
}Default, impl ::core::ops::Deref for UniformAllocationBindGroups {
    type Target =
        HashMap<SceneUnpackingBuffersKey,
        ViewPhaseUniformAllocationBindGroups>;
    fn deref(&self) -> &Self::Target { &self.0 }
}Deref, impl ::core::ops::DerefMut for UniformAllocationBindGroups {
    fn deref_mut(&mut self) -> &mut Self::Target { &mut self.0 }
}DerefMut)]
454pub struct UniformAllocationBindGroups(
455    pub HashMap<SceneUnpackingBuffersKey, ViewPhaseUniformAllocationBindGroups>,
456);
457
458/// The bind groups for the `allocate_uniforms` shader for a single (view,
459/// phase) combination.
460#[derive(#[automatically_derived]
impl ::core::clone::Clone for ViewPhaseUniformAllocationBindGroups {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            indexed: ::core::clone::Clone::clone(&self.indexed),
            non_indexed: ::core::clone::Clone::clone(&self.non_indexed),
        }
    }
}Clone)]
461pub struct ViewPhaseUniformAllocationBindGroups {
462    /// The bind groups for the indexed meshes, one for each batch set.
463    indexed: Vec<ViewPhaseUniformAllocationBindGroup>,
464    /// The bind groups for the non-indexed meshes, one for each batch set.
465    non_indexed: Vec<ViewPhaseUniformAllocationBindGroup>,
466}
467
468/// The bind group for the `allocate_uniforms` shader for a single combination
469/// of view, phase, and mesh indexed-ness.
470#[derive(#[automatically_derived]
impl ::core::clone::Clone for ViewPhaseUniformAllocationBindGroup {
    #[inline]
    fn clone(&self) -> Self {
        Self {
            metadata_index: ::core::clone::Clone::clone(&self.metadata_index),
            bind_group: ::core::clone::Clone::clone(&self.bind_group),
            bin_count: ::core::clone::Clone::clone(&self.bin_count),
        }
    }
}Clone)]
471pub struct ViewPhaseUniformAllocationBindGroup {
472    /// The index of the metadata in the
473    /// [`SceneUnpackingBuffers::uniform_allocation_metadata`] buffer.
474    pub metadata_index: UniformAllocationMetadataIndex,
475    /// The actual shader bind group.
476    pub bind_group: BindGroup,
477    /// The total number of bins in this batch set.
478    pub bin_count: u32,
479}
480
481/// Stops the `GpuPreprocessNode` attempting to generate the buffer for this view
482/// useful to avoid duplicating effort if the bind group is shared between views
483#[derive(impl bevy_ecs::component::Component for SkipGpuPreprocess where
    Self: ::core::marker::Send + ::core::marker::Sync + 'static {
    const STORAGE_TYPE: bevy_ecs::component::StorageType =
        bevy_ecs::component::StorageType::Table;
    type Mutability = bevy_ecs::component::Mutable;
    fn register_required_components(_requiree:
            bevy_ecs::component::ComponentId,
        required_components:
            &mut bevy_ecs::component::RequiredComponentsRegistrator) {}
    fn clone_behavior() -> bevy_ecs::component::ComponentCloneBehavior {
        use bevy_ecs::component::{
            DefaultCloneBehaviorBase, DefaultCloneBehaviorViaClone,
        };
        (&&&bevy_ecs::component::DefaultCloneBehaviorSpecialization::<Self>::default()).default_clone_behavior()
    }
    fn relationship_accessor()
        ->
            ::core::option::Option<bevy_ecs::relationship::ComponentRelationshipAccessor<Self>> {
        ::core::option::Option::None
    }
}Component, #[automatically_derived]
impl ::core::default::Default for SkipGpuPreprocess {
    #[inline]
    fn default() -> Self { Self }
}Default)]
484pub struct SkipGpuPreprocess;
485
486type WithAnyPrepass = Or<(
487    With<DepthPrepass>,
488    With<NormalPrepass>,
489    With<MotionVectorPrepass>,
490    With<DeferredPrepass>,
491)>;
492
493impl Plugin for GpuMeshPreprocessPlugin {
494    fn build(&self, app: &mut App) {
495        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                    "mesh_preprocess.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/render/gpu_preprocess.rs",
                "mesh_preprocess.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! GPU mesh transforming and culling.\n//!\n//! This is a compute shader that expands each `MeshInputUniform` out to a full\n//! `MeshUniform` for each view before rendering. (Thus `MeshInputUniform` and\n//! `MeshUniform` are in a 1:N relationship.) It runs in parallel for all meshes\n//! for all views. As part of this process, the shader gathers each mesh\'s\n//! transform on the previous frame and writes it into the `MeshUniform` so that\n//! TAA works. It also performs frustum culling and occlusion culling, if\n//! requested.\n//!\n//! If occlusion culling is on, this shader runs twice: once to prepare the\n//! meshes that were visible last frame, and once to prepare the meshes that\n//! weren\'t visible last frame but became visible this frame. The two invocations\n//! are known as *early mesh preprocessing* and *late mesh preprocessing*\n//! respectively.\n\nimport bevy_render::occlusion_culling::mesh_preprocess_types::{\n    IndirectParametersMetadata, MeshInput, PreviousMeshInput, PreprocessWorkItem\n};\nimport package::render::mesh_types::{\n    Mesh, MESH_FLAGS_AABB_BASED_VISIBILITY_RANGE_BIT, MESH_FLAGS_NO_FRUSTUM_CULLING_BIT,\n    MESH_FLAGS_VISIBILITY_RANGE_INDEX_BITS\n};\nimport package::render::mesh_view_bindings::view;\nimport package::render::occlusion_culling;\nimport package::prepass::bindings::previous_view_uniforms;\nimport package::render::view_transformations::{\n    position_world_to_ndc, position_world_to_view, ndc_to_uv, view_z_to_depth_ndc,\n    position_world_to_prev_ndc, position_world_to_prev_view, prev_view_z_to_depth_ndc\n};\nimport bevy_render::maths;\nimport bevy_render::view::View;\n\n/// Information about each mesh instance needed to cull it on GPU.\n///\n/// At the moment, this just consists of its axis-aligned bounding box (AABB).\nstruct MeshCullingData {\n    /// The 3D center of the AABB in model space, padded with an extra unused\n    /// float value.\n    aabb_center: vec4<f32>,\n    /// The 3D extents of the AABB in model space, divided by two, padded with\n    /// an extra unused float value.\n    aabb_half_extents: vec4<f32>,\n}\n\n/// The parameters for the indirect compute dispatch for the late mesh\n/// preprocessing phase.\nstruct LatePreprocessWorkItemIndirectParameters {\n    /// The number of workgroups we\'re going to dispatch.\n    ///\n    /// This value should always be equal to `ceil(work_item_count / 64)`.\n    ///\n    /// In the late phase this buffer is bound read-only, and the atomic built-ins\n    /// are only defined for `read_write` storage, so the fields are plain `u32`\n    /// there and `atomic<u32>` in the early phase. `u32` and `atomic<u32>` are\n    /// layout-identical, so the buffer\'s memory layout is unchanged either way.\n    @if(LATE_PHASE)\n    dispatch_x: u32,\n    @else   // LATE_PHASE\n    dispatch_x: atomic<u32>,\n    /// The number of workgroups in the Y direction; always 1.\n    dispatch_y: u32,\n    /// The number of workgroups in the Z direction; always 1.\n    dispatch_z: u32,\n    /// The precise number of work items.\n    @if(LATE_PHASE)\n    work_item_count: u32,\n    @else   // LATE_PHASE\n    work_item_count: atomic<u32>,\n    /// Padding.\n    ///\n    /// This isn\'t the usual structure padding; it\'s needed because some hardware\n    /// requires indirect compute dispatch parameters to be aligned on 64-byte\n    /// boundaries.\n    pad: vec4<u32>,\n}\n\n/// These have to be in a structure because of Naga limitations on DX12.\nstruct Immediates {\n    /// The offset into the `late_preprocess_work_item_indirect_parameters`\n    /// buffer.\n    late_preprocess_work_item_indirect_offset: u32,\n}\n\n/// The current frame\'s `MeshInput`.\n@group(0) @binding(3) var<storage> current_input: array<MeshInput>;\n/// The `MeshInput` values from the previous frame.\n@group(0) @binding(4) var<storage> previous_input: array<PreviousMeshInput>;\n/// Indices into the `MeshInput` buffer.\n///\n/// There may be many indices that map to the same `MeshInput`.\n@group(0) @binding(5) var<storage> work_items: array<PreprocessWorkItem>;\n/// The output array of `Mesh`es.\n@group(0) @binding(6) var<storage, read_write> output: array<Mesh>;\n\n@if(INDIRECT)\n/// The array of indirect parameters for drawcalls.\n@group(0) @binding(7) var<storage, read_write> indirect_parameters_metadata:\n    array<IndirectParametersMetadata>;\n\n@if(FRUSTUM_CULLING) {\n/// Data needed to cull the meshes.\n///\n/// At the moment, this consists only of AABBs.\n@group(0) @binding(9) var<storage> mesh_culling_data: array<MeshCullingData>;\n\n@group(0) @binding(10) var<storage> visibility_ranges: array<vec4<f32>>;\n}\n\n@if(OCCLUSION_CULLING)\n@group(0) @binding(11) var depth_pyramid: texture_2d<f32>;\n\n@if(OCCLUSION_CULLING && EARLY_PHASE) {\n@group(0) @binding(12) var<storage, read_write> late_preprocess_work_items:\n    array<PreprocessWorkItem>;\n\n@group(0) @binding(13) var<storage, read_write> late_preprocess_work_item_indirect_parameters:\n    array<LatePreprocessWorkItemIndirectParameters>;\n}\n\n@if(OCCLUSION_CULLING && LATE_PHASE)\n@group(0) @binding(13) var<storage, read> late_preprocess_work_item_indirect_parameters:\n    array<LatePreprocessWorkItemIndirectParameters>;\n\n@if(OCCLUSION_CULLING)\nvar<immediate> immediates: Immediates;\n\n@if(FRUSTUM_CULLING)\n/// Returns true if the view frustum intersects an oriented bounding box (OBB).\n///\n/// `aabb_center.w` should be 1.0.\nfn view_frustum_intersects_obb(\n    world_from_local: mat4x4<f32>,\n    aabb_center: vec4<f32>,\n    aabb_half_extents: vec3<f32>,\n) -> bool {\n\n    for (var i = 0; i < 5; i += 1) {\n        // Calculate relative radius of the sphere associated with this plane.\n        let plane_normal = view.frustum[i];\n        let relative_radius = dot(\n            abs(\n                vec3(\n                    dot(plane_normal.xyz, world_from_local[0].xyz),\n                    dot(plane_normal.xyz, world_from_local[1].xyz),\n                    dot(plane_normal.xyz, world_from_local[2].xyz),\n                )\n            ),\n            aabb_half_extents\n        );\n\n        // Check the frustum plane.\n        if (!maths::sphere_intersects_plane_half_space(\n                plane_normal, aabb_center, relative_radius)) {\n            return false;\n        }\n    }\n\n    return true;\n}\n\n@compute\n@workgroup_size(64)\nfn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {\n    // Figure out our instance index. If this thread doesn\'t correspond to any\n    // index, bail.\n    let instance_index = global_invocation_id.x;\n\n@if(LATE_PHASE)\n    if (instance_index >= late_preprocess_work_item_indirect_parameters[\n            immediates.late_preprocess_work_item_indirect_offset].work_item_count) {\n        return;\n    }\n@else   // LATE_PHASE\n    if (instance_index >= arrayLength(&work_items)) {\n        return;\n    }\n\n    // Unpack the work item.\n    let input_index = work_items[instance_index].input_index;\n@if(INDIRECT)\n    let indirect_parameters_index = work_items[instance_index].output_or_indirect_parameters_index;\n\n    // If we\'re the first mesh instance in this batch, write the index of our\n    // `MeshInput` into the appropriate slot so that the indirect parameters\n    // building shader can access it.\n@if(INDIRECT && !LATE_PHASE)\n    if (instance_index == 0u) || (work_items[instance_index - 1].output_or_indirect_parameters_index != indirect_parameters_index) {\n        indirect_parameters_metadata[indirect_parameters_index].mesh_index = input_index;\n    }\n\n@if(!INDIRECT)\n    let mesh_output_index = work_items[instance_index].output_or_indirect_parameters_index;\n\n    // Unpack the input matrix.\n    let world_from_local_affine_transpose = current_input[input_index].world_from_local;\n    let world_from_local = maths::affine3_to_square(world_from_local_affine_transpose);\n\n@if(FRUSTUM_CULLING) {\n    // Frustum cull if necessary.\n    if ((current_input[input_index].flags & MESH_FLAGS_NO_FRUSTUM_CULLING_BIT) == 0u) {\n        let aabb_center = mesh_culling_data[input_index].aabb_center.xyz;\n        let aabb_half_extents = mesh_culling_data[input_index].aabb_half_extents.xyz;\n\n        // Do an OBB-based frustum cull.\n        let model_center = world_from_local * vec4(aabb_center, 1.0);\n        if (!view_frustum_intersects_obb(world_from_local, model_center, aabb_half_extents)) {\n            return;\n        }\n    }\n\n    // Visibility range cull if necessary.\n    let visibility_buffer_array_len = arrayLength(&visibility_ranges);\n    let visibility_buffer_index =\n        current_input[input_index].flags & MESH_FLAGS_VISIBILITY_RANGE_INDEX_BITS;\n    if (visibility_buffer_index < visibility_buffer_array_len) {\n        let lod_range = visibility_ranges[visibility_buffer_index];\n\n        // If we\'re using the AABB as the mesh center, determine its world space position.\n        // Otherwise, just use the center of the transform.\n        var world_pos: vec3<f32>;\n        if ((current_input[input_index].flags & MESH_FLAGS_AABB_BASED_VISIBILITY_RANGE_BIT) != 0u) {\n            let aabb_center = mesh_culling_data[input_index].aabb_center.xyz;\n            world_pos = (world_from_local * vec4(aabb_center, 1.0)).xyz;\n        } else {\n            world_pos = world_from_local[3].xyz;\n        }\n\n        let camera_distance = length(world_pos - view.lod_view_world_position);\n        // `x` is the minimum range; `w` is the largest range.\n        if (camera_distance < lod_range.x || camera_distance >= lod_range.w) {\n            return;\n        }\n    }\n}\n\n    // See whether the `MeshInputUniform` was updated on this frame. If it\n    // wasn\'t, then we know the transforms of this mesh must be identical to\n    // those on the previous frame, and therefore we don\'t need to access the\n    // `previous_input_index` (in fact, we can\'t; that index are only valid for\n    // one frame and will be invalid).\n    let timestamp = current_input[input_index].timestamp;\n    let mesh_changed_this_frame = timestamp == view.frame_count;\n\n    // Look up the previous model matrix, if it could have been.\n    let previous_input_index = current_input[input_index].previous_input_index;\n    var previous_world_from_local_affine_transpose: mat3x4<f32>;\n    if (mesh_changed_this_frame && previous_input_index != 0xffffffffu) {\n        previous_world_from_local_affine_transpose =\n            previous_input[previous_input_index].world_from_local;\n    } else {\n        previous_world_from_local_affine_transpose = world_from_local_affine_transpose;\n    }\n    let previous_world_from_local =\n        maths::affine3_to_square(previous_world_from_local_affine_transpose);\n\n    // Occlusion cull if necessary. This is done by calculating the screen-space\n    // axis-aligned bounding box (AABB) of the mesh and testing it against the\n    // appropriate level of the depth pyramid (a.k.a. hierarchical Z-buffer). If\n    // no part of the AABB is in front of the corresponding pixel quad in the\n    // hierarchical Z-buffer, then this mesh must be occluded, and we can skip\n    // rendering it.\n@if(OCCLUSION_CULLING) {\n    let aabb_center = mesh_culling_data[input_index].aabb_center.xyz;\n    let aabb_half_extents = mesh_culling_data[input_index].aabb_half_extents.xyz;\n\n    // Initialize the AABB and the maximum depth.\n    let infinity = bitcast<f32>(0x7f800000u);\n    let neg_infinity = bitcast<f32>(0xff800000u);\n    var aabb = vec4(infinity, infinity, neg_infinity, neg_infinity);\n    var max_depth_view = neg_infinity;\n\n    // Build up the AABB by taking each corner of this mesh\'s OBB, transforming\n    // it, and updating the AABB and depth accordingly.\n    for (var i = 0u; i < 8u; i += 1u) {\n        let local_pos = aabb_center + select(\n            vec3(-1.0),\n            vec3(1.0),\n            vec3((i & 1) != 0, (i & 2) != 0, (i & 4) != 0)\n        ) * aabb_half_extents;\n\n        // If we\'re in the early phase, we\'re testing against the last frame\'s\n        // depth buffer, so we need to use the previous frame\'s transform.\n        @if(EARLY_PHASE)\n        let prev_world_pos = (previous_world_from_local * vec4(local_pos, 1.0)).xyz;\n        @if(EARLY_PHASE)\n        let view_pos = position_world_to_prev_view(prev_world_pos);\n        @if(EARLY_PHASE)\n        let ndc_pos = position_world_to_prev_ndc(prev_world_pos);\n        // Otherwise, if this is the late phase, we use the current frame\'s\n        // transform.\n        @if(!EARLY_PHASE)\n        let world_pos = (world_from_local * vec4(local_pos, 1.0)).xyz;\n        @if(!EARLY_PHASE)\n        let view_pos = position_world_to_view(world_pos);\n        @if(!EARLY_PHASE)\n        let ndc_pos = position_world_to_ndc(world_pos);\n\n        let uv_pos = ndc_to_uv(ndc_pos.xy);\n\n        // Update the AABB and maximum view-space depth.\n        aabb = vec4(min(aabb.xy, uv_pos), max(aabb.zw, uv_pos));\n        max_depth_view = max(max_depth_view, view_pos.z);\n    }\n\n    // Clip to the near plane to avoid the NDC depth becoming negative.\n    @if(EARLY_PHASE)\n    max_depth_view = min(-previous_view_uniforms.clip_from_view[3][2], max_depth_view);\n    @else   // EARLY_PHASE\n    max_depth_view = min(-view.clip_from_view[3][2], max_depth_view);\n\n    // Figure out the depth of the occluder, and compare it to our own depth.\n\n    let aabb_pixel_size = occlusion_culling::get_aabb_size_in_pixels(aabb, depth_pyramid);\n    let occluder_depth_ndc =\n        occlusion_culling::get_occluder_depth(aabb, aabb_pixel_size, depth_pyramid);\n\n    @if(EARLY_PHASE)\n    let max_depth_ndc = prev_view_z_to_depth_ndc(max_depth_view);\n    @else   // EARLY_PHASE\n    let max_depth_ndc = view_z_to_depth_ndc(max_depth_view);\n\n    // Are we culled out?\n    if (max_depth_ndc < occluder_depth_ndc) {\n@if(EARLY_PHASE) {\n        // If this is the early phase, we need to make a note of this mesh so\n        // that we examine it again in the late phase, so that we handle the\n        // case in which a mesh that was invisible last frame became visible in\n        // this frame.\n        let output_work_item_index = atomicAdd(&late_preprocess_work_item_indirect_parameters[\n            immediates.late_preprocess_work_item_indirect_offset].work_item_count, 1u);\n        if (output_work_item_index % 64u == 0u) {\n            // Our workgroup size is 64, and the indirect parameters for the\n            // late mesh preprocessing phase are counted in workgroups, so if\n            // we\'re the first thread in this workgroup, bump the workgroup\n            // count.\n            atomicAdd(&late_preprocess_work_item_indirect_parameters[\n                immediates.late_preprocess_work_item_indirect_offset].dispatch_x, 1u);\n        }\n\n        // Enqueue a work item for the late prepass phase.\n        late_preprocess_work_items[output_work_item_index].input_index = input_index;\n        late_preprocess_work_items[output_work_item_index].output_or_indirect_parameters_index =\n            indirect_parameters_index;\n}\n        // This mesh is culled. Skip it.\n        return;\n    }\n}\n\n    // Calculate inverse transpose.\n    let local_from_world_transpose = transpose(maths::inverse_affine3(transpose(\n        world_from_local_affine_transpose)));\n\n    // Pack inverse transpose.\n    let local_from_world_transpose_a = mat2x4<f32>(\n        vec4<f32>(local_from_world_transpose[0].xyz, local_from_world_transpose[1].x),\n        vec4<f32>(local_from_world_transpose[1].yz, local_from_world_transpose[2].xy));\n    let local_from_world_transpose_b = local_from_world_transpose[2].z;\n\n    // Figure out the output index. In indirect mode, this involves bumping the\n    // instance index in the indirect parameters metadata, which\n    // `build_indirect_params.wesl` will use to generate the actual indirect\n    // parameters. Otherwise, this index was directly supplied to us.\n@if(INDIRECT && LATE_PHASE)\n    let batch_output_index = atomicLoad(\n        &indirect_parameters_metadata[indirect_parameters_index].early_instance_count\n    ) + atomicAdd(\n        &indirect_parameters_metadata[indirect_parameters_index].late_instance_count,\n        1u\n    );\n@if(INDIRECT && !LATE_PHASE)\n    let batch_output_index = atomicAdd(\n        &indirect_parameters_metadata[indirect_parameters_index].early_instance_count,\n        1u\n    );\n\n@if(INDIRECT)\n    let mesh_output_index =\n        indirect_parameters_metadata[indirect_parameters_index].base_output_index +\n        batch_output_index;\n\n    // Write the output.\n    output[mesh_output_index].world_from_local = world_from_local_affine_transpose;\n    output[mesh_output_index].previous_world_from_local =\n        previous_world_from_local_affine_transpose;\n    output[mesh_output_index].local_from_world_transpose_a = local_from_world_transpose_a;\n    output[mesh_output_index].local_from_world_transpose_b = local_from_world_transpose_b;\n    output[mesh_output_index].flags = current_input[input_index].flags;\n    output[mesh_output_index].lightmap_uv_rect = current_input[input_index].lightmap_uv_rect;\n    output[mesh_output_index].first_vertex_index = current_input[input_index].first_vertex_index;\n    output[mesh_output_index].current_skin_index = current_input[input_index].current_skin_index;\n    output[mesh_output_index].material_and_lightmap_bind_group_slot =\n        current_input[input_index].material_and_lightmap_bind_group_slot;\n    output[mesh_output_index].tag = current_input[input_index].tag;\n    output[mesh_output_index].morph_descriptor_index = current_input[input_index].morph_descriptor_index;\n    output[mesh_output_index].metadata_index = current_input[input_index].metadata_index;\n}\n");
    }
};embedded_asset!(app, "mesh_preprocess.wesl");
496        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                    "reset_indirect_batch_sets.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/render/gpu_preprocess.rs",
                "reset_indirect_batch_sets.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! Resets the indirect draw counts to zero.\n//!\n//! This shader is needed because we reuse the same indirect batch set count\n//! buffer (i.e. the buffer that gets passed to `multi_draw_indirect_count` to\n//! determine how many objects to draw) between phases (early, late, and main).\n//! Before launching `build_indirect_params.wesl`, we need to reinitialize the\n//! value to 0.\n\nimport bevy_render::occlusion_culling::mesh_preprocess_types::IndirectBatchSet;\n\n@group(0) @binding(0) var<storage, read_write> indirect_batch_sets: array<IndirectBatchSet>;\n\n@compute\n@workgroup_size(64)\nfn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {\n    // Figure out our instance index. If this thread doesn\'t correspond to any\n    // index, bail.\n    let instance_index = global_invocation_id.x;\n    if (instance_index >= arrayLength(&indirect_batch_sets)) {\n        return;\n    }\n\n    // Reset the number of batch sets to 0.\n    atomicStore(&indirect_batch_sets[instance_index].indirect_parameters_count, 0u);\n}\n");
    }
};embedded_asset!(app, "reset_indirect_batch_sets.wesl");
497        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                    "build_indirect_params.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/render/gpu_preprocess.rs",
                "build_indirect_params.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! Builds GPU indirect draw parameters from metadata.\n//!\n//! This only runs when indirect drawing is enabled. It takes the output of\n//! `mesh_preprocess.wesl` and creates indirect parameters for the GPU.\n//!\n//! This shader runs separately for indexed and non-indexed meshes. Unlike\n//! `mesh_preprocess.wesl`, which runs one instance per mesh *instance*, one\n//! instance of this shader corresponds to a single *batch* which could contain\n//! arbitrarily many instances of a single mesh.\n\nimport bevy_render::occlusion_culling::mesh_preprocess_types::{\n    IndirectBatchSet,\n    IndirectParametersIndexed,\n    IndirectParametersNonIndexed,\n    IndirectParametersMetadata,\n    MeshInput\n};\n\n/// Specifies the batches that this shader invocation is to process.\nstruct IndirectParametersBuildJob {\n    /// The first batch index that this shader invocation should process\n    /// (inclusive).\n    first_batch_index: u32,\n    /// The last batch index that this shader invocation should process\n    /// (exclusive).\n    last_batch_index: u32,\n}\n\n/// The data for each mesh that the CPU supplied to the GPU.\n@group(0) @binding(0) var<storage> current_input: array<MeshInput>;\n\n/// Data that we use to generate the indirect parameters.\n///\n/// The `mesh_preprocess.wesl` shader emits these.\n@group(0) @binding(1) var<storage> indirect_parameters_metadata:\n    array<IndirectParametersMetadata>;\n\n/// Information about each batch set.\n///\n/// A *batch set* is a set of meshes that might be multi-drawn together.\n@group(0) @binding(3) var<storage, read_write> indirect_batch_sets: array<IndirectBatchSet>;\n\n/// Specifies the batches that this shader invocation is to process.\n@group(0) @binding(4) var<uniform> indirect_parameters_build_job: IndirectParametersBuildJob;\n\n@if(INDEXED)\n/// The buffer of indirect draw parameters that we generate, and that the GPU\n/// reads to issue the draws.\n///\n/// This buffer is for indexed meshes.\n@group(0) @binding(5) var<storage, read_write> indirect_parameters:\n    array<IndirectParametersIndexed>;\n@else   // INDEXED\n/// The buffer of indirect draw parameters that we generate, and that the GPU\n/// reads to issue the draws.\n///\n/// This buffer is for non-indexed meshes.\n@group(0) @binding(5) var<storage, read_write> indirect_parameters:\n    array<IndirectParametersNonIndexed>;\n\n@compute\n@workgroup_size(64)\nfn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {\n    // Figure out our instance index (i.e. batch index). If this thread doesn\'t\n    // correspond to a valid index in the range.\n    let instance_index = global_invocation_id.x + indirect_parameters_build_job.first_batch_index;\n    if (instance_index >= indirect_parameters_build_job.last_batch_index) {\n        return;\n    }\n\n    // Unpack the metadata for this batch.\n    let base_output_index = indirect_parameters_metadata[instance_index].base_output_index;\n    let batch_set_index = indirect_parameters_metadata[instance_index].batch_set_index;\n    let mesh_index = indirect_parameters_metadata[instance_index].mesh_index;\n\n    // If we aren\'t using `multi_draw_indirect_count`, we have a 1:1 fixed\n    // assignment of batches to slots in the indirect parameters buffer, so we\n    // can just use the instance index as the index of our indirect parameters.\n    let early_instance_count =\n        indirect_parameters_metadata[instance_index].early_instance_count;\n    let late_instance_count = indirect_parameters_metadata[instance_index].late_instance_count;\n\n    // If in the early phase, we draw only the early meshes. If in the late\n    // phase, we draw only the late meshes. If in the main phase, draw all the\n    // meshes.\n@if(EARLY_PHASE)\n    let instance_count = early_instance_count;\n@elif(LATE_PHASE)\n    let instance_count = late_instance_count;\n@else   // LATE_PHASE\n    let instance_count = early_instance_count + late_instance_count;\n\n    var indirect_parameters_index = instance_index;\n\n    // If the current hardware and driver support `multi_draw_indirect_count`,\n    // dynamically reserve an index for the indirect parameters we\'re to\n    // generate.\n@if(MULTI_DRAW_INDIRECT_COUNT_SUPPORTED)\n    // If this batch belongs to a batch set, then allocate space for the\n    // indirect commands in that batch set.\n    if (batch_set_index != 0xffffffffu) {\n        // Bail out now if there are no instances. Note that we can only bail if\n        // we\'re in a batch set. That\'s because only batch sets are drawn using\n        // `multi_draw_indirect_count`. If we aren\'t using\n        // `multi_draw_indirect_count`, then we need to continue in order to\n        // zero out the instance count; otherwise, it\'ll have garbage data in\n        // it.\n        if (instance_count == 0u) {\n            return;\n        }\n\n        let indirect_parameters_base =\n            indirect_batch_sets[batch_set_index].indirect_parameters_base;\n        let indirect_parameters_offset =\n            atomicAdd(&indirect_batch_sets[batch_set_index].indirect_parameters_count, 1u);\n\n        indirect_parameters_index = indirect_parameters_base + indirect_parameters_offset;\n    }\n\n    // Build up the indirect parameters. The structures for indexed and\n    // non-indexed meshes are slightly different.\n\n    indirect_parameters[indirect_parameters_index].instance_count = instance_count;\n\n@if(LATE_PHASE)\n    // The late mesh instances are stored after the early mesh instances, so we\n    // offset the output index by the number of early mesh instances.\n    indirect_parameters[indirect_parameters_index].first_instance =\n        base_output_index + early_instance_count;\n@else   // LATE_PHASE\n    indirect_parameters[indirect_parameters_index].first_instance = base_output_index;\n\n    indirect_parameters[indirect_parameters_index].base_vertex =\n        current_input[mesh_index].first_vertex_index;\n\n@if(INDEXED)\n    indirect_parameters[indirect_parameters_index].index_count =\n        current_input[mesh_index].index_count;\n@if(INDEXED)\n    indirect_parameters[indirect_parameters_index].first_index =\n        current_input[mesh_index].first_index_index;\n@if(!INDEXED)\n    indirect_parameters[indirect_parameters_index].vertex_count =\n        current_input[mesh_index].index_count;\n}\n");
    }
};embedded_asset!(app, "build_indirect_params.wesl");
498        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                    "unpack_bins.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/render/gpu_preprocess.rs",
                "unpack_bins.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! A compute shader that unpacks bins.\n//!\n//! This shader runs before mesh preprocessing in order to generate\n//! `PreprocessWorkItem`s from cached bins. A single dispatch of this shader\n//! corresponds to a single batch set, and one invocation of this shader\n//! corresponds to one binned entity. Each shader invocation builds the work item\n//! for its entity by copying over the mesh input uniform index and calculating\n//! the position of the command in the indirect parameters buffer that will draw\n//! that entity.\n\nimport bevy_render::occlusion_culling::mesh_preprocess_types::{BinMetadata, PreprocessWorkItem};\n\n/// Information needed to unpack bins belonging to a single batch set.\nstruct BinUnpackingMetadata {\n    /// The index of the first `PreprocessWorkItem` that this compute shader\n    /// dispatch is to write to.\n    base_output_work_item_index: u32,\n    /// The index of the first GPU indirect parameters command for this batch\n    /// set.\n    base_indirect_parameters_index: u32,\n    /// The number of binned mesh instances in the `binned_mesh_instances`\n    /// array.\n    binned_mesh_instance_count: u32,\n    /// Padding.\n    pad_a: u32,\n    /// Padding.\n    pad_b: array<vec4<u32>, 15>,\n};\n\n/// One mesh instance in a bin.\n///\n/// This corresponds to the CPU-side `GpuRenderBinnedMeshInstance` structure.\n///\n/// Note that this structure isn\'t sorted within the\n/// `binned_mesh_instances` buffer. Instances from the same bin aren\'t\n/// guaranteed to be adjacent to one another.\nstruct BinnedMeshInstance {\n    /// The index of the `MeshInputUniform` corresponding to this mesh instance\n    /// in the `MeshInputUniform` buffer.\n    input_uniform_index: u32,\n    /// The index of the bin that this mesh instance belongs to.\n    bin_index: u32,\n};\n\n/// Metadata for the entire batch set.\n@group(0) @binding(0) var<uniform> bin_unpacking_metadata: BinUnpackingMetadata;\n\n/// The input array of `BinnedMeshInstance`s.\n///\n/// Note that this array isn\'t sorted.\n@group(0) @binding(1) var<storage> binned_mesh_instances: array<BinnedMeshInstance>;\n\n/// The output list of `PreprocessWorkItem`s.\n@group(0) @binding(2) var<storage, read_write> preprocess_work_items: array<PreprocessWorkItem>;\n\n/// The bin metadata, which contains the index of the GPU indirect parameters for\n/// this bin, relative to the start of the indirect parameters for this batch\n/// set.\n///\n/// This is indexed by the bin metadata index below.\n@group(0) @binding(3) var<storage> bin_metadata: array<BinMetadata>;\n\n/// A mapping from the bin index to the metadata index within the `bin_metadata`\n/// buffer.\n@group(0) @binding(4) var<storage> bin_index_to_bin_metadata_index: array<u32>;\n\n@compute\n@workgroup_size(64)\nfn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {\n    // Figure out which instance we\'re looking at.\n    let global_id = global_invocation_id.x;\n    if (global_id >= bin_unpacking_metadata.binned_mesh_instance_count) {\n        return;\n    }\n\n    // Unpack the `BinnedMeshInstance`.\n    let input_uniform_index = binned_mesh_instances[global_id].input_uniform_index;\n    let bin_index = binned_mesh_instances[global_id].bin_index;\n\n    // Look up the indirect parameters index for this bin, relative to the first\n    // indirect parameters offset for this batch set.\n    let bin_metadata_index = bin_index_to_bin_metadata_index[bin_index];\n    let indirect_parameters_offset = bin_metadata[bin_metadata_index].indirect_parameters_offset;\n\n    // Determine the location we should write the work item to.\n    let output_index = bin_unpacking_metadata.base_output_work_item_index + global_id;\n\n    // Write out the resulting work item.\n    preprocess_work_items[output_index].input_index = input_uniform_index;\n    preprocess_work_items[output_index].output_or_indirect_parameters_index =\n        bin_unpacking_metadata.base_indirect_parameters_index + indirect_parameters_offset;\n}\n");
    }
};embedded_asset!(app, "unpack_bins.wesl");
499        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                    "allocate_uniforms.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/render/gpu_preprocess.rs",
                "allocate_uniforms.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! A compute shader that allocates `MeshUniform`s.\n//!\n//! This shader runs before mesh preprocessing in order to determine the\n//! positions of `MeshUniform`s. Unlike `MeshInputUniform`s, which are scattered\n//! throughout the buffer, `MeshUniform`s are indexed by instance ID, and so we\n//! must place instances of the same mesh together in the buffer. One dispatch\n//! call corresponds to one batch set (i.e. one multidraw operation), and one\n//! thread corresponds to one bin (a.k.a. draw, a.k.a. batch).\n//!\n//! Essentially, the goal of this shader is to perform a prefix sum, using the\n//! \"scan-then-fan\" approach. It has three phases:\n//!\n//! 1. *Local scan*: Perform a [Hillis-Steele scan] on each chunk of draws, where\n//! the size of each chunk (i.e. the number of draws) is equal to the workgroup\n//! size (256). Write the total size for this chunk to the fan buffer.\n//!\n//! 2. *Global scan*: Do a Hillis-Steele scan on the fan buffer. Now we know the\n//! running total for each chunk.\n//!\n//! 3. *Fan*: Copy the running total for each chunk to every element of that\n//! chunk.\n//!\n//! Note that, for batch sets (i.e. multidraw indirect calls) that have fewer\n//! than 256 batches in them, we only need step (1). This is the common case.\n//!\n//! [Hillis-Steele scan]: https://en.wikipedia.org/wiki/Prefix_sum#Algorithm_1:_Shorter_span,_more_parallel\n\nimport bevy_render::occlusion_culling::mesh_preprocess_types::{BinMetadata, IndirectParametersMetadata};\n\n/// Information needed to allocate `MeshUniform`s.\nstruct UniformAllocationMetadata {\n    /// The index of this batch set in the `IndirectBatchSet` array.\n    ///\n    /// We write this into the `indirect_parameters_metadata`.\n    batch_set_index: u32,\n\n    /// The number of bins (a.k.a. draws, a.k.a. batches) in this batch set.\n    bin_count: u32,\n\n    /// The index of the first set of indirect parameters for this batch set.\n    ///\n    /// This is also the index of the first `IndirectParametersMetadata`, as\n    /// that\'s a parallel array with the indirect parameters.\n    first_indirect_parameters_index: u32,\n\n    /// The index of the first `MeshUniform` slot for this batch set.\n    first_output_mesh_uniform_index: u32,\n\n    /// Padding.\n    pad: array<vec4<f32>, 15u>,\n};\n\n/// The number of threads in a workgroup.\nconst WORKGROUP_SIZE: u32 = 256u;\n\n/// Information needed to allocate `MeshUniform`s.\n@group(0) @binding(0) var<uniform> allocate_uniforms_metadata: UniformAllocationMetadata;\n\n/// Information for each bin, including the indirect parameters offset and the\n/// instance count.\n@group(0) @binding(1) var<storage> bin_metadata: array<BinMetadata>;\n\n/// The array of indirect parameters metadata that we fill out, one for each\n/// batch.\n@group(0) @binding(2) var<storage, read_write> indirect_parameters_metadata:\n    array<IndirectParametersMetadata>;\n\n/// A temporary buffer that stores the mesh uniform index of the last instance\n/// plus one for each workgroup (i.e. for each 256-bin chunk).\n///\n/// This is accumulated in the second stage and written out in the third.\n@group(0) @binding(3) var<storage, read_write> fan_buffer: array<u32>;\n\n/// Scratch memory that stores the prefix sum for every element in our chunk.\nvar<workgroup> output_offsets: array<u32, 256>;\n\n/// The first step of the prefix sum. This computes the prefix sum for each\n/// 256-element chunk.\n///\n/// Note that this will be the *only* step in the operation if the total number\n/// of bins in this batch set is 256 or fewer. Thus we must fill in the indirect\n/// parameters metadata for each batch here, as we can\'t guarantee that the\n/// following two steps will be run at all.\n@compute @workgroup_size(256, 1, 1)\nfn allocate_local_scan(\n    @builtin(local_invocation_id) local_id: vec3<u32>,\n    @builtin(workgroup_id) group_id: vec3<u32>,\n    @builtin(global_invocation_id) global_id: vec3<u32>\n) {\n    let bin_count = allocate_uniforms_metadata.bin_count;\n\n    let block_start = group_id.x * WORKGROUP_SIZE;\n    let block_end = min(block_start + WORKGROUP_SIZE, bin_count);\n\n    // If this is the first workgroup, take the first output index from the\n    // metadata into account. But if this is the second chunk or beyond, don\'t\n    // do that, as the second and third phases will add it in and we don\'t want\n    // to double-count it.\n    if (group_id.x == 0u) {\n        output_offsets[local_id.x] = allocate_uniforms_metadata.first_output_mesh_uniform_index;\n    } else {\n        output_offsets[local_id.x] = 0u;\n    }\n    workgroupBarrier();\n\n    // We\'re doing an inclusive sum, so put the instance count in the *next* bin.\n    if (global_id.x < block_end && local_id.x < WORKGROUP_SIZE - 1u) {\n        output_offsets[local_id.x + 1] = bin_metadata[global_id.x].instance_count;\n    }\n    workgroupBarrier();\n\n    // Prefix sum within our workgroup.\n    hillis_steele_scan(local_id.x);\n\n    // Now write the indirect parameters metadata for this batch. We fill in the\n    // `base_output_index` with the value of the prefix sum (which might be\n    // incomplete if this isn\'t the first chunk). We also populate a few\n    // bookkeeping fields for later rendering passes to use.\n    if (global_id.x < block_end) {\n        let indirect_parameters_offset =\n            allocate_uniforms_metadata.first_indirect_parameters_index +\n            bin_metadata[global_id.x].indirect_parameters_offset;\n        indirect_parameters_metadata[indirect_parameters_offset].base_output_index =\n            output_offsets[local_id.x];\n        indirect_parameters_metadata[indirect_parameters_offset].batch_set_index =\n            allocate_uniforms_metadata.batch_set_index;\n        // These parameters get filled in later. Initialize them to zero for now.\n        // This is required in the case of the early/late instance counts\n        // because the mesh preprocessing shader will atomically increment them.\n        indirect_parameters_metadata[indirect_parameters_offset].mesh_index = 0u;\n        indirect_parameters_metadata[indirect_parameters_offset].early_instance_count = 0u;\n        indirect_parameters_metadata[indirect_parameters_offset].late_instance_count = 0u;\n    }\n\n    // If this is the last element in the workgroup, put the total number of\n    // instances (plus the first output mesh uniform index if we\'re the first\n    // workgroup) in the fan buffer in preparation for the next phase.\n    if (local_id.x == WORKGROUP_SIZE - 1u) {\n        var chunk_total = output_offsets[WORKGROUP_SIZE - 1u];\n        if (global_id.x < block_end) {\n            chunk_total += bin_metadata[global_id.x].instance_count;\n        }\n        fan_buffer[group_id.x] = chunk_total;\n    }\n}\n\n/// The second step of the prefix sum.\n///\n/// This step takes the intermediate fan values computed in the previous step\n/// (i.e. the sum going out of each chunk) and performs one or more Hillis-Steele\n/// scans in order to compute the fan value going into each chunk.\n///\n/// This step is omitted if there are 256 or fewer total draws.\n@compute @workgroup_size(256, 1, 1)\nfn allocate_global_scan(@builtin(local_invocation_id) local_id: vec3<u32>) {\n    var sum = 0u;\n    let chunk_count = div_ceil(allocate_uniforms_metadata.bin_count, WORKGROUP_SIZE);\n\n    // Do a sequential loop over each block of 256 chunks. Because each\n    // iteration of this loop covers 64K meshes, the fact that it\'s sequential\n    // isn\'t going to be a problem in practice.\n    for (var block_start = 0u; block_start < chunk_count; block_start += WORKGROUP_SIZE) {\n        // Set up the Hillis-Steele scan.\n        let block_end = min(block_start + WORKGROUP_SIZE, chunk_count);\n        let global_id = block_start + local_id.x;\n        if (global_id < block_end) {\n            output_offsets[local_id.x] = fan_buffer[global_id];\n        }\n        workgroupBarrier();\n\n        // Perform the scan.\n        hillis_steele_scan(local_id.x);\n\n        // Write the value back.\n        // Note that we don\'t need a workgroup barrier here because\n        // `hillis_steele_scan` already did one.\n        if (global_id < block_end) {\n            fan_buffer[global_id] = sum + output_offsets[local_id.x];\n        }\n\n        // Save the sum coming out of this block for the next one.\n        sum += output_offsets[WORKGROUP_SIZE - 1u];\n        workgroupBarrier();\n    }\n}\n\n/// The third step of the prefix sum.\n///\n/// We take the summed fan value computed in the previous step and add it in to\n/// each value of each chunk beyond the first. We dispatch one fewer workgroup\n/// here than in step (1), because there\'s nothing to do for the first chunk.\n///\n/// This step is omitted if there are 256 or fewer total draws.\n@compute @workgroup_size(256, 1, 1)\nfn allocate_fan(\n    @builtin(workgroup_id) group_id: vec3<u32>,\n    @builtin(global_invocation_id) global_id: vec3<u32>\n) {\n    let id = global_id.x + WORKGROUP_SIZE;\n    let bin_count = allocate_uniforms_metadata.bin_count;\n    if (id >= bin_count) {\n        return;\n    }\n\n    let fan_value = fan_buffer[group_id.x];\n    let indirect_parameters_offset =\n        allocate_uniforms_metadata.first_indirect_parameters_index +\n        bin_metadata[id].indirect_parameters_offset;\n    indirect_parameters_metadata[indirect_parameters_offset].base_output_index += fan_value;\n}\n\n/// Calculates a running exclusive sum.\n/// https://en.wikipedia.org/wiki/Prefix_sum#Algorithm_1:_Shorter_span,_more_parallel\nfn hillis_steele_scan(local_id: u32) {\n    for (var offset = 1u; offset < WORKGROUP_SIZE; offset *= 2u) {\n        var term = 0u;\n        if (local_id >= offset) {\n            term = output_offsets[local_id - offset];\n        }\n        workgroupBarrier();\n        output_offsets[local_id] += term;\n        workgroupBarrier();\n    }\n}\n\n/// Divides unsigned integer a by b, rounding up.\nfn div_ceil(a: u32, b: u32) -> u32 {\n    return (a + b - 1u) / b;\n}\n");
    }
};embedded_asset!(app, "allocate_uniforms.wesl");
500    }
501
502    fn finish(&self, app: &mut App) {
503        let Some(render_app) = app.get_sub_app_mut(RenderApp) else {
504            return;
505        };
506
507        // This plugin does nothing if GPU instance buffer building isn't in
508        // use.
509        let gpu_preprocessing_support = render_app.world().resource::<GpuPreprocessingSupport>();
510        if !self.use_gpu_instance_buffer_builder || !gpu_preprocessing_support.is_available() {
511            return;
512        }
513
514        render_app
515            .init_gpu_resource::<BinUnpackingBindGroups>()
516            .init_gpu_resource::<UniformAllocationBindGroups>()
517            .init_gpu_resource::<PreprocessPipelines>()
518            .init_gpu_resource::<SpecializedComputePipelines<PreprocessPipeline>>()
519            .init_gpu_resource::<SpecializedComputePipelines<ResetIndirectBatchSetsPipeline>>()
520            .init_gpu_resource::<SpecializedComputePipelines<BuildIndirectParametersPipeline>>()
521            .init_gpu_resource::<SpecializedComputePipelines<BinUnpackingPipeline>>()
522            .init_gpu_resource::<SpecializedComputePipelines<UniformAllocationLocalScanPipeline>>()
523            .init_gpu_resource::<SpecializedComputePipelines<UniformAllocationGlobalScanPipeline>>()
524            .init_gpu_resource::<SpecializedComputePipelines<UniformAllocationFanPipeline>>()
525            .add_systems(
526                Render,
527                (
528                    clear_scene_unpacking_buffers.in_set(RenderSystems::PrepareResources),
529                    prepare_preprocess_pipelines.in_set(RenderSystems::Prepare),
530                    prepare_preprocess_bind_groups
531                        .run_if(resource_exists::<BatchedInstanceBuffers<
532                            MeshUniform,
533                            MeshInputUniform
534                        >>)
535                        .in_set(RenderSystems::PrepareBindGroups)
536                        .after(prepare_preprocess_pipelines),
537                    write_mesh_culling_data_buffer.in_set(RenderSystems::PrepareResourcesFlush),
538                ),
539            )
540            .add_systems(
541                Core3d,
542                (
543                    (
544                        allocate_uniforms,
545                        unpack_bins,
546                        early_gpu_preprocess,
547                        early_prepass_build_indirect_parameters.run_if(any_match_filter::<(
548                            With<PreprocessBindGroups>,
549                            Without<SkipGpuPreprocess>,
550                            Without<NoIndirectDrawing>,
551                            Or<(WithAnyPrepass, With<ShadowView>)>,
552                        )>),
553                    )
554                        .chain()
555                        .before(early_prepass),
556                    (
557                        late_gpu_preprocess,
558                        late_prepass_build_indirect_parameters.run_if(any_match_filter::<(
559                            With<PreprocessBindGroups>,
560                            Without<SkipGpuPreprocess>,
561                            Without<NoIndirectDrawing>,
562                            Or<(WithAnyPrepass, With<ShadowView>)>,
563                            With<OcclusionCulling>,
564                        )>),
565                    )
566                        .chain()
567                        .after(early_downsample_depth)
568                        .before(late_prepass),
569                    main_build_indirect_parameters
570                        .run_if(any_match_filter::<(
571                            With<PreprocessBindGroups>,
572                            Without<SkipGpuPreprocess>,
573                            Without<NoIndirectDrawing>,
574                        )>)
575                        .after(late_prepass_build_indirect_parameters)
576                        .after(late_deferred_prepass)
577                        .before(Core3dSystems::MainPass),
578                ),
579            );
580    }
581}
582
583/// A rendering system that invokes a compute shader for each batch set in order
584/// to determine where `MeshUniform`s should be placed.
585///
586/// This shader exists because a single batch set could contain many meshes. By
587/// performing this on the GPU, we avoid having to traverse every visible mesh
588/// on the CPU every frame.
589pub fn allocate_uniforms(
590    current_view: ViewQuery<Option<&ViewLightEntities>, Without<SkipGpuPreprocess>>,
591    view_query: Query<&ExtractedView, Without<SkipGpuPreprocess>>,
592    light_query: Query<&LightEntity>,
593    batched_instance_buffers: Res<BatchedInstanceBuffers<MeshUniform, MeshInputUniform>>,
594    pipeline_cache: Res<PipelineCache>,
595    preprocess_pipelines: Res<PreprocessPipelines>,
596    uniform_allocation_bind_groups: Res<UniformAllocationBindGroups>,
597    mut render_context: RenderContext,
598) {
599    let diagnostics = render_context.diagnostic_recorder();
600    let diagnostics = diagnostics.as_deref();
601
602    // Don't run if the shaders haven't been compiled yet.
603
604    let (
605        Some(uniform_allocation_local_scan_pipeline_id),
606        Some(uniform_allocation_global_scan_pipeline_id),
607        Some(uniform_allocation_fan_pipeline_id),
608    ) = (
609        preprocess_pipelines
610            .uniform_allocation
611            .local_scan
612            .pipeline_id_local_scan,
613        preprocess_pipelines
614            .uniform_allocation
615            .global_scan
616            .pipeline_id_global_scan,
617        preprocess_pipelines.uniform_allocation.fan.pipeline_id_fan,
618    )
619    else {
620        return;
621    };
622
623    let (
624        Some(uniform_allocation_local_scan_pipeline),
625        Some(uniform_allocation_global_scan_pipeline),
626        Some(uniform_allocation_fan_pipeline),
627    ) = (
628        pipeline_cache.get_compute_pipeline(uniform_allocation_local_scan_pipeline_id),
629        pipeline_cache.get_compute_pipeline(uniform_allocation_global_scan_pipeline_id),
630        pipeline_cache.get_compute_pipeline(uniform_allocation_fan_pipeline_id),
631    )
632    else {
633        return;
634    };
635
636    let command_encoder = render_context.command_encoder();
637    let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
638        label: Some("uniform allocation"),
639        timestamp_writes: None,
640    });
641
642    let pass_span = diagnostics.pass_span(&mut compute_pass, "uniform_allocation");
643
644    // Gather up all views.
645    let view_entity = current_view.entity();
646    let shadow_cascade_views = current_view.into_inner();
647    let all_views =
648        gather_shadow_cascades_for_view(view_entity, shadow_cascade_views, &light_query);
649
650    // Loop over each view…
651    for view_entity in all_views {
652        let Ok(view) = view_query.get(view_entity) else {
653            continue;
654        };
655
656        // …and each phase within each view.
657        for phase_type_id in batched_instance_buffers.phase_instance_buffers.keys() {
658            let uniform_allocation_buffers_key = SceneUnpackingBuffersKey {
659                phase: *phase_type_id,
660                view: view.retained_view_entity,
661            };
662
663            // Fetch the bind groups for this (view, phase) combination.
664            let Some(phase_uniform_allocation_bind_groups) =
665                uniform_allocation_bind_groups.get(&uniform_allocation_buffers_key)
666            else {
667                continue;
668            };
669
670            // Invoke the shader for all batch sets corresponding to indexed
671            // meshes and then for all batch sets corresponding to
672            // non-indexed meshes.
673            for uniform_allocation_bind_group in phase_uniform_allocation_bind_groups
674                .indexed
675                .iter()
676                .chain(phase_uniform_allocation_bind_groups.non_indexed.iter())
677            {
678                // Invoke the local scan (step 1).
679                compute_pass.set_pipeline(uniform_allocation_local_scan_pipeline);
680                compute_pass.set_bind_group(0, &uniform_allocation_bind_group.bind_group, &[]);
681                let local_scan_workgroup_count = uniform_allocation_bind_group
682                    .bin_count
683                    .div_ceil(UNIFORM_ALLOCATION_WORKGROUP_SIZE);
684                if local_scan_workgroup_count > 0 {
685                    compute_pass.dispatch_workgroups(local_scan_workgroup_count, 1, 1);
686                }
687
688                // If there are 256 or fewer draws in this batch, we're
689                // done. Otherwise, perform the other two steps.
690                if local_scan_workgroup_count > 1 {
691                    // Invoke the global scan (step 2).
692                    compute_pass.set_pipeline(uniform_allocation_global_scan_pipeline);
693                    compute_pass.dispatch_workgroups(1, 1, 1);
694
695                    // Perform the fan operation (step 3).
696                    compute_pass.set_pipeline(uniform_allocation_fan_pipeline);
697                    let fan_workgroup_count = local_scan_workgroup_count - 1;
698                    compute_pass.dispatch_workgroups(fan_workgroup_count, 1, 1);
699                }
700            }
701        }
702    }
703
704    pass_span.end(&mut compute_pass);
705}
706
707/// A rendering system that invokes a compute shader for each batch set in order
708/// to generate preprocessing jobs for the subsequent mesh preprocessing shader.
709///
710/// This shader exists because performing the unpack operation on the CPU is
711/// slow when there are many entities. By caching the bins on the GPU from frame
712/// to frame, we avoid having to perform a CPU-side traversal of every mesh
713/// instance every frame.
714pub fn unpack_bins(
715    current_view: ViewQuery<Option<&ViewLightEntities>, Without<SkipGpuPreprocess>>,
716    view_query: Query<&ExtractedView, Without<SkipGpuPreprocess>>,
717    light_query: Query<&LightEntity>,
718    batched_instance_buffers: Res<BatchedInstanceBuffers<MeshUniform, MeshInputUniform>>,
719    pipeline_cache: Res<PipelineCache>,
720    preprocess_pipelines: Res<PreprocessPipelines>,
721    bin_unpacking_bind_groups: Res<BinUnpackingBindGroups>,
722    mut render_context: RenderContext,
723) {
724    let diagnostics = render_context.diagnostic_recorder();
725    let diagnostics = diagnostics.as_deref();
726
727    let command_encoder = render_context.command_encoder();
728    let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
729        label: Some("bin unpacking"),
730        timestamp_writes: None,
731    });
732
733    let pass_span = diagnostics.pass_span(&mut compute_pass, "bin_unpacking");
734
735    // Gather up all views.
736    let view_entity = current_view.entity();
737    let shadow_cascade_views = current_view.into_inner();
738    let all_views =
739        gather_shadow_cascades_for_view(view_entity, shadow_cascade_views, &light_query);
740
741    // Don't run if the shaders haven't been compiled yet.
742    if let Some(bin_unpacking_pipeline_id) = preprocess_pipelines.bin_unpacking.pipeline_id
743        && let Some(bin_unpacking_pipeline) =
744            pipeline_cache.get_compute_pipeline(bin_unpacking_pipeline_id)
745    {
746        compute_pass.set_pipeline(bin_unpacking_pipeline);
747
748        // Loop over each view…
749        for view_entity in all_views {
750            let Ok(view) = view_query.get(view_entity) else {
751                continue;
752            };
753
754            // …and each phase within each view.
755            for phase_type_id in batched_instance_buffers.phase_instance_buffers.keys() {
756                let scene_unpacking_buffers_key = SceneUnpackingBuffersKey {
757                    phase: *phase_type_id,
758                    view: view.retained_view_entity,
759                };
760
761                // Fetch the bind groups for this (view, phase) combination.
762                let Some(phase_bin_unpacking_bind_groups) =
763                    bin_unpacking_bind_groups.get(&scene_unpacking_buffers_key)
764                else {
765                    continue;
766                };
767
768                // Invoke the shader for all batch sets corresponding to indexed
769                // meshes and then for all batch sets corresponding to
770                // non-indexed meshes.
771                for bin_unpacking_bind_group in phase_bin_unpacking_bind_groups
772                    .indexed
773                    .iter()
774                    .chain(phase_bin_unpacking_bind_groups.non_indexed.iter())
775                {
776                    compute_pass.set_bind_group(0, &bin_unpacking_bind_group.bind_group, &[]);
777                    let workgroup_count = (bin_unpacking_bind_group.mesh_instance_count as usize)
778                        .div_ceil(WORKGROUP_SIZE);
779                    if workgroup_count > 0 {
780                        compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
781                    }
782                }
783            }
784        }
785    }
786
787    pass_span.end(&mut compute_pass);
788}
789
790pub fn early_gpu_preprocess(
791    current_view: ViewQuery<Option<&ViewLightEntities>, Without<SkipGpuPreprocess>>,
792    view_query: Query<
793        (
794            &ExtractedView,
795            Option<&PreprocessBindGroups>,
796            Option<&ViewUniformOffset>,
797            Has<NoIndirectDrawing>,
798            Has<OcclusionCulling>,
799        ),
800        Without<SkipGpuPreprocess>,
801    >,
802    light_query: Query<&LightEntity>,
803    batched_instance_buffers: Res<BatchedInstanceBuffers<MeshUniform, MeshInputUniform>>,
804    pipeline_cache: Res<PipelineCache>,
805    preprocess_pipelines: Res<PreprocessPipelines>,
806    mut ctx: RenderContext,
807) {
808    let diagnostics = ctx.diagnostic_recorder();
809    let diagnostics = diagnostics.as_deref();
810
811    let command_encoder = ctx.command_encoder();
812
813    let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
814        label: Some("early_mesh_preprocessing"),
815        timestamp_writes: None,
816    });
817
818    let pass_span = diagnostics.pass_span(&mut compute_pass, "early_mesh_preprocessing");
819
820    let view_entity = current_view.entity();
821    let shadow_cascade_views = current_view.into_inner();
822    let all_views =
823        gather_shadow_cascades_for_view(view_entity, shadow_cascade_views, &light_query);
824
825    // Run the compute passes.
826    for view_entity in all_views {
827        let Ok((view, bind_groups, view_uniform_offset, no_indirect_drawing, occlusion_culling)) =
828            view_query.get(view_entity)
829        else {
830            continue;
831        };
832
833        let Some(bind_groups) = bind_groups else {
834            continue;
835        };
836        let Some(view_uniform_offset) = view_uniform_offset else {
837            continue;
838        };
839
840        // Select the right pipeline, depending on whether GPU culling is in
841        // use.
842        let maybe_pipeline_id = if no_indirect_drawing {
843            preprocess_pipelines.direct_preprocess.pipeline_id
844        } else if occlusion_culling {
845            preprocess_pipelines
846                .early_gpu_occlusion_culling_preprocess
847                .pipeline_id
848        } else {
849            preprocess_pipelines
850                .gpu_frustum_culling_preprocess
851                .pipeline_id
852        };
853
854        // Fetch the pipeline.
855        let Some(preprocess_pipeline_id) = maybe_pipeline_id else {
856            {
    use ::tracing::__macro_support::Callsite as _;
    static __CALLSITE: ::tracing::callsite::DefaultCallsite =
        {
            static META: ::tracing::Metadata<'static> =
                {
                    ::tracing_core::metadata::Metadata::new("event src/render/gpu_preprocess.rs:856",
                        "bevy_pbr::render::gpu_preprocess", ::tracing::Level::WARN,
                        ::tracing_core::__macro_support::Option::Some("src/render/gpu_preprocess.rs"),
                        ::tracing_core::__macro_support::Option::Some(856u32),
                        ::tracing_core::__macro_support::Option::Some("bevy_pbr::render::gpu_preprocess"),
                        ::tracing_core::field::FieldSet::new(&["message"],
                            ::tracing_core::callsite::Identifier(&__CALLSITE)),
                        ::tracing::metadata::Kind::EVENT)
                };
            ::tracing::callsite::DefaultCallsite::new(&META)
        };
    let enabled =
        ::tracing::Level::WARN <= ::tracing::level_filters::STATIC_MAX_LEVEL
                &&
                ::tracing::Level::WARN <=
                    ::tracing::level_filters::LevelFilter::current() &&
            {
                let interest = __CALLSITE.interest();
                !interest.is_never() &&
                    ::tracing::__macro_support::__is_enabled(__CALLSITE.metadata(),
                        interest)
            };
    if enabled {
        (|value_set: ::tracing::field::ValueSet|
                    {
                        let meta = __CALLSITE.metadata();
                        ::tracing::Event::dispatch(meta, &value_set);
                        ;
                    })({
                #[allow(unused_imports)]
                use ::tracing::field::{debug, display, Value};
                __CALLSITE.metadata().fields().value_set_all(&[(::tracing::__macro_support::Option::Some(&format_args!("The build mesh uniforms pipeline wasn\'t ready")
                                            as &dyn ::tracing::field::Value))])
            });
    } else { ; }
};warn!("The build mesh uniforms pipeline wasn't ready");
857            continue;
858        };
859
860        let Some(preprocess_pipeline) = pipeline_cache.get_compute_pipeline(preprocess_pipeline_id)
861        else {
862            // This will happen while the pipeline is being compiled and is fine.
863            continue;
864        };
865
866        compute_pass.set_pipeline(preprocess_pipeline);
867
868        // Loop over each render phase.
869        for (phase_type_id, batched_phase_instance_buffers) in
870            &batched_instance_buffers.phase_instance_buffers
871        {
872            // Grab the work item buffers for this view.
873            let Some(work_item_buffers) = batched_phase_instance_buffers
874                .work_item_buffers
875                .get(&view.retained_view_entity)
876            else {
877                continue;
878            };
879
880            // Fetch the bind group for the render phase.
881            let Some(phase_bind_groups) = bind_groups.get(phase_type_id) else {
882                continue;
883            };
884
885            // Make sure the mesh preprocessing shader has access to the
886            // view info it needs to do culling and motion vector
887            // computation.
888            let dynamic_offsets = [view_uniform_offset.offset];
889
890            // Are we drawing directly or indirectly?
891            match *phase_bind_groups {
892                PhasePreprocessBindGroups::Direct(ref bind_group) => {
893                    // Invoke the mesh preprocessing shader to transform
894                    // meshes only, but not cull.
895                    let PreprocessWorkItemBuffers::Direct(work_item_buffer) = work_item_buffers
896                    else {
897                        continue;
898                    };
899                    compute_pass.set_bind_group(0, bind_group, &dynamic_offsets);
900                    let workgroup_count = work_item_buffer.len().div_ceil(WORKGROUP_SIZE);
901                    if workgroup_count > 0 {
902                        compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
903                    }
904                }
905
906                PhasePreprocessBindGroups::IndirectFrustumCulling {
907                    indexed: ref maybe_indexed_bind_group,
908                    non_indexed: ref maybe_non_indexed_bind_group,
909                }
910                | PhasePreprocessBindGroups::IndirectOcclusionCulling {
911                    early_indexed: ref maybe_indexed_bind_group,
912                    early_non_indexed: ref maybe_non_indexed_bind_group,
913                    ..
914                } => {
915                    // Invoke the mesh preprocessing shader to transform and
916                    // cull the meshes.
917                    let PreprocessWorkItemBuffers::Indirect {
918                        indexed: indexed_buffer,
919                        non_indexed: non_indexed_buffer,
920                        ..
921                    } = work_item_buffers
922                    else {
923                        continue;
924                    };
925
926                    // Transform and cull indexed meshes if there are any.
927                    if let Some(indexed_bind_group) = maybe_indexed_bind_group {
928                        if let PreprocessWorkItemBuffers::Indirect {
929                            gpu_occlusion_culling:
930                                Some(GpuOcclusionCullingWorkItemBuffers {
931                                    late_indirect_parameters_indexed_offset,
932                                    ..
933                                }),
934                            ..
935                        } = *work_item_buffers
936                        {
937                            compute_pass.set_immediates(
938                                0,
939                                bytemuck::bytes_of(&late_indirect_parameters_indexed_offset),
940                            );
941                        }
942
943                        compute_pass.set_bind_group(0, indexed_bind_group, &dynamic_offsets);
944                        let workgroup_count = indexed_buffer.len().div_ceil(WORKGROUP_SIZE);
945                        if workgroup_count > 0 {
946                            compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
947                        }
948                    }
949
950                    // Transform and cull non-indexed meshes if there are any.
951                    if let Some(non_indexed_bind_group) = maybe_non_indexed_bind_group {
952                        if let PreprocessWorkItemBuffers::Indirect {
953                            gpu_occlusion_culling:
954                                Some(GpuOcclusionCullingWorkItemBuffers {
955                                    late_indirect_parameters_non_indexed_offset,
956                                    ..
957                                }),
958                            ..
959                        } = *work_item_buffers
960                        {
961                            compute_pass.set_immediates(
962                                0,
963                                bytemuck::bytes_of(&late_indirect_parameters_non_indexed_offset),
964                            );
965                        }
966
967                        compute_pass.set_bind_group(0, non_indexed_bind_group, &dynamic_offsets);
968                        let workgroup_count = non_indexed_buffer.len().div_ceil(WORKGROUP_SIZE);
969                        if workgroup_count > 0 {
970                            compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
971                        }
972                    }
973                }
974            }
975        }
976    }
977
978    pass_span.end(&mut compute_pass);
979}
980
981/// A helper function that returns all the shadow cascades that need to be
982/// rendered for the given view, as well as the view itself.
983fn gather_shadow_cascades_for_view(
984    view_entity: Entity,
985    shadow_cascade_views: Option<&ViewLightEntities>,
986    light_query: &Query<&LightEntity>,
987) -> SmallVec<[Entity; 8]> {
988    let mut all_views: SmallVec<[_; 8]> = SmallVec::new();
989    all_views.push(view_entity);
990    if let Some(shadow_cascade_views) = shadow_cascade_views {
991        all_views.extend(
992            shadow_cascade_views
993                .lights
994                .iter()
995                .filter(|light_entity| {
996                    light_query.get(**light_entity).is_ok_and(|light_entity| {
997                        #[allow(non_exhaustive_omitted_patterns)] match *light_entity {
    LightEntity::Directional { .. } => true,
    _ => false,
}matches!(*light_entity, LightEntity::Directional { .. })
998                    })
999                })
1000                .copied(),
1001        );
1002    }
1003    all_views
1004}
1005
1006pub fn late_gpu_preprocess(
1007    current_view: ViewQuery<
1008        (&ExtractedView, &PreprocessBindGroups, &ViewUniformOffset),
1009        (
1010            Without<SkipGpuPreprocess>,
1011            Without<NoIndirectDrawing>,
1012            With<OcclusionCulling>,
1013            With<DepthPrepass>,
1014        ),
1015    >,
1016    batched_instance_buffers: Res<BatchedInstanceBuffers<MeshUniform, MeshInputUniform>>,
1017    pipeline_cache: Res<PipelineCache>,
1018    preprocess_pipelines: Res<PreprocessPipelines>,
1019    mut ctx: RenderContext,
1020) {
1021    let (view, bind_groups, view_uniform_offset) = current_view.into_inner();
1022
1023    // Fetch the pipeline BEFORE starting diagnostic spans to avoid panic on early return
1024    let maybe_pipeline_id = preprocess_pipelines
1025        .late_gpu_occlusion_culling_preprocess
1026        .pipeline_id;
1027
1028    let Some(preprocess_pipeline_id) = maybe_pipeline_id else {
1029        {
    {
        static SHOULD_FIRE: ::bevy_utils::OnceFlag =
            ::bevy_utils::OnceFlag::new();
        if SHOULD_FIRE.set() {
            {
                use ::tracing::__macro_support::Callsite as _;
                static __CALLSITE: ::tracing::callsite::DefaultCallsite =
                    {
                        static META: ::tracing::Metadata<'static> =
                            {
                                ::tracing_core::metadata::Metadata::new("event src/render/gpu_preprocess.rs:1029",
                                    "bevy_pbr::render::gpu_preprocess", ::tracing::Level::WARN,
                                    ::tracing_core::__macro_support::Option::Some("src/render/gpu_preprocess.rs"),
                                    ::tracing_core::__macro_support::Option::Some(1029u32),
                                    ::tracing_core::__macro_support::Option::Some("bevy_pbr::render::gpu_preprocess"),
                                    ::tracing_core::field::FieldSet::new(&["message"],
                                        ::tracing_core::callsite::Identifier(&__CALLSITE)),
                                    ::tracing::metadata::Kind::EVENT)
                            };
                        ::tracing::callsite::DefaultCallsite::new(&META)
                    };
                let enabled =
                    ::tracing::Level::WARN <=
                                ::tracing::level_filters::STATIC_MAX_LEVEL &&
                            ::tracing::Level::WARN <=
                                ::tracing::level_filters::LevelFilter::current() &&
                        {
                            let interest = __CALLSITE.interest();
                            !interest.is_never() &&
                                ::tracing::__macro_support::__is_enabled(__CALLSITE.metadata(),
                                    interest)
                        };
                if enabled {
                    (|value_set: ::tracing::field::ValueSet|
                                {
                                    let meta = __CALLSITE.metadata();
                                    ::tracing::Event::dispatch(meta, &value_set);
                                    ;
                                })({
                            #[allow(unused_imports)]
                            use ::tracing::field::{debug, display, Value};
                            __CALLSITE.metadata().fields().value_set_all(&[(::tracing::__macro_support::Option::Some(&format_args!("The build mesh uniforms pipeline wasn\'t ready")
                                                        as &dyn ::tracing::field::Value))])
                        });
                } else { ; }
            };
        }
    }
};warn_once!("The build mesh uniforms pipeline wasn't ready");
1030        return;
1031    };
1032
1033    let Some(preprocess_pipeline) = pipeline_cache.get_compute_pipeline(preprocess_pipeline_id)
1034    else {
1035        // This will happen while the pipeline is being compiled and is fine.
1036        return;
1037    };
1038
1039    let diagnostics = ctx.diagnostic_recorder();
1040    let diagnostics = diagnostics.as_deref();
1041
1042    let command_encoder = ctx.command_encoder();
1043
1044    let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
1045        label: Some("late_mesh_preprocessing"),
1046        timestamp_writes: None,
1047    });
1048
1049    let pass_span = diagnostics.pass_span(&mut compute_pass, "late_mesh_preprocessing");
1050
1051    compute_pass.set_pipeline(preprocess_pipeline);
1052
1053    // Loop over each phase. Because we built the phases in parallel,
1054    // each phase has a separate set of instance buffers.
1055    for (phase_type_id, batched_phase_instance_buffers) in
1056        &batched_instance_buffers.phase_instance_buffers
1057    {
1058        let UntypedPhaseBatchedInstanceBuffers {
1059            ref work_item_buffers,
1060            ref late_indexed_indirect_parameters_buffer,
1061            ref late_non_indexed_indirect_parameters_buffer,
1062            ..
1063        } = *batched_phase_instance_buffers;
1064
1065        // Grab the work item buffers for this view.
1066        let Some(phase_work_item_buffers) = work_item_buffers.get(&view.retained_view_entity)
1067        else {
1068            continue;
1069        };
1070
1071        let (
1072            PreprocessWorkItemBuffers::Indirect {
1073                gpu_occlusion_culling:
1074                    Some(GpuOcclusionCullingWorkItemBuffers {
1075                        late_indirect_parameters_indexed_offset,
1076                        late_indirect_parameters_non_indexed_offset,
1077                        ..
1078                    }),
1079                ..
1080            },
1081            Some(PhasePreprocessBindGroups::IndirectOcclusionCulling {
1082                late_indexed: maybe_late_indexed_bind_group,
1083                late_non_indexed: maybe_late_non_indexed_bind_group,
1084                ..
1085            }),
1086            Some(late_indexed_indirect_parameters_buffer),
1087            Some(late_non_indexed_indirect_parameters_buffer),
1088        ) = (
1089            phase_work_item_buffers,
1090            bind_groups.get(phase_type_id),
1091            late_indexed_indirect_parameters_buffer.buffer(),
1092            late_non_indexed_indirect_parameters_buffer.buffer(),
1093        )
1094        else {
1095            continue;
1096        };
1097
1098        let mut dynamic_offsets: SmallVec<[u32; 1]> = ::smallvec::SmallVec::new()smallvec![];
1099        dynamic_offsets.push(view_uniform_offset.offset);
1100
1101        // If there's no space reserved for work items, then don't
1102        // bother doing the dispatch, as there can't possibly be any
1103        // meshes of the given class (indexed or non-indexed) in this
1104        // phase.
1105
1106        // Transform and cull indexed meshes if there are any.
1107        if let Some(late_indexed_bind_group) = maybe_late_indexed_bind_group {
1108            compute_pass.set_immediates(
1109                0,
1110                bytemuck::bytes_of(late_indirect_parameters_indexed_offset),
1111            );
1112
1113            compute_pass.set_bind_group(0, late_indexed_bind_group, &dynamic_offsets);
1114            compute_pass.dispatch_workgroups_indirect(
1115                late_indexed_indirect_parameters_buffer,
1116                (*late_indirect_parameters_indexed_offset as u64)
1117                    * (size_of::<LatePreprocessWorkItemIndirectParameters>() as u64),
1118            );
1119        }
1120
1121        // Transform and cull non-indexed meshes if there are any.
1122        if let Some(late_non_indexed_bind_group) = maybe_late_non_indexed_bind_group {
1123            compute_pass.set_immediates(
1124                0,
1125                bytemuck::bytes_of(late_indirect_parameters_non_indexed_offset),
1126            );
1127
1128            compute_pass.set_bind_group(0, late_non_indexed_bind_group, &dynamic_offsets);
1129            compute_pass.dispatch_workgroups_indirect(
1130                late_non_indexed_indirect_parameters_buffer,
1131                (*late_indirect_parameters_non_indexed_offset as u64)
1132                    * (size_of::<LatePreprocessWorkItemIndirectParameters>() as u64),
1133            );
1134        }
1135    }
1136
1137    pass_span.end(&mut compute_pass);
1138}
1139
1140/// A render graph system, run for each view, that builds the indirect
1141/// parameters for multi-draw indirect calls for the early prepass.
1142///
1143/// The early prepass is the prepass that draws objects that were visible in the
1144/// previous frame.
1145pub fn early_prepass_build_indirect_parameters(
1146    current_view: ViewQuery<&ExtractedView>,
1147    preprocess_pipelines: Res<PreprocessPipelines>,
1148    build_indirect_params_bind_groups: Option<Res<BuildIndirectParametersBindGroups>>,
1149    pipeline_cache: Res<PipelineCache>,
1150    indirect_parameters_buffers: Option<Res<IndirectParametersBuffers>>,
1151    build_indirect_parameters_uniform_indices: Res<BuildIndirectParametersMetadata>,
1152    mut ctx: RenderContext,
1153) {
1154    run_build_indirect_parameters(
1155        &mut ctx,
1156        current_view.into_inner().retained_view_entity,
1157        build_indirect_params_bind_groups.as_deref(),
1158        &pipeline_cache,
1159        indirect_parameters_buffers.as_deref(),
1160        &build_indirect_parameters_uniform_indices,
1161        &preprocess_pipelines.early_phase,
1162        "early_prepass_indirect_parameters_building",
1163    );
1164}
1165
1166/// A render graph system, run for each view, that builds the indirect
1167/// parameters for multi-draw indirect calls for the late prepass.
1168///
1169/// The late prepass is the prepass that draws objects that weren't visible in
1170/// the previous frame but became visible this frame (disocclusions). It'll be
1171/// skipped if occlusion culling is disabled.
1172pub fn late_prepass_build_indirect_parameters(
1173    current_view: ViewQuery<&ExtractedView>,
1174    preprocess_pipelines: Res<PreprocessPipelines>,
1175    build_indirect_params_bind_groups: Option<Res<BuildIndirectParametersBindGroups>>,
1176    pipeline_cache: Res<PipelineCache>,
1177    indirect_parameters_buffers: Option<Res<IndirectParametersBuffers>>,
1178    build_indirect_parameters_uniform_indices: Res<BuildIndirectParametersMetadata>,
1179    mut ctx: RenderContext,
1180) {
1181    run_build_indirect_parameters(
1182        &mut ctx,
1183        current_view.into_inner().retained_view_entity,
1184        build_indirect_params_bind_groups.as_deref(),
1185        &pipeline_cache,
1186        indirect_parameters_buffers.as_deref(),
1187        &build_indirect_parameters_uniform_indices,
1188        &preprocess_pipelines.late_phase,
1189        "late_prepass_indirect_parameters_building",
1190    );
1191}
1192
1193/// A render graph system, run for each view, that builds the indirect
1194/// parameters for multi-draw indirect calls for the main opaque and transparent
1195/// passes.
1196pub fn main_build_indirect_parameters(
1197    current_view: ViewQuery<&ExtractedView, Without<ShadowView>>,
1198    preprocess_pipelines: Res<PreprocessPipelines>,
1199    build_indirect_params_bind_groups: Option<Res<BuildIndirectParametersBindGroups>>,
1200    pipeline_cache: Res<PipelineCache>,
1201    indirect_parameters_buffers: Option<Res<IndirectParametersBuffers>>,
1202    build_indirect_parameters_uniform_indices: Res<BuildIndirectParametersMetadata>,
1203    mut ctx: RenderContext,
1204) {
1205    run_build_indirect_parameters(
1206        &mut ctx,
1207        current_view.into_inner().retained_view_entity,
1208        build_indirect_params_bind_groups.as_deref(),
1209        &pipeline_cache,
1210        indirect_parameters_buffers.as_deref(),
1211        &build_indirect_parameters_uniform_indices,
1212        &preprocess_pipelines.main_phase,
1213        "main_indirect_parameters_building",
1214    );
1215}
1216
1217/// Shared logic common to all render graph systems that build indirect
1218/// parameters for multi-draw indirect calls.
1219pub(crate) fn run_build_indirect_parameters(
1220    ctx: &mut RenderContext,
1221    retained_view_entity: RetainedViewEntity,
1222    build_indirect_params_bind_groups: Option<&BuildIndirectParametersBindGroups>,
1223    pipeline_cache: &PipelineCache,
1224    indirect_parameters_buffers: Option<&IndirectParametersBuffers>,
1225    build_indirect_parameters_uniform_indices: &BuildIndirectParametersMetadata,
1226    preprocess_phase_pipelines: &PreprocessPhasePipelines,
1227    label: &'static str,
1228) {
1229    let Some(build_indirect_params_bind_groups) = build_indirect_params_bind_groups else {
1230        return;
1231    };
1232    let Some(indirect_parameters_buffers) = indirect_parameters_buffers else {
1233        return;
1234    };
1235    let Some(view_build_indirect_parameters_uniform_indices) =
1236        build_indirect_parameters_uniform_indices.get(&retained_view_entity)
1237    else {
1238        return;
1239    };
1240
1241    let command_encoder = ctx.command_encoder();
1242
1243    let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
1244        label: Some(label),
1245        timestamp_writes: None,
1246    });
1247
1248    // Fetch the pipeline.
1249    let (
1250        Some(reset_indirect_batch_sets_pipeline_id),
1251        Some(build_indexed_indirect_params_pipeline_id),
1252        Some(build_non_indexed_indirect_params_pipeline_id),
1253    ) = (
1254        preprocess_phase_pipelines
1255            .reset_indirect_batch_sets
1256            .pipeline_id,
1257        preprocess_phase_pipelines
1258            .gpu_occlusion_culling_build_indexed_indirect_params
1259            .pipeline_id,
1260        preprocess_phase_pipelines
1261            .gpu_occlusion_culling_build_non_indexed_indirect_params
1262            .pipeline_id,
1263    )
1264    else {
1265        {
    use ::tracing::__macro_support::Callsite as _;
    static __CALLSITE: ::tracing::callsite::DefaultCallsite =
        {
            static META: ::tracing::Metadata<'static> =
                {
                    ::tracing_core::metadata::Metadata::new("event src/render/gpu_preprocess.rs:1265",
                        "bevy_pbr::render::gpu_preprocess", ::tracing::Level::WARN,
                        ::tracing_core::__macro_support::Option::Some("src/render/gpu_preprocess.rs"),
                        ::tracing_core::__macro_support::Option::Some(1265u32),
                        ::tracing_core::__macro_support::Option::Some("bevy_pbr::render::gpu_preprocess"),
                        ::tracing_core::field::FieldSet::new(&["message"],
                            ::tracing_core::callsite::Identifier(&__CALLSITE)),
                        ::tracing::metadata::Kind::EVENT)
                };
            ::tracing::callsite::DefaultCallsite::new(&META)
        };
    let enabled =
        ::tracing::Level::WARN <= ::tracing::level_filters::STATIC_MAX_LEVEL
                &&
                ::tracing::Level::WARN <=
                    ::tracing::level_filters::LevelFilter::current() &&
            {
                let interest = __CALLSITE.interest();
                !interest.is_never() &&
                    ::tracing::__macro_support::__is_enabled(__CALLSITE.metadata(),
                        interest)
            };
    if enabled {
        (|value_set: ::tracing::field::ValueSet|
                    {
                        let meta = __CALLSITE.metadata();
                        ::tracing::Event::dispatch(meta, &value_set);
                        ;
                    })({
                #[allow(unused_imports)]
                use ::tracing::field::{debug, display, Value};
                __CALLSITE.metadata().fields().value_set_all(&[(::tracing::__macro_support::Option::Some(&format_args!("The build indirect parameters pipelines weren\'t ready")
                                            as &dyn ::tracing::field::Value))])
            });
    } else { ; }
};warn!("The build indirect parameters pipelines weren't ready");
1266        return;
1267    };
1268
1269    let (
1270        Some(reset_indirect_batch_sets_pipeline),
1271        Some(build_indexed_indirect_params_pipeline),
1272        Some(build_non_indexed_indirect_params_pipeline),
1273    ) = (
1274        pipeline_cache.get_compute_pipeline(reset_indirect_batch_sets_pipeline_id),
1275        pipeline_cache.get_compute_pipeline(build_indexed_indirect_params_pipeline_id),
1276        pipeline_cache.get_compute_pipeline(build_non_indexed_indirect_params_pipeline_id),
1277    )
1278    else {
1279        // This will happen while the pipeline is being compiled and is fine.
1280        return;
1281    };
1282
1283    // Loop over each phase. As each has as separate set of buffers, we need to
1284    // build indirect parameters individually for each phase.
1285    for (phase_type_id, phase_build_indirect_params_bind_groups) in
1286        build_indirect_params_bind_groups.iter()
1287    {
1288        let Some(phase_indirect_parameters_buffers) =
1289            indirect_parameters_buffers.get(phase_type_id)
1290        else {
1291            continue;
1292        };
1293        let Some(build_indirect_parameters_uniform_index) =
1294            view_build_indirect_parameters_uniform_indices.get(phase_type_id)
1295        else {
1296            continue;
1297        };
1298
1299        // Build indexed indirect parameters.
1300        if let (
1301            Some(reset_indexed_indirect_batch_sets_bind_group),
1302            Some(build_indirect_indexed_params_bind_group),
1303        ) = (
1304            &phase_build_indirect_params_bind_groups.reset_indexed_indirect_batch_sets,
1305            &phase_build_indirect_params_bind_groups.build_indexed_indirect,
1306        ) {
1307            compute_pass.set_pipeline(reset_indirect_batch_sets_pipeline);
1308            compute_pass.set_bind_group(0, reset_indexed_indirect_batch_sets_bind_group, &[]);
1309            let workgroup_count = phase_indirect_parameters_buffers
1310                .batch_set_count(true)
1311                .div_ceil(WORKGROUP_SIZE);
1312            if workgroup_count > 0 {
1313                compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
1314            }
1315
1316            compute_pass.set_pipeline(build_indexed_indirect_params_pipeline);
1317
1318            for indexed_build_indirect_parameters_metadata in
1319                &build_indirect_parameters_uniform_index.indexed
1320            {
1321                compute_pass.set_bind_group(
1322                    0,
1323                    build_indirect_indexed_params_bind_group,
1324                    &[indexed_build_indirect_parameters_metadata.uniform_offset],
1325                );
1326                let workgroup_count = indexed_build_indirect_parameters_metadata
1327                    .batch_count
1328                    .div_ceil(WORKGROUP_SIZE as u32);
1329                if workgroup_count > 0 {
1330                    compute_pass.dispatch_workgroups(workgroup_count, 1, 1);
1331                }
1332            }
1333        }
1334
1335        // Build non-indexed indirect parameters.
1336        if let (
1337            Some(reset_non_indexed_indirect_batch_sets_bind_group),
1338            Some(build_indirect_non_indexed_params_bind_group),
1339        ) = (
1340            &phase_build_indirect_params_bind_groups.reset_non_indexed_indirect_batch_sets,
1341            &phase_build_indirect_params_bind_groups.build_non_indexed_indirect,
1342        ) {
1343            compute_pass.set_pipeline(reset_indirect_batch_sets_pipeline);
1344            compute_pass.set_bind_group(0, reset_non_indexed_indirect_batch_sets_bind_group, &[]);
1345            let workgroup_count = phase_indirect_parameters_buffers
1346                .batch_set_count(false)
1347                .div_ceil(WORKGROUP_SIZE);
1348            if workgroup_count > 0 {
1349                compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
1350            }
1351
1352            compute_pass.set_pipeline(build_non_indexed_indirect_params_pipeline);
1353
1354            for non_indexed_build_indirect_parameters_metadata in
1355                &build_indirect_parameters_uniform_index.non_indexed
1356            {
1357                compute_pass.set_bind_group(
1358                    0,
1359                    build_indirect_non_indexed_params_bind_group,
1360                    &[non_indexed_build_indirect_parameters_metadata.uniform_offset],
1361                );
1362                let workgroup_count = non_indexed_build_indirect_parameters_metadata
1363                    .batch_count
1364                    .div_ceil(WORKGROUP_SIZE as u32);
1365                if workgroup_count > 0 {
1366                    compute_pass.dispatch_workgroups(workgroup_count, 1, 1);
1367                }
1368            }
1369        }
1370    }
1371}
1372
1373impl PreprocessPipelines {
1374    /// Returns true if the preprocessing and indirect parameters pipelines have
1375    /// been loaded or false otherwise.
1376    pub(crate) fn pipelines_are_loaded(
1377        &self,
1378        pipeline_cache: &PipelineCache,
1379        preprocessing_support: &GpuPreprocessingSupport,
1380    ) -> bool {
1381        match preprocessing_support.max_supported_mode {
1382            GpuPreprocessingMode::None => false,
1383            GpuPreprocessingMode::PreprocessingOnly => {
1384                self.direct_preprocess.is_loaded(pipeline_cache)
1385                    && self
1386                        .gpu_frustum_culling_preprocess
1387                        .is_loaded(pipeline_cache)
1388            }
1389            GpuPreprocessingMode::Culling => {
1390                self.direct_preprocess.is_loaded(pipeline_cache)
1391                    && self
1392                        .gpu_frustum_culling_preprocess
1393                        .is_loaded(pipeline_cache)
1394                    && self
1395                        .early_gpu_occlusion_culling_preprocess
1396                        .is_loaded(pipeline_cache)
1397                    && self
1398                        .late_gpu_occlusion_culling_preprocess
1399                        .is_loaded(pipeline_cache)
1400                    && self
1401                        .gpu_frustum_culling_build_indexed_indirect_params
1402                        .is_loaded(pipeline_cache)
1403                    && self
1404                        .gpu_frustum_culling_build_non_indexed_indirect_params
1405                        .is_loaded(pipeline_cache)
1406                    && self.early_phase.is_loaded(pipeline_cache)
1407                    && self.late_phase.is_loaded(pipeline_cache)
1408                    && self.main_phase.is_loaded(pipeline_cache)
1409            }
1410        }
1411    }
1412}
1413
1414impl PreprocessPhasePipelines {
1415    fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1416        self.reset_indirect_batch_sets.is_loaded(pipeline_cache)
1417            && self
1418                .gpu_occlusion_culling_build_indexed_indirect_params
1419                .is_loaded(pipeline_cache)
1420            && self
1421                .gpu_occlusion_culling_build_non_indexed_indirect_params
1422                .is_loaded(pipeline_cache)
1423    }
1424}
1425
1426impl PreprocessPipeline {
1427    fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1428        self.pipeline_id
1429            .is_some_and(|pipeline_id| pipeline_cache.get_compute_pipeline(pipeline_id).is_some())
1430    }
1431}
1432
1433impl ResetIndirectBatchSetsPipeline {
1434    fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1435        self.pipeline_id
1436            .is_some_and(|pipeline_id| pipeline_cache.get_compute_pipeline(pipeline_id).is_some())
1437    }
1438}
1439
1440impl BuildIndirectParametersPipeline {
1441    /// Returns true if this pipeline has been loaded into the pipeline cache or
1442    /// false otherwise.
1443    fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1444        self.pipeline_id
1445            .is_some_and(|pipeline_id| pipeline_cache.get_compute_pipeline(pipeline_id).is_some())
1446    }
1447}
1448
1449impl SpecializedComputePipeline for PreprocessPipeline {
1450    type Key = PreprocessPipelineKey;
1451
1452    fn specialize(&self, key: Self::Key) -> ComputePipelineDescriptor {
1453        let mut shader_defs = ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        ["WRITE_INDIRECT_PARAMETERS_METADATA".into()]))vec!["WRITE_INDIRECT_PARAMETERS_METADATA".into()];
1454        if key.contains(PreprocessPipelineKey::FRUSTUM_CULLING) {
1455            shader_defs.push("INDIRECT".into());
1456            shader_defs.push("FRUSTUM_CULLING".into());
1457        }
1458        if key.contains(PreprocessPipelineKey::OCCLUSION_CULLING) {
1459            shader_defs.push("OCCLUSION_CULLING".into());
1460            if key.contains(PreprocessPipelineKey::EARLY_PHASE) {
1461                shader_defs.push("EARLY_PHASE".into());
1462            } else {
1463                shader_defs.push("LATE_PHASE".into());
1464            }
1465        }
1466
1467        ComputePipelineDescriptor {
1468            label: Some(
1469                ::alloc::__export::must_use({
        ::alloc::fmt::format(format_args!("mesh preprocessing ({0})",
                if key.contains(PreprocessPipelineKey::OCCLUSION_CULLING |
                            PreprocessPipelineKey::EARLY_PHASE) {
                    "early GPU occlusion culling"
                } else if key.contains(PreprocessPipelineKey::OCCLUSION_CULLING)
                    {
                    "late GPU occlusion culling"
                } else if key.contains(PreprocessPipelineKey::FRUSTUM_CULLING)
                    {
                    "GPU frustum culling"
                } else { "direct" }))
    })format!(
1470                    "mesh preprocessing ({})",
1471                    if key.contains(
1472                        PreprocessPipelineKey::OCCLUSION_CULLING
1473                            | PreprocessPipelineKey::EARLY_PHASE
1474                    ) {
1475                        "early GPU occlusion culling"
1476                    } else if key.contains(PreprocessPipelineKey::OCCLUSION_CULLING) {
1477                        "late GPU occlusion culling"
1478                    } else if key.contains(PreprocessPipelineKey::FRUSTUM_CULLING) {
1479                        "GPU frustum culling"
1480                    } else {
1481                        "direct"
1482                    }
1483                )
1484                .into(),
1485            ),
1486            layout: ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        [self.bind_group_layout.clone()]))vec![self.bind_group_layout.clone()],
1487            immediate_size: if key.contains(PreprocessPipelineKey::OCCLUSION_CULLING) {
1488                4
1489            } else {
1490                0
1491            },
1492            shader: self.shader.clone(),
1493            shader_defs,
1494            ..default()
1495        }
1496    }
1497}
1498
1499impl FromWorld for PreprocessPipelines {
1500    fn from_world(world: &mut World) -> Self {
1501        // GPU culling bind group parameters are a superset of those in the CPU
1502        // culling (direct) shader.
1503        let direct_bind_group_layout_entries = preprocess_direct_bind_group_layout_entries();
1504        let gpu_frustum_culling_bind_group_layout_entries = gpu_culling_bind_group_layout_entries();
1505        let gpu_early_occlusion_culling_bind_group_layout_entries =
1506            gpu_occlusion_culling_bind_group_layout_entries().extend_with_indices((
1507                (
1508                    12,
1509                    storage_buffer::<PreprocessWorkItem>(/*has_dynamic_offset=*/ false),
1510                ),
1511                (
1512                    13,
1513                    storage_buffer::<LatePreprocessWorkItemIndirectParameters>(
1514                        /*has_dynamic_offset=*/ false,
1515                    ),
1516                ),
1517            ));
1518        let gpu_late_occlusion_culling_bind_group_layout_entries =
1519            gpu_occlusion_culling_bind_group_layout_entries().extend_with_indices(((
1520                13,
1521                storage_buffer_read_only::<LatePreprocessWorkItemIndirectParameters>(
1522                    /*has_dynamic_offset=*/ false,
1523                ),
1524            ),));
1525
1526        let reset_indirect_batch_sets_bind_group_layout_entries =
1527            DynamicBindGroupLayoutEntries::sequential(
1528                ShaderStages::COMPUTE,
1529                (storage_buffer::<IndirectBatchSet>(false),),
1530            );
1531
1532        // Indexed and non-indexed bind group parameters share all the bind
1533        // group layout entries except the final one.
1534        let build_indexed_indirect_params_bind_group_layout_entries =
1535            build_indirect_params_bind_group_layout_entries()
1536                .extend_sequential((storage_buffer::<IndirectParametersIndexed>(false),));
1537        let build_non_indexed_indirect_params_bind_group_layout_entries =
1538            build_indirect_params_bind_group_layout_entries()
1539                .extend_sequential((storage_buffer::<IndirectParametersNonIndexed>(false),));
1540
1541        let bin_unpacking_bind_group_layout_entries = bin_unpacking_bind_group_layout_entries();
1542        let uniform_allocation_bind_group_layout_entries =
1543            uniform_allocation_bind_group_layout_entries();
1544
1545        // Create the bind group layouts.
1546        let direct_bind_group_layout = BindGroupLayoutDescriptor::new(
1547            "build mesh uniforms direct bind group layout",
1548            &direct_bind_group_layout_entries,
1549        );
1550        let gpu_frustum_culling_bind_group_layout = BindGroupLayoutDescriptor::new(
1551            "build mesh uniforms GPU frustum culling bind group layout",
1552            &gpu_frustum_culling_bind_group_layout_entries,
1553        );
1554        let gpu_early_occlusion_culling_bind_group_layout = BindGroupLayoutDescriptor::new(
1555            "build mesh uniforms GPU early occlusion culling bind group layout",
1556            &gpu_early_occlusion_culling_bind_group_layout_entries,
1557        );
1558        let gpu_late_occlusion_culling_bind_group_layout = BindGroupLayoutDescriptor::new(
1559            "build mesh uniforms GPU late occlusion culling bind group layout",
1560            &gpu_late_occlusion_culling_bind_group_layout_entries,
1561        );
1562        let reset_indirect_batch_sets_bind_group_layout = BindGroupLayoutDescriptor::new(
1563            "reset indirect batch sets bind group layout",
1564            &reset_indirect_batch_sets_bind_group_layout_entries,
1565        );
1566        let build_indexed_indirect_params_bind_group_layout = BindGroupLayoutDescriptor::new(
1567            "build indexed indirect parameters bind group layout",
1568            &build_indexed_indirect_params_bind_group_layout_entries,
1569        );
1570        let build_non_indexed_indirect_params_bind_group_layout = BindGroupLayoutDescriptor::new(
1571            "build non-indexed indirect parameters bind group layout",
1572            &build_non_indexed_indirect_params_bind_group_layout_entries,
1573        );
1574        let bin_unpacking_bind_group_layout = BindGroupLayoutDescriptor::new(
1575            "bin unpacking bind group layout",
1576            &bin_unpacking_bind_group_layout_entries,
1577        );
1578        let uniform_allocation_bind_group_layout = BindGroupLayoutDescriptor::new(
1579            "uniform allocation bind group layout",
1580            &uniform_allocation_bind_group_layout_entries,
1581        );
1582
1583        let preprocess_shader = {
    let (path, asset_server) =
        {
            let path =
                {
                    {
                        let crate_name =
                            "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                        ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                            "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                            "mesh_preprocess.wesl".as_ref())
                    }
                };
            let path =
                ::bevy_asset::AssetPath::from_path_buf(path).with_source("embedded");
            let asset_server =
                ::bevy_asset::io::embedded::GetAssetServer::get_asset_server(world);
            (path, asset_server)
        };
    asset_server.load(path)
}load_embedded_asset!(world, "mesh_preprocess.wesl");
1584        let reset_indirect_batch_sets_shader =
1585            {
    let (path, asset_server) =
        {
            let path =
                {
                    {
                        let crate_name =
                            "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                        ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                            "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                            "reset_indirect_batch_sets.wesl".as_ref())
                    }
                };
            let path =
                ::bevy_asset::AssetPath::from_path_buf(path).with_source("embedded");
            let asset_server =
                ::bevy_asset::io::embedded::GetAssetServer::get_asset_server(world);
            (path, asset_server)
        };
    asset_server.load(path)
}load_embedded_asset!(world, "reset_indirect_batch_sets.wesl");
1586        let build_indirect_params_shader =
1587            {
    let (path, asset_server) =
        {
            let path =
                {
                    {
                        let crate_name =
                            "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                        ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                            "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                            "build_indirect_params.wesl".as_ref())
                    }
                };
            let path =
                ::bevy_asset::AssetPath::from_path_buf(path).with_source("embedded");
            let asset_server =
                ::bevy_asset::io::embedded::GetAssetServer::get_asset_server(world);
            (path, asset_server)
        };
    asset_server.load(path)
}load_embedded_asset!(world, "build_indirect_params.wesl");
1588        let bin_unpacking_shader = {
    let (path, asset_server) =
        {
            let path =
                {
                    {
                        let crate_name =
                            "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                        ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                            "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                            "unpack_bins.wesl".as_ref())
                    }
                };
            let path =
                ::bevy_asset::AssetPath::from_path_buf(path).with_source("embedded");
            let asset_server =
                ::bevy_asset::io::embedded::GetAssetServer::get_asset_server(world);
            (path, asset_server)
        };
    asset_server.load(path)
}load_embedded_asset!(world, "unpack_bins.wesl");
1589        let uniform_allocation_shader = {
    let (path, asset_server) =
        {
            let path =
                {
                    {
                        let crate_name =
                            "bevy_pbr::render::gpu_preprocess".split(':').next().unwrap();
                        ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                            "src".as_ref(), "src/render/gpu_preprocess.rs".as_ref(),
                            "allocate_uniforms.wesl".as_ref())
                    }
                };
            let path =
                ::bevy_asset::AssetPath::from_path_buf(path).with_source("embedded");
            let asset_server =
                ::bevy_asset::io::embedded::GetAssetServer::get_asset_server(world);
            (path, asset_server)
        };
    asset_server.load(path)
}load_embedded_asset!(world, "allocate_uniforms.wesl");
1590
1591        let preprocess_phase_pipelines = PreprocessPhasePipelines {
1592            reset_indirect_batch_sets: ResetIndirectBatchSetsPipeline {
1593                bind_group_layout: reset_indirect_batch_sets_bind_group_layout.clone(),
1594                shader: reset_indirect_batch_sets_shader,
1595                pipeline_id: None,
1596            },
1597            gpu_occlusion_culling_build_indexed_indirect_params: BuildIndirectParametersPipeline {
1598                bind_group_layout: build_indexed_indirect_params_bind_group_layout.clone(),
1599                shader: build_indirect_params_shader.clone(),
1600                pipeline_id: None,
1601            },
1602            gpu_occlusion_culling_build_non_indexed_indirect_params:
1603                BuildIndirectParametersPipeline {
1604                    bind_group_layout: build_non_indexed_indirect_params_bind_group_layout.clone(),
1605                    shader: build_indirect_params_shader.clone(),
1606                    pipeline_id: None,
1607                },
1608        };
1609
1610        PreprocessPipelines {
1611            direct_preprocess: PreprocessPipeline {
1612                bind_group_layout: direct_bind_group_layout,
1613                shader: preprocess_shader.clone(),
1614                pipeline_id: None,
1615            },
1616            gpu_frustum_culling_preprocess: PreprocessPipeline {
1617                bind_group_layout: gpu_frustum_culling_bind_group_layout,
1618                shader: preprocess_shader.clone(),
1619                pipeline_id: None,
1620            },
1621            early_gpu_occlusion_culling_preprocess: PreprocessPipeline {
1622                bind_group_layout: gpu_early_occlusion_culling_bind_group_layout,
1623                shader: preprocess_shader.clone(),
1624                pipeline_id: None,
1625            },
1626            late_gpu_occlusion_culling_preprocess: PreprocessPipeline {
1627                bind_group_layout: gpu_late_occlusion_culling_bind_group_layout,
1628                shader: preprocess_shader,
1629                pipeline_id: None,
1630            },
1631            gpu_frustum_culling_build_indexed_indirect_params: BuildIndirectParametersPipeline {
1632                bind_group_layout: build_indexed_indirect_params_bind_group_layout.clone(),
1633                shader: build_indirect_params_shader.clone(),
1634                pipeline_id: None,
1635            },
1636            gpu_frustum_culling_build_non_indexed_indirect_params:
1637                BuildIndirectParametersPipeline {
1638                    bind_group_layout: build_non_indexed_indirect_params_bind_group_layout.clone(),
1639                    shader: build_indirect_params_shader,
1640                    pipeline_id: None,
1641                },
1642            early_phase: preprocess_phase_pipelines.clone(),
1643            late_phase: preprocess_phase_pipelines.clone(),
1644            main_phase: preprocess_phase_pipelines.clone(),
1645            bin_unpacking: BinUnpackingPipeline {
1646                bind_group_layout: bin_unpacking_bind_group_layout,
1647                shader: bin_unpacking_shader,
1648                pipeline_id: None,
1649            },
1650            uniform_allocation: UniformAllocationPipelines {
1651                local_scan: UniformAllocationLocalScanPipeline {
1652                    bind_group_layout: uniform_allocation_bind_group_layout.clone(),
1653                    shader: uniform_allocation_shader.clone(),
1654                    pipeline_id_local_scan: None,
1655                },
1656                global_scan: UniformAllocationGlobalScanPipeline {
1657                    bind_group_layout: uniform_allocation_bind_group_layout.clone(),
1658                    shader: uniform_allocation_shader.clone(),
1659                    pipeline_id_global_scan: None,
1660                },
1661                fan: UniformAllocationFanPipeline {
1662                    bind_group_layout: uniform_allocation_bind_group_layout.clone(),
1663                    shader: uniform_allocation_shader.clone(),
1664                    pipeline_id_fan: None,
1665                },
1666            },
1667        }
1668    }
1669}
1670
1671fn preprocess_direct_bind_group_layout_entries() -> DynamicBindGroupLayoutEntries {
1672    DynamicBindGroupLayoutEntries::new_with_indices(
1673        ShaderStages::COMPUTE,
1674        (
1675            // `view`
1676            (
1677                0,
1678                uniform_buffer::<ViewUniform>(/* has_dynamic_offset= */ true),
1679            ),
1680            // `current_input`
1681            (3, storage_buffer_read_only::<MeshInputUniform>(false)),
1682            // `previous_input`
1683            (
1684                4,
1685                storage_buffer_read_only::<PreviousMeshInputUniform>(false),
1686            ),
1687            // `indices`
1688            (5, storage_buffer_read_only::<PreprocessWorkItem>(false)),
1689            // `output`
1690            (6, storage_buffer::<MeshUniform>(false)),
1691        ),
1692    )
1693}
1694
1695// Returns the first 5 bind group layout entries shared between all invocations
1696// of the indirect parameters building shader.
1697fn build_indirect_params_bind_group_layout_entries() -> DynamicBindGroupLayoutEntries {
1698    DynamicBindGroupLayoutEntries::new_with_indices(
1699        ShaderStages::COMPUTE,
1700        (
1701            // @group(0) @binding(0) var<storage> current_input:
1702            // array<MeshInput>;
1703            (0, storage_buffer_read_only::<MeshInputUniform>(false)),
1704            // @group(0) @binding(1) var<storage> indirect_parameters_metadata:
1705            // array<IndirectParametersMetadata>;
1706            (
1707                1,
1708                storage_buffer_read_only::<IndirectParametersMetadata>(false),
1709            ),
1710            // @group(0) @binding(3) var<storage, read_write>
1711            // indirect_batch_sets: array<IndirectBatchSet>;
1712            (3, storage_buffer::<IndirectBatchSet>(false)),
1713            // @group(0) @binding(4) var<uniform> indirect_parameters_build_job:
1714            // IndirectParametersBuildJob;
1715            (4, uniform_buffer::<IndirectParametersBuildJob>(true)),
1716        ),
1717    )
1718}
1719
1720/// A system that specializes the `mesh_preprocess.wesl` and
1721/// `build_indirect_params.wesl` pipelines if necessary.
1722fn gpu_culling_bind_group_layout_entries() -> DynamicBindGroupLayoutEntries {
1723    // GPU culling bind group parameters are a superset of those in the CPU
1724    // culling (direct) shader.
1725    preprocess_direct_bind_group_layout_entries().extend_with_indices((
1726        // @group(0) @binding(7) var<storage> indirect_parameters_metadata:
1727        // array<IndirectParametersMetadata>;
1728        (
1729            7,
1730            storage_buffer::<IndirectParametersMetadata>(/* has_dynamic_offset= */ false),
1731        ),
1732        // `mesh_culling_data`
1733        (
1734            9,
1735            storage_buffer_read_only::<MeshCullingData>(/* has_dynamic_offset= */ false),
1736        ),
1737        // `visibility_ranges`
1738        (
1739            10,
1740            storage_buffer_read_only::<Vec4>(/* has_dynamic_offset= */ false),
1741        ),
1742    ))
1743}
1744
1745fn gpu_occlusion_culling_bind_group_layout_entries() -> DynamicBindGroupLayoutEntries {
1746    gpu_culling_bind_group_layout_entries().extend_with_indices((
1747        (
1748            2,
1749            uniform_buffer::<PreviousViewData>(/*has_dynamic_offset=*/ false),
1750        ),
1751        (
1752            11,
1753            texture_2d(TextureSampleType::Float { filterable: true }),
1754        ),
1755    ))
1756}
1757
1758/// Creates and returns bind group layout entries for the GPU bin unpacking
1759/// shader (`unpack_bins`).
1760fn bin_unpacking_bind_group_layout_entries() -> BindGroupLayoutEntries<5> {
1761    BindGroupLayoutEntries::sequential(
1762        ShaderStages::COMPUTE,
1763        (
1764            // @group(0) @binding(0) var<uniform> bin_unpacking_metadata:
1765            // BinUnpackingMetadata;
1766            uniform_buffer::<GpuBinUnpackingMetadata>(false),
1767            // @group(0) @binding(1) var<storage> binned_mesh_instances:
1768            // array<BinnedMeshInstance>;
1769            storage_buffer_read_only::<GpuRenderBinnedMeshInstance>(false),
1770            // @group(0) @binding(2) var<storage, read_write>
1771            // preprocess_work_items: array<PreprocessWorkItem>;
1772            storage_buffer::<PreprocessWorkItem>(false),
1773            // @group(0) @binding(3) var<storage> bin_metadata:
1774            // array<GpuBinMetadata>;
1775            storage_buffer_read_only::<GpuBinMetadata>(false),
1776            // @group(0) @binding(4) var<storage>
1777            // bin_index_to_bin_metadata_index: array<u32>;
1778            storage_buffer_read_only::<u32>(false),
1779        ),
1780    )
1781}
1782
1783/// Creates and returns bind group layout entries for the GPU uniform allocation
1784/// shader (`allocate_uniforms`).
1785fn uniform_allocation_bind_group_layout_entries() -> BindGroupLayoutEntries<4> {
1786    BindGroupLayoutEntries::sequential(
1787        ShaderStages::COMPUTE,
1788        (
1789            // @group(0) @binding(0) var<uniform> allocate_uniforms_metadata:
1790            // AllocateUniformsMetadata;
1791            uniform_buffer::<GpuUniformAllocationMetadata>(false),
1792            // @group(0) @binding(1) var<storage> bin_metadata: array<BinMetadata>;
1793            storage_buffer_read_only::<GpuBinMetadata>(false),
1794            // @group(0) @binding(2) var<storage, read_write>
1795            // indirect_parameters_metadata: array<IndirectParametersMetadata>;
1796            storage_buffer::<IndirectParametersMetadata>(false),
1797            // @group(0) @binding(3) var<storage, read_write> fan_buffer:
1798            // array<u32>;
1799            storage_buffer::<u32>(false),
1800        ),
1801    )
1802}
1803
1804/// A system that specializes the pipelines relating to mesh preprocessing if
1805/// necessary.
1806///
1807/// These pipelines include those corresponding to the mesh preprocessing shader
1808/// itself, in addition to those corresponding to the indirect batch set
1809/// resetting shader, the indirect parameters building shader, and the bin
1810/// unpacking shader.
1811pub fn prepare_preprocess_pipelines(
1812    pipeline_cache: Res<PipelineCache>,
1813    render_device: Res<RenderDevice>,
1814    mut specialized_preprocess_pipelines: ResMut<SpecializedComputePipelines<PreprocessPipeline>>,
1815    mut specialized_reset_indirect_batch_sets_pipelines: ResMut<
1816        SpecializedComputePipelines<ResetIndirectBatchSetsPipeline>,
1817    >,
1818    mut specialized_build_indirect_parameters_pipelines: ResMut<
1819        SpecializedComputePipelines<BuildIndirectParametersPipeline>,
1820    >,
1821    mut specialized_bin_unpacking_pipelines: ResMut<
1822        SpecializedComputePipelines<BinUnpackingPipeline>,
1823    >,
1824    mut specialized_uniform_allocation_local_scan_pipelines: ResMut<
1825        SpecializedComputePipelines<UniformAllocationLocalScanPipeline>,
1826    >,
1827    mut specialized_uniform_allocation_global_scan_pipelines: ResMut<
1828        SpecializedComputePipelines<UniformAllocationGlobalScanPipeline>,
1829    >,
1830    mut specialized_uniform_allocation_fan_pipelines: ResMut<
1831        SpecializedComputePipelines<UniformAllocationFanPipeline>,
1832    >,
1833    preprocess_pipelines: ResMut<PreprocessPipelines>,
1834    gpu_preprocessing_support: Res<GpuPreprocessingSupport>,
1835) {
1836    let preprocess_pipelines = preprocess_pipelines.into_inner();
1837
1838    preprocess_pipelines.direct_preprocess.prepare(
1839        &pipeline_cache,
1840        &mut specialized_preprocess_pipelines,
1841        PreprocessPipelineKey::empty(),
1842    );
1843    preprocess_pipelines.gpu_frustum_culling_preprocess.prepare(
1844        &pipeline_cache,
1845        &mut specialized_preprocess_pipelines,
1846        PreprocessPipelineKey::FRUSTUM_CULLING,
1847    );
1848
1849    if gpu_preprocessing_support.is_culling_supported() {
1850        preprocess_pipelines
1851            .early_gpu_occlusion_culling_preprocess
1852            .prepare(
1853                &pipeline_cache,
1854                &mut specialized_preprocess_pipelines,
1855                PreprocessPipelineKey::FRUSTUM_CULLING
1856                    | PreprocessPipelineKey::OCCLUSION_CULLING
1857                    | PreprocessPipelineKey::EARLY_PHASE,
1858            );
1859        preprocess_pipelines
1860            .late_gpu_occlusion_culling_preprocess
1861            .prepare(
1862                &pipeline_cache,
1863                &mut specialized_preprocess_pipelines,
1864                PreprocessPipelineKey::FRUSTUM_CULLING | PreprocessPipelineKey::OCCLUSION_CULLING,
1865            );
1866    }
1867
1868    let mut build_indirect_parameters_pipeline_key = BuildIndirectParametersPipelineKey::empty();
1869
1870    // If the GPU and driver support `multi_draw_indirect_count`, tell the
1871    // shader that.
1872    if render_device
1873        .wgpu_device()
1874        .features()
1875        .contains(WgpuFeatures::MULTI_DRAW_INDIRECT_COUNT)
1876    {
1877        build_indirect_parameters_pipeline_key
1878            .insert(BuildIndirectParametersPipelineKey::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED);
1879    }
1880
1881    preprocess_pipelines
1882        .gpu_frustum_culling_build_indexed_indirect_params
1883        .prepare(
1884            &pipeline_cache,
1885            &mut specialized_build_indirect_parameters_pipelines,
1886            build_indirect_parameters_pipeline_key | BuildIndirectParametersPipelineKey::INDEXED,
1887        );
1888    preprocess_pipelines
1889        .gpu_frustum_culling_build_non_indexed_indirect_params
1890        .prepare(
1891            &pipeline_cache,
1892            &mut specialized_build_indirect_parameters_pipelines,
1893            build_indirect_parameters_pipeline_key,
1894        );
1895
1896    if !gpu_preprocessing_support.is_culling_supported() {
1897        return;
1898    }
1899
1900    for (preprocess_phase_pipelines, build_indirect_parameters_phase_pipeline_key) in [
1901        (
1902            &mut preprocess_pipelines.early_phase,
1903            BuildIndirectParametersPipelineKey::EARLY_PHASE,
1904        ),
1905        (
1906            &mut preprocess_pipelines.late_phase,
1907            BuildIndirectParametersPipelineKey::LATE_PHASE,
1908        ),
1909        (
1910            &mut preprocess_pipelines.main_phase,
1911            BuildIndirectParametersPipelineKey::MAIN_PHASE,
1912        ),
1913    ] {
1914        preprocess_phase_pipelines
1915            .reset_indirect_batch_sets
1916            .prepare(
1917                &pipeline_cache,
1918                &mut specialized_reset_indirect_batch_sets_pipelines,
1919            );
1920        preprocess_phase_pipelines
1921            .gpu_occlusion_culling_build_indexed_indirect_params
1922            .prepare(
1923                &pipeline_cache,
1924                &mut specialized_build_indirect_parameters_pipelines,
1925                build_indirect_parameters_pipeline_key
1926                    | build_indirect_parameters_phase_pipeline_key
1927                    | BuildIndirectParametersPipelineKey::INDEXED
1928                    | BuildIndirectParametersPipelineKey::OCCLUSION_CULLING,
1929            );
1930        preprocess_phase_pipelines
1931            .gpu_occlusion_culling_build_non_indexed_indirect_params
1932            .prepare(
1933                &pipeline_cache,
1934                &mut specialized_build_indirect_parameters_pipelines,
1935                build_indirect_parameters_pipeline_key
1936                    | build_indirect_parameters_phase_pipeline_key
1937                    | BuildIndirectParametersPipelineKey::OCCLUSION_CULLING,
1938            );
1939    }
1940
1941    // Prepare the bin unpacking compute pipeline.
1942    preprocess_pipelines
1943        .bin_unpacking
1944        .prepare(&pipeline_cache, &mut specialized_bin_unpacking_pipelines);
1945
1946    // Prepare the uniform allocation compute pipeline.
1947    preprocess_pipelines.uniform_allocation.prepare(
1948        &pipeline_cache,
1949        &mut specialized_uniform_allocation_local_scan_pipelines,
1950        &mut specialized_uniform_allocation_global_scan_pipelines,
1951        &mut specialized_uniform_allocation_fan_pipelines,
1952    );
1953}
1954
1955impl PreprocessPipeline {
1956    fn prepare(
1957        &mut self,
1958        pipeline_cache: &PipelineCache,
1959        pipelines: &mut SpecializedComputePipelines<PreprocessPipeline>,
1960        key: PreprocessPipelineKey,
1961    ) {
1962        if self.pipeline_id.is_some() {
1963            return;
1964        }
1965
1966        let preprocess_pipeline_id = pipelines.specialize(pipeline_cache, self, key);
1967        self.pipeline_id = Some(preprocess_pipeline_id);
1968    }
1969}
1970
1971impl SpecializedComputePipeline for ResetIndirectBatchSetsPipeline {
1972    type Key = ();
1973
1974    fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
1975        ComputePipelineDescriptor {
1976            label: Some("reset indirect batch sets".into()),
1977            layout: ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        [self.bind_group_layout.clone()]))vec![self.bind_group_layout.clone()],
1978            shader: self.shader.clone(),
1979            ..default()
1980        }
1981    }
1982}
1983
1984impl SpecializedComputePipeline for BuildIndirectParametersPipeline {
1985    type Key = BuildIndirectParametersPipelineKey;
1986
1987    fn specialize(&self, key: Self::Key) -> ComputePipelineDescriptor {
1988        let mut shader_defs = ::alloc::vec::Vec::new()vec![];
1989        if key.contains(BuildIndirectParametersPipelineKey::INDEXED) {
1990            shader_defs.push("INDEXED".into());
1991        }
1992        if key.contains(BuildIndirectParametersPipelineKey::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED) {
1993            shader_defs.push("MULTI_DRAW_INDIRECT_COUNT_SUPPORTED".into());
1994        }
1995        if key.contains(BuildIndirectParametersPipelineKey::OCCLUSION_CULLING) {
1996            shader_defs.push("OCCLUSION_CULLING".into());
1997        }
1998        if key.contains(BuildIndirectParametersPipelineKey::EARLY_PHASE) {
1999            shader_defs.push("EARLY_PHASE".into());
2000        }
2001        if key.contains(BuildIndirectParametersPipelineKey::LATE_PHASE) {
2002            shader_defs.push("LATE_PHASE".into());
2003        }
2004        if key.contains(BuildIndirectParametersPipelineKey::MAIN_PHASE) {
2005            shader_defs.push("MAIN_PHASE".into());
2006        }
2007
2008        let label = ::alloc::__export::must_use({
        ::alloc::fmt::format(format_args!("{0} build {1}indexed indirect parameters",
                if !key.contains(BuildIndirectParametersPipelineKey::OCCLUSION_CULLING)
                    {
                    "frustum culling"
                } else if key.contains(BuildIndirectParametersPipelineKey::EARLY_PHASE)
                    {
                    "early occlusion culling"
                } else if key.contains(BuildIndirectParametersPipelineKey::LATE_PHASE)
                    {
                    "late occlusion culling"
                } else { "main occlusion culling" },
                if key.contains(BuildIndirectParametersPipelineKey::INDEXED) {
                    ""
                } else { "non-" }))
    })format!(
2009            "{} build {}indexed indirect parameters",
2010            if !key.contains(BuildIndirectParametersPipelineKey::OCCLUSION_CULLING) {
2011                "frustum culling"
2012            } else if key.contains(BuildIndirectParametersPipelineKey::EARLY_PHASE) {
2013                "early occlusion culling"
2014            } else if key.contains(BuildIndirectParametersPipelineKey::LATE_PHASE) {
2015                "late occlusion culling"
2016            } else {
2017                "main occlusion culling"
2018            },
2019            if key.contains(BuildIndirectParametersPipelineKey::INDEXED) {
2020                ""
2021            } else {
2022                "non-"
2023            }
2024        );
2025
2026        ComputePipelineDescriptor {
2027            label: Some(label.into()),
2028            layout: ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        [self.bind_group_layout.clone()]))vec![self.bind_group_layout.clone()],
2029            shader: self.shader.clone(),
2030            shader_defs,
2031            ..default()
2032        }
2033    }
2034}
2035
2036impl SpecializedComputePipeline for BinUnpackingPipeline {
2037    type Key = ();
2038
2039    fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2040        ComputePipelineDescriptor {
2041            label: Some("bin unpacking".into()),
2042            layout: ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        [self.bind_group_layout.clone()]))vec![self.bind_group_layout.clone()],
2043            shader: self.shader.clone(),
2044            shader_defs: ::alloc::vec::Vec::new()vec![],
2045            ..default()
2046        }
2047    }
2048}
2049
2050impl SpecializedComputePipeline for UniformAllocationLocalScanPipeline {
2051    type Key = ();
2052
2053    fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2054        ComputePipelineDescriptor {
2055            label: Some("uniform allocation, local scan".into()),
2056            layout: ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        [self.bind_group_layout.clone()]))vec![self.bind_group_layout.clone()],
2057            shader: self.shader.clone(),
2058            shader_defs: ::alloc::vec::Vec::new()vec![],
2059            entry_point: Some("allocate_local_scan".into()),
2060            ..Default::default()
2061        }
2062    }
2063}
2064
2065impl SpecializedComputePipeline for UniformAllocationGlobalScanPipeline {
2066    type Key = ();
2067
2068    fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2069        ComputePipelineDescriptor {
2070            label: Some("uniform allocation, global scan".into()),
2071            layout: ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        [self.bind_group_layout.clone()]))vec![self.bind_group_layout.clone()],
2072            shader: self.shader.clone(),
2073            shader_defs: ::alloc::vec::Vec::new()vec![],
2074            entry_point: Some("allocate_global_scan".into()),
2075            ..Default::default()
2076        }
2077    }
2078}
2079
2080impl SpecializedComputePipeline for UniformAllocationFanPipeline {
2081    type Key = ();
2082
2083    fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2084        ComputePipelineDescriptor {
2085            label: Some("uniform allocation, fan".into()),
2086            layout: ::alloc::boxed::box_assume_init_into_vec_unsafe(::alloc::intrinsics::write_box_via_move(::alloc::boxed::Box::new_uninit(),
        [self.bind_group_layout.clone()]))vec![self.bind_group_layout.clone()],
2087            shader: self.shader.clone(),
2088            shader_defs: ::alloc::vec::Vec::new()vec![],
2089            entry_point: Some("allocate_fan".into()),
2090            ..Default::default()
2091        }
2092    }
2093}
2094
2095impl ResetIndirectBatchSetsPipeline {
2096    fn prepare(
2097        &mut self,
2098        pipeline_cache: &PipelineCache,
2099        pipelines: &mut SpecializedComputePipelines<ResetIndirectBatchSetsPipeline>,
2100    ) {
2101        if self.pipeline_id.is_some() {
2102            return;
2103        }
2104
2105        let reset_indirect_batch_sets_pipeline_id = pipelines.specialize(pipeline_cache, self, ());
2106        self.pipeline_id = Some(reset_indirect_batch_sets_pipeline_id);
2107    }
2108}
2109
2110impl BuildIndirectParametersPipeline {
2111    fn prepare(
2112        &mut self,
2113        pipeline_cache: &PipelineCache,
2114        pipelines: &mut SpecializedComputePipelines<BuildIndirectParametersPipeline>,
2115        key: BuildIndirectParametersPipelineKey,
2116    ) {
2117        if self.pipeline_id.is_some() {
2118            return;
2119        }
2120
2121        let build_indirect_parameters_pipeline_id = pipelines.specialize(pipeline_cache, self, key);
2122        self.pipeline_id = Some(build_indirect_parameters_pipeline_id);
2123    }
2124}
2125
2126impl BinUnpackingPipeline {
2127    /// Specializes a single pipeline for the bin unpacking shader.
2128    fn prepare(
2129        &mut self,
2130        pipeline_cache: &PipelineCache,
2131        pipelines: &mut SpecializedComputePipelines<BinUnpackingPipeline>,
2132    ) {
2133        if self.pipeline_id.is_some() {
2134            return;
2135        }
2136
2137        let bin_unpacking_pipeline_id = pipelines.specialize(pipeline_cache, self, ());
2138        self.pipeline_id = Some(bin_unpacking_pipeline_id);
2139    }
2140}
2141
2142impl UniformAllocationPipelines {
2143    /// Specializes all three pipelines that use the uniform allocation shader.
2144    fn prepare(
2145        &mut self,
2146        pipeline_cache: &PipelineCache,
2147        uniform_allocation_local_scan_pipelines: &mut SpecializedComputePipelines<
2148            UniformAllocationLocalScanPipeline,
2149        >,
2150        uniform_allocation_global_scan_pipelines: &mut SpecializedComputePipelines<
2151            UniformAllocationGlobalScanPipeline,
2152        >,
2153        uniform_allocation_fan_pipelines: &mut SpecializedComputePipelines<
2154            UniformAllocationFanPipeline,
2155        >,
2156    ) {
2157        if self.local_scan.pipeline_id_local_scan.is_none() {
2158            self.local_scan.pipeline_id_local_scan =
2159                Some(uniform_allocation_local_scan_pipelines.specialize(
2160                    pipeline_cache,
2161                    &self.local_scan,
2162                    (),
2163                ));
2164        }
2165
2166        if self.global_scan.pipeline_id_global_scan.is_none() {
2167            self.global_scan.pipeline_id_global_scan =
2168                Some(uniform_allocation_global_scan_pipelines.specialize(
2169                    pipeline_cache,
2170                    &self.global_scan,
2171                    (),
2172                ));
2173        }
2174
2175        if self.fan.pipeline_id_fan.is_none() {
2176            self.fan.pipeline_id_fan =
2177                Some(uniform_allocation_fan_pipelines.specialize(pipeline_cache, &self.fan, ()));
2178        }
2179    }
2180}
2181
2182/// A system that attaches buffers to bind groups for the variants of the
2183/// compute shaders relating to mesh preprocessing.
2184#[expect(
2185    clippy::too_many_arguments,
2186    reason = "it's a system that needs a lot of arguments"
2187)]
2188pub fn prepare_preprocess_bind_groups(
2189    mut commands: Commands,
2190    views: Query<(Entity, &ExtractedView)>,
2191    view_depth_pyramids: Query<(&ViewDepthPyramid, &PreviousViewUniformOffset)>,
2192    render_device: Res<RenderDevice>,
2193    pipeline_cache: Res<PipelineCache>,
2194    batched_instance_buffers: Res<BatchedInstanceBuffers<MeshUniform, MeshInputUniform>>,
2195    indirect_parameters_buffers: Res<IndirectParametersBuffers>,
2196    indirect_parameters_build_jobs: Res<IndirectParametersBuildJobs>,
2197    scene_unpacking_buffers: Res<SceneUnpackingBuffers>,
2198    mesh_culling_data_buffer: Res<MeshCullingDataBuffer>,
2199    visibility_ranges: Res<RenderVisibilityRanges>,
2200    view_uniforms: Res<ViewUniforms>,
2201    previous_view_uniforms: Res<PreviousViewUniforms>,
2202    pipelines: Res<PreprocessPipelines>,
2203    mut bin_unpacking_bind_groups: ResMut<BinUnpackingBindGroups>,
2204    mut uniform_allocation_bind_groups: ResMut<UniformAllocationBindGroups>,
2205) {
2206    // Grab the `BatchedInstanceBuffers`.
2207    let BatchedInstanceBuffers {
2208        current_input_buffer: current_input_buffer_vec,
2209        previous_input_buffer: previous_input_buffer_vec,
2210        phase_instance_buffers,
2211    } = batched_instance_buffers.into_inner();
2212
2213    let (Some(current_input_buffer), Some(previous_input_buffer)) = (
2214        current_input_buffer_vec.buffer().buffer(),
2215        previous_input_buffer_vec.buffer(),
2216    ) else {
2217        return;
2218    };
2219
2220    // Record whether we have any meshes that are to be drawn indirectly. If we
2221    // don't, then we can skip building indirect parameters.
2222    let mut any_indirect = false;
2223
2224    // Loop over each view.
2225    for (view_entity, view) in &views {
2226        let mut bind_groups = TypeIdHashMap::default();
2227
2228        // Loop over each phase.
2229        for (phase_type_id, phase_instance_buffers) in phase_instance_buffers {
2230            let UntypedPhaseBatchedInstanceBuffers {
2231                data_buffer: ref data_buffer_vec,
2232                ref work_item_buffers,
2233                ref late_indexed_indirect_parameters_buffer,
2234                ref late_non_indexed_indirect_parameters_buffer,
2235            } = *phase_instance_buffers;
2236
2237            let Some(data_buffer) = data_buffer_vec.buffer() else {
2238                continue;
2239            };
2240
2241            // Grab the indirect parameters buffers for this phase.
2242            let Some(phase_indirect_parameters_buffers) =
2243                indirect_parameters_buffers.get(phase_type_id)
2244            else {
2245                continue;
2246            };
2247
2248            let Some(work_item_buffers) = work_item_buffers.get(&view.retained_view_entity) else {
2249                continue;
2250            };
2251
2252            // Create the `PreprocessBindGroupBuilder`.
2253            let preprocess_bind_group_builder = PreprocessBindGroupBuilder {
2254                view: view_entity,
2255                late_indexed_indirect_parameters_buffer,
2256                late_non_indexed_indirect_parameters_buffer,
2257                render_device: &render_device,
2258                pipeline_cache: &pipeline_cache,
2259                phase_indirect_parameters_buffers,
2260                mesh_culling_data_buffer: &mesh_culling_data_buffer,
2261                visibility_range_data_buffer: visibility_ranges.buffer(),
2262                view_uniforms: &view_uniforms,
2263                previous_view_uniforms: &previous_view_uniforms,
2264                pipelines: &pipelines,
2265                current_input_buffer,
2266                previous_input_buffer,
2267                data_buffer,
2268            };
2269
2270            // Depending on the type of work items we have, construct the
2271            // appropriate bind groups.
2272            let (was_indirect, bind_group) = match *work_item_buffers {
2273                PreprocessWorkItemBuffers::Direct(ref work_item_buffer) => (
2274                    false,
2275                    preprocess_bind_group_builder
2276                        .create_direct_preprocess_bind_groups(work_item_buffer),
2277                ),
2278
2279                PreprocessWorkItemBuffers::Indirect {
2280                    indexed: ref indexed_work_item_buffer,
2281                    non_indexed: ref non_indexed_work_item_buffer,
2282                    gpu_occlusion_culling: Some(ref gpu_occlusion_culling_work_item_buffers),
2283                } => (
2284                    true,
2285                    preprocess_bind_group_builder
2286                        .create_indirect_occlusion_culling_preprocess_bind_groups(
2287                            &view_depth_pyramids,
2288                            indexed_work_item_buffer,
2289                            non_indexed_work_item_buffer,
2290                            gpu_occlusion_culling_work_item_buffers,
2291                        ),
2292                ),
2293
2294                PreprocessWorkItemBuffers::Indirect {
2295                    indexed: ref indexed_work_item_buffer,
2296                    non_indexed: ref non_indexed_work_item_buffer,
2297                    gpu_occlusion_culling: None,
2298                } => (
2299                    true,
2300                    preprocess_bind_group_builder
2301                        .create_indirect_frustum_culling_preprocess_bind_groups(
2302                            indexed_work_item_buffer,
2303                            non_indexed_work_item_buffer,
2304                        ),
2305                ),
2306            };
2307
2308            // Write that bind group in.
2309            if let Some(bind_group) = bind_group {
2310                any_indirect = any_indirect || was_indirect;
2311                bind_groups.insert(*phase_type_id, bind_group);
2312            }
2313        }
2314
2315        // Save the bind groups.
2316        commands
2317            .entity(view_entity)
2318            .insert(PreprocessBindGroups(bind_groups));
2319    }
2320
2321    // Now, if there were any indirect draw commands, create the bind groups for
2322    // the indirect parameters building shader.
2323    if any_indirect {
2324        create_build_indirect_parameters_bind_groups(
2325            &mut commands,
2326            &render_device,
2327            &pipeline_cache,
2328            &pipelines,
2329            current_input_buffer,
2330            &indirect_parameters_buffers,
2331            &indirect_parameters_build_jobs,
2332        );
2333    }
2334
2335    // Create the bind groups we'll need for each dispatch of the bin unpacking
2336    // (`unpack_bins`) and uniform allocation (`allocate_uniforms`) shaders.
2337    for (_, view) in &views {
2338        create_bin_unpacking_bind_groups(
2339            &mut bin_unpacking_bind_groups,
2340            &render_device,
2341            &pipeline_cache,
2342            &pipelines,
2343            &indirect_parameters_buffers,
2344            phase_instance_buffers,
2345            &scene_unpacking_buffers,
2346            &view.retained_view_entity,
2347        );
2348        create_uniform_allocation_bind_groups(
2349            &mut uniform_allocation_bind_groups,
2350            &render_device,
2351            &pipeline_cache,
2352            &pipelines,
2353            &indirect_parameters_buffers,
2354            &scene_unpacking_buffers,
2355            &view.retained_view_entity,
2356        );
2357    }
2358}
2359
2360/// A temporary structure that stores all the information needed to construct
2361/// bind groups for the mesh preprocessing shader.
2362struct PreprocessBindGroupBuilder<'a> {
2363    /// The render-world entity corresponding to the current view.
2364    view: Entity,
2365    /// The indirect compute dispatch parameters buffer for indexed meshes in
2366    /// the late prepass.
2367    late_indexed_indirect_parameters_buffer:
2368        &'a RawBufferVec<LatePreprocessWorkItemIndirectParameters>,
2369    /// The indirect compute dispatch parameters buffer for non-indexed meshes
2370    /// in the late prepass.
2371    late_non_indexed_indirect_parameters_buffer:
2372        &'a RawBufferVec<LatePreprocessWorkItemIndirectParameters>,
2373    /// The device.
2374    render_device: &'a RenderDevice,
2375    /// The pipeline cache
2376    pipeline_cache: &'a PipelineCache,
2377    /// The buffers that store indirect draw parameters.
2378    phase_indirect_parameters_buffers: &'a UntypedPhaseIndirectParametersBuffers,
2379    /// The GPU buffer that stores the information needed to cull each mesh.
2380    mesh_culling_data_buffer: &'a MeshCullingDataBuffer,
2381    /// The device buffer that stores the information needed to process
2382    /// visibility ranges on the GPU.
2383    visibility_range_data_buffer: &'a BufferVec<Vec4>,
2384    /// The GPU buffer that stores information about the view.
2385    view_uniforms: &'a ViewUniforms,
2386    /// The GPU buffer that stores information about the view from last frame.
2387    previous_view_uniforms: &'a PreviousViewUniforms,
2388    /// The pipelines for the mesh preprocessing shader.
2389    pipelines: &'a PreprocessPipelines,
2390    /// The GPU buffer containing the list of [`MeshInputUniform`]s for the
2391    /// current frame.
2392    current_input_buffer: &'a Buffer,
2393    /// The GPU buffer containing the list of [`MeshInputUniform`]s for the
2394    /// previous frame.
2395    previous_input_buffer: &'a Buffer,
2396    /// The GPU buffer containing the list of [`MeshUniform`]s for the current
2397    /// frame.
2398    ///
2399    /// This is the buffer containing the mesh's final transforms that the
2400    /// shaders will write to.
2401    data_buffer: &'a Buffer,
2402}
2403
2404impl<'a> PreprocessBindGroupBuilder<'a> {
2405    /// Creates the bind groups for mesh preprocessing when GPU frustum culling
2406    /// and GPU occlusion culling are both disabled.
2407    fn create_direct_preprocess_bind_groups(
2408        &self,
2409        work_item_buffer: &RawBufferVec<PreprocessWorkItem>,
2410    ) -> Option<PhasePreprocessBindGroups> {
2411        // Don't use `as_entire_binding()` here; the shader reads the array
2412        // length and the underlying buffer may be longer than the actual size
2413        // of the vector.
2414        let work_item_buffer_size = NonZero::<u64>::try_from(
2415            work_item_buffer.len() as u64 * u64::from(PreprocessWorkItem::min_size()),
2416        )
2417        .ok();
2418
2419        Some(PhasePreprocessBindGroups::Direct(
2420            self.render_device.create_bind_group(
2421                "preprocess_direct_bind_group",
2422                &self
2423                    .pipeline_cache
2424                    .get_bind_group_layout(&self.pipelines.direct_preprocess.bind_group_layout),
2425                &BindGroupEntries::with_indices((
2426                    (0, self.view_uniforms.uniforms.binding()?),
2427                    (3, self.current_input_buffer.as_entire_binding()),
2428                    (4, self.previous_input_buffer.as_entire_binding()),
2429                    (
2430                        5,
2431                        BindingResource::Buffer(BufferBinding {
2432                            buffer: work_item_buffer.buffer()?,
2433                            offset: 0,
2434                            size: work_item_buffer_size,
2435                        }),
2436                    ),
2437                    (6, self.data_buffer.as_entire_binding()),
2438                )),
2439            ),
2440        ))
2441    }
2442
2443    /// Creates the bind groups for mesh preprocessing when GPU occlusion
2444    /// culling is enabled.
2445    fn create_indirect_occlusion_culling_preprocess_bind_groups(
2446        &self,
2447        view_depth_pyramids: &Query<(&ViewDepthPyramid, &PreviousViewUniformOffset)>,
2448        indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2449        non_indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2450        gpu_occlusion_culling_work_item_buffers: &GpuOcclusionCullingWorkItemBuffers,
2451    ) -> Option<PhasePreprocessBindGroups> {
2452        let GpuOcclusionCullingWorkItemBuffers {
2453            late_indexed: ref late_indexed_work_item_buffer,
2454            late_non_indexed: ref late_non_indexed_work_item_buffer,
2455            ..
2456        } = *gpu_occlusion_culling_work_item_buffers;
2457
2458        let (view_depth_pyramid, previous_view_uniform_offset) =
2459            view_depth_pyramids.get(self.view).ok()?;
2460
2461        Some(PhasePreprocessBindGroups::IndirectOcclusionCulling {
2462            early_indexed: self.create_indirect_occlusion_culling_early_indexed_bind_group(
2463                view_depth_pyramid,
2464                previous_view_uniform_offset,
2465                indexed_work_item_buffer,
2466                late_indexed_work_item_buffer,
2467            ),
2468
2469            early_non_indexed: self.create_indirect_occlusion_culling_early_non_indexed_bind_group(
2470                view_depth_pyramid,
2471                previous_view_uniform_offset,
2472                non_indexed_work_item_buffer,
2473                late_non_indexed_work_item_buffer,
2474            ),
2475
2476            late_indexed: self.create_indirect_occlusion_culling_late_indexed_bind_group(
2477                view_depth_pyramid,
2478                previous_view_uniform_offset,
2479                late_indexed_work_item_buffer,
2480            ),
2481
2482            late_non_indexed: self.create_indirect_occlusion_culling_late_non_indexed_bind_group(
2483                view_depth_pyramid,
2484                previous_view_uniform_offset,
2485                late_non_indexed_work_item_buffer,
2486            ),
2487        })
2488    }
2489
2490    /// Creates the bind group for the first phase of mesh preprocessing of
2491    /// indexed meshes when GPU occlusion culling is enabled.
2492    fn create_indirect_occlusion_culling_early_indexed_bind_group(
2493        &self,
2494        view_depth_pyramid: &ViewDepthPyramid,
2495        previous_view_uniform_offset: &PreviousViewUniformOffset,
2496        indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2497        late_indexed_work_item_buffer: &UninitBufferVec<PreprocessWorkItem>,
2498    ) -> Option<BindGroup> {
2499        let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2500        let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2501        let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2502        let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
2503
2504        match (
2505            self.phase_indirect_parameters_buffers
2506                .indexed
2507                .metadata_buffer(),
2508            indexed_work_item_buffer.buffer(),
2509            late_indexed_work_item_buffer.buffer(),
2510            self.late_indexed_indirect_parameters_buffer.buffer(),
2511        ) {
2512            (
2513                Some(indexed_metadata_buffer),
2514                Some(indexed_work_item_gpu_buffer),
2515                Some(late_indexed_work_item_gpu_buffer),
2516                Some(late_indexed_indirect_parameters_buffer),
2517            ) => {
2518                // Don't use `as_entire_binding()` here; the shader reads the array
2519                // length and the underlying buffer may be longer than the actual size
2520                // of the vector.
2521                let indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2522                    indexed_work_item_buffer.len() as u64
2523                        * u64::from(PreprocessWorkItem::min_size()),
2524                )
2525                .ok();
2526
2527                Some(
2528                    self.render_device.create_bind_group(
2529                        "preprocess_early_indexed_gpu_occlusion_culling_bind_group",
2530                        &self.pipeline_cache.get_bind_group_layout(
2531                            &self
2532                                .pipelines
2533                                .early_gpu_occlusion_culling_preprocess
2534                                .bind_group_layout,
2535                        ),
2536                        &BindGroupEntries::with_indices((
2537                            // @group(0) @binding(3) var<storage> current_input:
2538                            // array<MeshInput>;
2539                            (3, self.current_input_buffer.as_entire_binding()),
2540                            // @group(0) @binding(4) var<storage>
2541                            // previous_input: array<MeshInput>;
2542                            (4, self.previous_input_buffer.as_entire_binding()),
2543                            // @group(0) @binding(5) var<storage> work_items:
2544                            // array<PreprocessWorkItem>;
2545                            (
2546                                5,
2547                                BindingResource::Buffer(BufferBinding {
2548                                    buffer: indexed_work_item_gpu_buffer,
2549                                    offset: 0,
2550                                    size: indexed_work_item_buffer_size,
2551                                }),
2552                            ),
2553                            // @group(0) @binding(6) var<storage, read_write>
2554                            // output: array<Mesh>;
2555                            (6, self.data_buffer.as_entire_binding()),
2556                            // @group(0) @binding(7) var<storage>
2557                            // indirect_parameters_metadata:
2558                            // array<IndirectParametersMetadata>;
2559                            (7, indexed_metadata_buffer.as_entire_binding()),
2560                            // @group(0) @binding(9) var<storage>
2561                            // mesh_culling_data: array<MeshCullingData>;
2562                            (9, mesh_culling_data_buffer.as_entire_binding()),
2563                            // @group(0) @binding(10) var<storage>
2564                            // visibility_ranges: array<vec4<f32>>;
2565                            (10, visibility_range_binding.clone()),
2566                            // @group(0) @binding(0) var<uniform> view: View;
2567                            (0, view_uniforms_binding.clone()),
2568                            // @group(0) @binding(11) var depth_pyramid:
2569                            // texture_2d<f32>;
2570                            (11, &view_depth_pyramid.all_mips),
2571                            // @group(0) @binding(2) var<uniform>
2572                            // previous_view_uniforms: PreviousViewUniforms;
2573                            (
2574                                2,
2575                                BufferBinding {
2576                                    buffer: previous_view_buffer,
2577                                    offset: previous_view_uniform_offset.offset as u64,
2578                                    size: NonZeroU64::new(size_of::<PreviousViewData>() as u64),
2579                                },
2580                            ),
2581                            // @group(0) @binding(12) var<storage, read_write>
2582                            // late_preprocess_work_items:
2583                            // array<PreprocessWorkItem>;
2584                            (
2585                                12,
2586                                BufferBinding {
2587                                    buffer: late_indexed_work_item_gpu_buffer,
2588                                    offset: 0,
2589                                    size: indexed_work_item_buffer_size,
2590                                },
2591                            ),
2592                            // @group(0) @binding(13) var<storage, read_write>
2593                            // late_preprocess_work_item_indirect_parameters:
2594                            // array<LatePreprocessWorkItemIndirectParameters>;
2595                            (
2596                                13,
2597                                BufferBinding {
2598                                    buffer: late_indexed_indirect_parameters_buffer,
2599                                    offset: 0,
2600                                    size: NonZeroU64::new(
2601                                        late_indexed_indirect_parameters_buffer.size(),
2602                                    ),
2603                                },
2604                            ),
2605                        )),
2606                    ),
2607                )
2608            }
2609            _ => None,
2610        }
2611    }
2612
2613    /// Creates the bind group for the first phase of mesh preprocessing of
2614    /// non-indexed meshes when GPU occlusion culling is enabled.
2615    fn create_indirect_occlusion_culling_early_non_indexed_bind_group(
2616        &self,
2617        view_depth_pyramid: &ViewDepthPyramid,
2618        previous_view_uniform_offset: &PreviousViewUniformOffset,
2619        non_indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2620        late_non_indexed_work_item_buffer: &UninitBufferVec<PreprocessWorkItem>,
2621    ) -> Option<BindGroup> {
2622        let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2623        let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2624        let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2625        let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
2626
2627        match (
2628            self.phase_indirect_parameters_buffers
2629                .non_indexed
2630                .metadata_buffer(),
2631            non_indexed_work_item_buffer.buffer(),
2632            late_non_indexed_work_item_buffer.buffer(),
2633            self.late_non_indexed_indirect_parameters_buffer.buffer(),
2634        ) {
2635            (
2636                Some(non_indexed_metadata_buffer),
2637                Some(non_indexed_work_item_gpu_buffer),
2638                Some(late_non_indexed_work_item_buffer),
2639                Some(late_non_indexed_indirect_parameters_buffer),
2640            ) => {
2641                // Don't use `as_entire_binding()` here; the shader reads the array
2642                // length and the underlying buffer may be longer than the actual size
2643                // of the vector.
2644                let non_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2645                    non_indexed_work_item_buffer.len() as u64
2646                        * u64::from(PreprocessWorkItem::min_size()),
2647                )
2648                .ok();
2649
2650                Some(
2651                    self.render_device.create_bind_group(
2652                        "preprocess_early_non_indexed_gpu_occlusion_culling_bind_group",
2653                        &self.pipeline_cache.get_bind_group_layout(
2654                            &self
2655                                .pipelines
2656                                .early_gpu_occlusion_culling_preprocess
2657                                .bind_group_layout,
2658                        ),
2659                        &BindGroupEntries::with_indices((
2660                            // @group(0) @binding(3) var<storage> current_input:
2661                            // array<MeshInput>;
2662                            (3, self.current_input_buffer.as_entire_binding()),
2663                            // @group(0) @binding(4) var<storage>
2664                            // previous_input: array<MeshInput>;
2665                            (4, self.previous_input_buffer.as_entire_binding()),
2666                            // @group(0) @binding(5) var<storage> work_items:
2667                            // array<PreprocessWorkItem>;
2668                            (
2669                                5,
2670                                BindingResource::Buffer(BufferBinding {
2671                                    buffer: non_indexed_work_item_gpu_buffer,
2672                                    offset: 0,
2673                                    size: non_indexed_work_item_buffer_size,
2674                                }),
2675                            ),
2676                            (6, self.data_buffer.as_entire_binding()),
2677                            // @group(0) @binding(7) var<storage>
2678                            // indirect_parameters_metadata:
2679                            // array<IndirectParametersMetadata>;
2680                            (7, non_indexed_metadata_buffer.as_entire_binding()),
2681                            // @group(0) @binding(9) var<storage>
2682                            // mesh_culling_data: array<MeshCullingData>;
2683                            (9, mesh_culling_data_buffer.as_entire_binding()),
2684                            // @group(0) @binding(10) var<storage>
2685                            // visibility_ranges: array<vec4<f32>>;
2686                            (10, visibility_range_binding.clone()),
2687                            // @group(0) @binding(0) var<uniform> view: View;
2688                            (0, view_uniforms_binding.clone()),
2689                            // @group(0) @binding(11) var depth_pyramid:
2690                            // texture_2d<f32>;
2691                            (11, &view_depth_pyramid.all_mips),
2692                            // @group(0) @binding(2) var<uniform>
2693                            // previous_view_uniforms: PreviousViewUniforms;
2694                            (
2695                                2,
2696                                BufferBinding {
2697                                    buffer: previous_view_buffer,
2698                                    offset: previous_view_uniform_offset.offset as u64,
2699                                    size: NonZeroU64::new(size_of::<PreviousViewData>() as u64),
2700                                },
2701                            ),
2702                            // @group(0) @binding(12) var<storage, read_write>
2703                            // late_preprocess_work_items:
2704                            // array<PreprocessWorkItem>;
2705                            (
2706                                12,
2707                                BufferBinding {
2708                                    buffer: late_non_indexed_work_item_buffer,
2709                                    offset: 0,
2710                                    size: non_indexed_work_item_buffer_size,
2711                                },
2712                            ),
2713                            // @group(0) @binding(13) var<storage, read_write>
2714                            // late_preprocess_work_item_indirect_parameters:
2715                            // array<LatePreprocessWorkItemIndirectParameters>;
2716                            (
2717                                13,
2718                                BufferBinding {
2719                                    buffer: late_non_indexed_indirect_parameters_buffer,
2720                                    offset: 0,
2721                                    size: NonZeroU64::new(
2722                                        late_non_indexed_indirect_parameters_buffer.size(),
2723                                    ),
2724                                },
2725                            ),
2726                        )),
2727                    ),
2728                )
2729            }
2730            _ => None,
2731        }
2732    }
2733
2734    /// Creates the bind group for the second phase of mesh preprocessing of
2735    /// indexed meshes when GPU occlusion culling is enabled.
2736    fn create_indirect_occlusion_culling_late_indexed_bind_group(
2737        &self,
2738        view_depth_pyramid: &ViewDepthPyramid,
2739        previous_view_uniform_offset: &PreviousViewUniformOffset,
2740        late_indexed_work_item_buffer: &UninitBufferVec<PreprocessWorkItem>,
2741    ) -> Option<BindGroup> {
2742        let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2743        let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2744        let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2745        let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
2746
2747        match (
2748            self.phase_indirect_parameters_buffers
2749                .indexed
2750                .metadata_buffer(),
2751            late_indexed_work_item_buffer.buffer(),
2752            self.late_indexed_indirect_parameters_buffer.buffer(),
2753        ) {
2754            (
2755                Some(indexed_metadata_buffer),
2756                Some(late_indexed_work_item_gpu_buffer),
2757                Some(late_indexed_indirect_parameters_buffer),
2758            ) => {
2759                // Don't use `as_entire_binding()` here; the shader reads the array
2760                // length and the underlying buffer may be longer than the actual size
2761                // of the vector.
2762                let late_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2763                    late_indexed_work_item_buffer.len() as u64
2764                        * u64::from(PreprocessWorkItem::min_size()),
2765                )
2766                .ok();
2767
2768                Some(
2769                    self.render_device.create_bind_group(
2770                        "preprocess_late_indexed_gpu_occlusion_culling_bind_group",
2771                        &self.pipeline_cache.get_bind_group_layout(
2772                            &self
2773                                .pipelines
2774                                .late_gpu_occlusion_culling_preprocess
2775                                .bind_group_layout,
2776                        ),
2777                        &BindGroupEntries::with_indices((
2778                            // @group(0) @binding(3) var<storage> current_input:
2779                            // array<MeshInput>;
2780                            (3, self.current_input_buffer.as_entire_binding()),
2781                            // @group(0) @binding(4) var<storage>
2782                            // previous_input: array<MeshInput>;
2783                            (4, self.previous_input_buffer.as_entire_binding()),
2784                            // @group(0) @binding(5) var<storage> work_items:
2785                            // array<PreprocessWorkItem>;
2786                            (
2787                                5,
2788                                BindingResource::Buffer(BufferBinding {
2789                                    buffer: late_indexed_work_item_gpu_buffer,
2790                                    offset: 0,
2791                                    size: late_indexed_work_item_buffer_size,
2792                                }),
2793                            ),
2794                            // @group(0) @binding(6) var<storage, read_write>
2795                            // output: array<Mesh>;
2796                            (6, self.data_buffer.as_entire_binding()),
2797                            // @group(0) @binding(7) var<storage>
2798                            // indirect_parameters_metadata:
2799                            // array<IndirectParametersMetadata>;
2800                            (7, indexed_metadata_buffer.as_entire_binding()),
2801                            // @group(0) @binding(9) var<storage>
2802                            // mesh_culling_data: array<MeshCullingData>;
2803                            (9, mesh_culling_data_buffer.as_entire_binding()),
2804                            // @group(0) @binding(10) var<storage>
2805                            // visibility_ranges: array<vec4<f32>>;
2806                            (10, visibility_range_binding.clone()),
2807                            // @group(0) @binding(0) var<uniform> view: View;
2808                            (0, view_uniforms_binding.clone()),
2809                            // @group(0) @binding(11) var depth_pyramid:
2810                            // texture_2d<f32>;
2811                            (11, &view_depth_pyramid.all_mips),
2812                            // @group(0) @binding(2) var<uniform>
2813                            // previous_view_uniforms: PreviousViewUniforms;
2814                            (
2815                                2,
2816                                BufferBinding {
2817                                    buffer: previous_view_buffer,
2818                                    offset: previous_view_uniform_offset.offset as u64,
2819                                    size: NonZeroU64::new(size_of::<PreviousViewData>() as u64),
2820                                },
2821                            ),
2822                            // @group(0) @binding(13) var<storage, read_write>
2823                            // late_preprocess_work_item_indirect_parameters:
2824                            // array<LatePreprocessWorkItemIndirectParameters>;
2825                            (
2826                                13,
2827                                BufferBinding {
2828                                    buffer: late_indexed_indirect_parameters_buffer,
2829                                    offset: 0,
2830                                    size: NonZeroU64::new(
2831                                        late_indexed_indirect_parameters_buffer.size(),
2832                                    ),
2833                                },
2834                            ),
2835                        )),
2836                    ),
2837                )
2838            }
2839            _ => None,
2840        }
2841    }
2842
2843    /// Creates the bind group for the second phase of mesh preprocessing of
2844    /// non-indexed meshes when GPU occlusion culling is enabled.
2845    fn create_indirect_occlusion_culling_late_non_indexed_bind_group(
2846        &self,
2847        view_depth_pyramid: &ViewDepthPyramid,
2848        previous_view_uniform_offset: &PreviousViewUniformOffset,
2849        late_non_indexed_work_item_buffer: &UninitBufferVec<PreprocessWorkItem>,
2850    ) -> Option<BindGroup> {
2851        let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2852        let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2853        let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2854        let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
2855
2856        match (
2857            self.phase_indirect_parameters_buffers
2858                .non_indexed
2859                .metadata_buffer(),
2860            late_non_indexed_work_item_buffer.buffer(),
2861            self.late_non_indexed_indirect_parameters_buffer.buffer(),
2862        ) {
2863            (
2864                Some(non_indexed_metadata_buffer),
2865                Some(non_indexed_work_item_gpu_buffer),
2866                Some(late_non_indexed_indirect_parameters_buffer),
2867            ) => {
2868                // Don't use `as_entire_binding()` here; the shader reads the array
2869                // length and the underlying buffer may be longer than the actual size
2870                // of the vector.
2871                let non_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2872                    late_non_indexed_work_item_buffer.len() as u64
2873                        * u64::from(PreprocessWorkItem::min_size()),
2874                )
2875                .ok();
2876
2877                Some(
2878                    self.render_device.create_bind_group(
2879                        "preprocess_late_non_indexed_gpu_occlusion_culling_bind_group",
2880                        &self.pipeline_cache.get_bind_group_layout(
2881                            &self
2882                                .pipelines
2883                                .late_gpu_occlusion_culling_preprocess
2884                                .bind_group_layout,
2885                        ),
2886                        &BindGroupEntries::with_indices((
2887                            // @group(0) @binding(3) var<storage> current_input:
2888                            // array<MeshInput>;
2889                            (3, self.current_input_buffer.as_entire_binding()),
2890                            // @group(0) @binding(4) var<storage>
2891                            // previous_input: array<MeshInput>;
2892                            (4, self.previous_input_buffer.as_entire_binding()),
2893                            // @group(0) @binding(5) var<storage> work_items:
2894                            // array<PreprocessWorkItem>;
2895                            (
2896                                5,
2897                                BindingResource::Buffer(BufferBinding {
2898                                    buffer: non_indexed_work_item_gpu_buffer,
2899                                    offset: 0,
2900                                    size: non_indexed_work_item_buffer_size,
2901                                }),
2902                            ),
2903                            // @group(0) @binding(6) var<storage, read_write>
2904                            // output: array<Mesh>;
2905                            (6, self.data_buffer.as_entire_binding()),
2906                            // @group(0) @binding(7) var<storage>
2907                            // indirect_parameters_metadata:
2908                            // array<IndirectParametersMetadata>;
2909                            (7, non_indexed_metadata_buffer.as_entire_binding()),
2910                            // @group(0) @binding(9) var<storage>
2911                            // mesh_culling_data: array<MeshCullingData>;
2912                            (9, mesh_culling_data_buffer.as_entire_binding()),
2913                            // @group(0) @binding(10) var<storage>
2914                            // visibility_ranges: array<vec4<f32>>;
2915                            (10, visibility_range_binding.clone()),
2916                            // @group(0) @binding(0) var<uniform> view: View;
2917                            (0, view_uniforms_binding.clone()),
2918                            // @group(0) @binding(11) var depth_pyramid:
2919                            // texture_2d<f32>;
2920                            (11, &view_depth_pyramid.all_mips),
2921                            // @group(0) @binding(2) var<uniform>
2922                            // previous_view_uniforms: PreviousViewUniforms;
2923                            (
2924                                2,
2925                                BufferBinding {
2926                                    buffer: previous_view_buffer,
2927                                    offset: previous_view_uniform_offset.offset as u64,
2928                                    size: NonZeroU64::new(size_of::<PreviousViewData>() as u64),
2929                                },
2930                            ),
2931                            // @group(0) @binding(13) var<storage, read>
2932                            // late_preprocess_work_item_indirect_parameters:
2933                            // array<LatePreprocessWorkItemIndirectParameters>;
2934                            (
2935                                13,
2936                                BufferBinding {
2937                                    buffer: late_non_indexed_indirect_parameters_buffer,
2938                                    offset: 0,
2939                                    size: NonZeroU64::new(
2940                                        late_non_indexed_indirect_parameters_buffer.size(),
2941                                    ),
2942                                },
2943                            ),
2944                        )),
2945                    ),
2946                )
2947            }
2948            _ => None,
2949        }
2950    }
2951
2952    /// Creates the bind groups for mesh preprocessing when GPU frustum culling
2953    /// is enabled, but GPU occlusion culling is disabled.
2954    fn create_indirect_frustum_culling_preprocess_bind_groups(
2955        &self,
2956        indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2957        non_indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2958    ) -> Option<PhasePreprocessBindGroups> {
2959        Some(PhasePreprocessBindGroups::IndirectFrustumCulling {
2960            indexed: self
2961                .create_indirect_frustum_culling_indexed_bind_group(indexed_work_item_buffer),
2962            non_indexed: self.create_indirect_frustum_culling_non_indexed_bind_group(
2963                non_indexed_work_item_buffer,
2964            ),
2965        })
2966    }
2967
2968    /// Creates the bind group for mesh preprocessing of indexed meshes when GPU
2969    /// frustum culling is enabled, but GPU occlusion culling is disabled.
2970    fn create_indirect_frustum_culling_indexed_bind_group(
2971        &self,
2972        indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2973    ) -> Option<BindGroup> {
2974        let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2975        let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2976        let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2977
2978        match (
2979            self.phase_indirect_parameters_buffers
2980                .indexed
2981                .metadata_buffer(),
2982            indexed_work_item_buffer.buffer(),
2983        ) {
2984            (Some(indexed_metadata_buffer), Some(indexed_work_item_gpu_buffer)) => {
2985                // Don't use `as_entire_binding()` here; the shader reads the array
2986                // length and the underlying buffer may be longer than the actual size
2987                // of the vector.
2988                let indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2989                    indexed_work_item_buffer.len() as u64
2990                        * u64::from(PreprocessWorkItem::min_size()),
2991                )
2992                .ok();
2993
2994                Some(
2995                    self.render_device.create_bind_group(
2996                        "preprocess_gpu_indexed_frustum_culling_bind_group",
2997                        &self.pipeline_cache.get_bind_group_layout(
2998                            &self
2999                                .pipelines
3000                                .gpu_frustum_culling_preprocess
3001                                .bind_group_layout,
3002                        ),
3003                        &BindGroupEntries::with_indices((
3004                            (3, self.current_input_buffer.as_entire_binding()),
3005                            (4, self.previous_input_buffer.as_entire_binding()),
3006                            (
3007                                5,
3008                                BindingResource::Buffer(BufferBinding {
3009                                    buffer: indexed_work_item_gpu_buffer,
3010                                    offset: 0,
3011                                    size: indexed_work_item_buffer_size,
3012                                }),
3013                            ),
3014                            (6, self.data_buffer.as_entire_binding()),
3015                            (7, indexed_metadata_buffer.as_entire_binding()),
3016                            (9, mesh_culling_data_buffer.as_entire_binding()),
3017                            (10, visibility_range_binding.clone()),
3018                            (0, view_uniforms_binding.clone()),
3019                        )),
3020                    ),
3021                )
3022            }
3023            _ => None,
3024        }
3025    }
3026
3027    /// Creates the bind group for mesh preprocessing of non-indexed meshes when
3028    /// GPU frustum culling is enabled, but GPU occlusion culling is disabled.
3029    fn create_indirect_frustum_culling_non_indexed_bind_group(
3030        &self,
3031        non_indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
3032    ) -> Option<BindGroup> {
3033        let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
3034        let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
3035        let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
3036
3037        match (
3038            self.phase_indirect_parameters_buffers
3039                .non_indexed
3040                .metadata_buffer(),
3041            non_indexed_work_item_buffer.buffer(),
3042        ) {
3043            (Some(non_indexed_metadata_buffer), Some(non_indexed_work_item_gpu_buffer)) => {
3044                // Don't use `as_entire_binding()` here; the shader reads the array
3045                // length and the underlying buffer may be longer than the actual size
3046                // of the vector.
3047                let non_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
3048                    non_indexed_work_item_buffer.len() as u64
3049                        * u64::from(PreprocessWorkItem::min_size()),
3050                )
3051                .ok();
3052
3053                Some(
3054                    self.render_device.create_bind_group(
3055                        "preprocess_gpu_non_indexed_frustum_culling_bind_group",
3056                        &self.pipeline_cache.get_bind_group_layout(
3057                            &self
3058                                .pipelines
3059                                .gpu_frustum_culling_preprocess
3060                                .bind_group_layout,
3061                        ),
3062                        &BindGroupEntries::with_indices((
3063                            // @group(0) @binding(3) var<storage> current_input:
3064                            // array<MeshInput>;
3065                            (3, self.current_input_buffer.as_entire_binding()),
3066                            // @group(0) @binding(4) var<storage>
3067                            // previous_input: array<MeshInput>;
3068                            (4, self.previous_input_buffer.as_entire_binding()),
3069                            // @group(0) @binding(5) var<storage> work_items:
3070                            // array<PreprocessWorkItem>;
3071                            (
3072                                5,
3073                                BindingResource::Buffer(BufferBinding {
3074                                    buffer: non_indexed_work_item_gpu_buffer,
3075                                    offset: 0,
3076                                    size: non_indexed_work_item_buffer_size,
3077                                }),
3078                            ),
3079                            // @group(0) @binding(6) var<storage, read_write>
3080                            // output: array<Mesh>;
3081                            (6, self.data_buffer.as_entire_binding()),
3082                            // @group(0) @binding(7) var<storage>
3083                            // indirect_parameters_metadata:
3084                            // array<IndirectParametersMetadata>;
3085                            (7, non_indexed_metadata_buffer.as_entire_binding()),
3086                            // @group(0) @binding(9) var<storage>
3087                            // mesh_culling_data: array<MeshCullingData>;
3088                            (9, mesh_culling_data_buffer.as_entire_binding()),
3089                            // @group(0) @binding(10) var<storage>
3090                            // visibility_ranges: array<vec4<f32>>;
3091                            (10, visibility_range_binding.clone()),
3092                            // @group(0) @binding(0) var<uniform> view: View;
3093                            (0, view_uniforms_binding.clone()),
3094                        )),
3095                    ),
3096                )
3097            }
3098            _ => None,
3099        }
3100    }
3101}
3102
3103/// A system that creates bind groups from the indirect parameters metadata and
3104/// data buffers for the indirect batch set reset shader and the indirect
3105/// parameter building shader.
3106fn create_build_indirect_parameters_bind_groups(
3107    commands: &mut Commands,
3108    render_device: &RenderDevice,
3109    pipeline_cache: &PipelineCache,
3110    pipelines: &PreprocessPipelines,
3111    current_input_buffer: &Buffer,
3112    indirect_parameters_buffers: &IndirectParametersBuffers,
3113    indirect_parameters_build_jobs: &IndirectParametersBuildJobs,
3114) {
3115    let mut build_indirect_parameters_bind_groups = BuildIndirectParametersBindGroups::new();
3116
3117    for (phase_type_id, phase_indirect_parameters_buffer) in indirect_parameters_buffers.iter() {
3118        build_indirect_parameters_bind_groups.insert(
3119            *phase_type_id,
3120            PhaseBuildIndirectParametersBindGroups {
3121                reset_indexed_indirect_batch_sets: phase_indirect_parameters_buffer
3122                    .indexed
3123                    .batch_sets_buffer()
3124                    .map(|indexed_batch_sets_buffer| {
3125                        render_device.create_bind_group(
3126                            "reset_indexed_indirect_batch_sets_bind_group",
3127                            // The early bind group is good for the main phase and late
3128                            // phase too. They bind the same buffers.
3129                            &pipeline_cache.get_bind_group_layout(
3130                                &pipelines
3131                                    .early_phase
3132                                    .reset_indirect_batch_sets
3133                                    .bind_group_layout,
3134                            ),
3135                            &BindGroupEntries::sequential((
3136                                indexed_batch_sets_buffer.as_entire_binding(),
3137                            )),
3138                        )
3139                    }),
3140
3141                reset_non_indexed_indirect_batch_sets: phase_indirect_parameters_buffer
3142                    .non_indexed
3143                    .batch_sets_buffer()
3144                    .map(|non_indexed_batch_sets_buffer| {
3145                        render_device.create_bind_group(
3146                            "reset_non_indexed_indirect_batch_sets_bind_group",
3147                            // The early bind group is good for the main phase and late
3148                            // phase too. They bind the same buffers.
3149                            &pipeline_cache.get_bind_group_layout(
3150                                &pipelines
3151                                    .early_phase
3152                                    .reset_indirect_batch_sets
3153                                    .bind_group_layout,
3154                            ),
3155                            &BindGroupEntries::sequential((
3156                                non_indexed_batch_sets_buffer.as_entire_binding(),
3157                            )),
3158                        )
3159                    }),
3160
3161                build_indexed_indirect: match (
3162                    phase_indirect_parameters_buffer.indexed.metadata_buffer(),
3163                    phase_indirect_parameters_buffer.indexed.data_buffer(),
3164                    phase_indirect_parameters_buffer.indexed.batch_sets_buffer(),
3165                    indirect_parameters_build_jobs.buffer(),
3166                ) {
3167                    (
3168                        Some(indexed_indirect_parameters_metadata_buffer),
3169                        Some(indexed_indirect_parameters_data_buffer),
3170                        Some(indexed_batch_sets_buffer),
3171                        Some(indirect_parameters_build_job_buffer),
3172                    ) => Some(
3173                        render_device.create_bind_group(
3174                            "build_indexed_indirect_parameters_bind_group",
3175                            // The frustum culling bind group is good for occlusion culling
3176                            // too. They bind the same buffers.
3177                            &pipeline_cache.get_bind_group_layout(
3178                                &pipelines
3179                                    .gpu_frustum_culling_build_indexed_indirect_params
3180                                    .bind_group_layout,
3181                            ),
3182                            &BindGroupEntries::with_indices((
3183                                // @group(0) @binding(0) var<storage>
3184                                // current_input: array<MeshInput>;
3185                                (0, current_input_buffer.as_entire_binding()),
3186                                // @group(0) @binding(1) var<storage>
3187                                // indirect_parameters_metadata:
3188                                // array<IndirectParametersMetadata>;
3189                                (
3190                                    1,
3191                                    indexed_indirect_parameters_metadata_buffer.as_entire_binding(),
3192                                ),
3193                                // @group(0) @binding(3) var<storage,
3194                                // read_write> indirect_batch_sets:
3195                                // array<IndirectBatchSet>;
3196                                (3, indexed_batch_sets_buffer.as_entire_binding()),
3197                                // @group(0) @binding(4) var<uniform>
3198                                // indirect_parameters_build_job:
3199                                // IndirectParametersBuildJob;
3200                                (
3201                                    4,
3202                                    BindingResource::Buffer(BufferBinding {
3203                                        buffer: indirect_parameters_build_job_buffer,
3204                                        offset: 0,
3205                                        size: NonZeroU64::new(
3206                                            size_of::<IndirectParametersBuildJob>() as u64,
3207                                        ),
3208                                    }),
3209                                ),
3210                                // @group(0) @binding(5) var<storage,
3211                                // read_write> indirect_parameters:
3212                                // array<IndirectParametersIndexed>;
3213                                (
3214                                    5,
3215                                    indexed_indirect_parameters_data_buffer.as_entire_binding(),
3216                                ),
3217                            )),
3218                        ),
3219                    ),
3220                    _ => None,
3221                },
3222
3223                build_non_indexed_indirect: match (
3224                    phase_indirect_parameters_buffer
3225                        .non_indexed
3226                        .metadata_buffer(),
3227                    phase_indirect_parameters_buffer.non_indexed.data_buffer(),
3228                    phase_indirect_parameters_buffer
3229                        .non_indexed
3230                        .batch_sets_buffer(),
3231                    indirect_parameters_build_jobs.buffer(),
3232                ) {
3233                    (
3234                        Some(non_indexed_indirect_parameters_metadata_buffer),
3235                        Some(non_indexed_indirect_parameters_data_buffer),
3236                        Some(non_indexed_batch_sets_buffer),
3237                        Some(indirect_parameters_build_job_buffer),
3238                    ) => Some(
3239                        render_device.create_bind_group(
3240                            "build_non_indexed_indirect_parameters_bind_group",
3241                            // The frustum culling bind group is good for occlusion culling
3242                            // too. They bind the same buffers.
3243                            &pipeline_cache.get_bind_group_layout(
3244                                &pipelines
3245                                    .gpu_frustum_culling_build_non_indexed_indirect_params
3246                                    .bind_group_layout,
3247                            ),
3248                            &BindGroupEntries::with_indices((
3249                                // @group(0) @binding(0) var<storage>
3250                                // current_input: array<MeshInput>;
3251                                (0, current_input_buffer.as_entire_binding()),
3252                                // @group(0) @binding(1) var<storage>
3253                                // indirect_parameters_metadata:
3254                                // array<IndirectParametersMetadata>;
3255                                //
3256                                // Don't use `as_entire_binding` here; the shader reads
3257                                // the length and `RawBufferVec` overallocates.
3258                                (
3259                                    1,
3260                                    BufferBinding {
3261                                        buffer: non_indexed_indirect_parameters_metadata_buffer,
3262                                        offset: 0,
3263                                        size: NonZeroU64::new(
3264                                            phase_indirect_parameters_buffer
3265                                                .non_indexed
3266                                                .batch_count()
3267                                                as u64
3268                                                * size_of::<IndirectParametersMetadata>() as u64,
3269                                        ),
3270                                    },
3271                                ),
3272                                // @group(0) @binding(3) var<storage,
3273                                // read_write> indirect_batch_sets:
3274                                // array<IndirectBatchSet>;
3275                                (3, non_indexed_batch_sets_buffer.as_entire_binding()),
3276                                // @group(0) @binding(4) var<uniform>
3277                                // indirect_parameters_build_job:
3278                                // IndirectParametersBuildJob;
3279                                (
3280                                    4,
3281                                    BindingResource::Buffer(BufferBinding {
3282                                        buffer: indirect_parameters_build_job_buffer,
3283                                        offset: 0,
3284                                        size: NonZeroU64::new(
3285                                            size_of::<IndirectParametersBuildJob>() as u64,
3286                                        ),
3287                                    }),
3288                                ),
3289                                // @group(0) @binding(5) var<storage,
3290                                // read_write> indirect_parameters:
3291                                // array<IndirectParametersNonIndexed>;
3292                                (
3293                                    5,
3294                                    non_indexed_indirect_parameters_data_buffer.as_entire_binding(),
3295                                ),
3296                            )),
3297                        ),
3298                    ),
3299                    _ => None,
3300                },
3301            },
3302        );
3303    }
3304
3305    commands.insert_resource(build_indirect_parameters_bind_groups);
3306}
3307
3308/// Creates all bind groups needed to run the `unpack_bins` shader for all the
3309/// phases for a single view.
3310fn create_bin_unpacking_bind_groups(
3311    bin_unpacking_bind_groups: &mut BinUnpackingBindGroups,
3312    render_device: &RenderDevice,
3313    pipeline_cache: &PipelineCache,
3314    preprocess_pipelines: &PreprocessPipelines,
3315    indirect_parameters_buffers: &IndirectParametersBuffers,
3316    phase_instance_buffers: &TypeIdHashMap<UntypedPhaseBatchedInstanceBuffers<MeshUniform>>,
3317    scene_unpacking_buffers: &SceneUnpackingBuffers,
3318    view_entity: &RetainedViewEntity,
3319) {
3320    let Some(bin_unpacking_metadata_buffer) =
3321        scene_unpacking_buffers.bin_unpacking_metadata.buffer()
3322    else {
3323        return;
3324    };
3325
3326    // We run the bin unpacking shader once per phase, so loop over all phases.
3327    for phase_type_id in indirect_parameters_buffers.keys() {
3328        // Fetch the buffers we need.
3329        let Some(phase_batched_instance_buffers) = phase_instance_buffers.get(phase_type_id) else {
3330            continue;
3331        };
3332        let Some(work_item_buffers) = phase_batched_instance_buffers
3333            .work_item_buffers
3334            .get(view_entity)
3335        else {
3336            continue;
3337        };
3338        let Some(view_phase_bin_unpacking_buffers) = scene_unpacking_buffers
3339            .view_phase_buffers
3340            .get(&SceneUnpackingBuffersKey {
3341                phase: *phase_type_id,
3342                view: *view_entity,
3343            })
3344        else {
3345            continue;
3346        };
3347
3348        // Fetch the work item buffers.
3349        let maybe_indexed_work_item_buffer = match *work_item_buffers {
3350            PreprocessWorkItemBuffers::Direct(ref raw_buffer_vec) => raw_buffer_vec.buffer(),
3351            PreprocessWorkItemBuffers::Indirect { ref indexed, .. } => indexed.buffer(),
3352        };
3353        let maybe_non_indexed_work_item_buffer = match *work_item_buffers {
3354            PreprocessWorkItemBuffers::Direct(ref raw_buffer_vec) => raw_buffer_vec.buffer(),
3355            PreprocessWorkItemBuffers::Indirect {
3356                ref non_indexed, ..
3357            } => non_indexed.buffer(),
3358        };
3359
3360        // Create the actual bind groups.
3361        bin_unpacking_bind_groups.insert(
3362            SceneUnpackingBuffersKey {
3363                phase: *phase_type_id,
3364                view: *view_entity,
3365            },
3366            ViewPhaseBinUnpackingBindGroups {
3367                indexed: match maybe_indexed_work_item_buffer {
3368                    Some(indexed_work_item_buffer) => view_phase_bin_unpacking_buffers
3369                        .indexed_unpacking_jobs
3370                        .iter()
3371                        .map(|job| {
3372                            create_bin_unpacking_bind_group(
3373                                render_device,
3374                                preprocess_pipelines,
3375                                pipeline_cache,
3376                                job,
3377                                bin_unpacking_metadata_buffer,
3378                                indexed_work_item_buffer,
3379                                true,
3380                            )
3381                        })
3382                        .collect(),
3383                    None => ::alloc::vec::Vec::new()vec![],
3384                },
3385                non_indexed: match maybe_non_indexed_work_item_buffer {
3386                    Some(non_indexed_work_item_buffer) => view_phase_bin_unpacking_buffers
3387                        .non_indexed_unpacking_jobs
3388                        .iter()
3389                        .map(|job| {
3390                            create_bin_unpacking_bind_group(
3391                                render_device,
3392                                preprocess_pipelines,
3393                                pipeline_cache,
3394                                job,
3395                                bin_unpacking_metadata_buffer,
3396                                non_indexed_work_item_buffer,
3397                                false,
3398                            )
3399                        })
3400                        .collect(),
3401                    None => ::alloc::vec::Vec::new()vec![],
3402                },
3403            },
3404        );
3405    }
3406}
3407
3408/// Creates a bind group for the bin unpacking shader for a single (view, phase,
3409/// mesh indexed-ness) combination.
3410fn create_bin_unpacking_bind_group(
3411    render_device: &RenderDevice,
3412    preprocess_pipelines: &PreprocessPipelines,
3413    pipeline_cache: &PipelineCache,
3414    job: &SceneUnpackingJob,
3415    bin_unpacking_metadata_buffer: &Buffer,
3416    work_item_buffer: &Buffer,
3417    indexed: bool,
3418) -> ViewPhaseBinUnpackingBindGroup {
3419    let bind_group = render_device.create_bind_group(
3420        if indexed {
3421            "bin unpacking indexed bind group"
3422        } else {
3423            "bin unpacking non-indexed bind group"
3424        },
3425        &pipeline_cache
3426            .get_bind_group_layout(&preprocess_pipelines.bin_unpacking.bind_group_layout),
3427        &BindGroupEntries::sequential((
3428            // @group(0) @binding(0) var<uniform>
3429            // bin_unpacking_metadata:
3430            // BinUnpackingMetadata;
3431            BindingResource::Buffer(BufferBinding {
3432                buffer: bin_unpacking_metadata_buffer,
3433                offset: job.bin_unpacking_metadata_index.uniform_offset() as u64,
3434                size: NonZeroU64::new(size_of::<GpuBinUnpackingMetadata>() as u64),
3435            }),
3436            // @group(0) @binding(1) var<storage>
3437            // binned_mesh_instances:
3438            // array<BinnedMeshInstance>;
3439            job.render_binned_mesh_instance_buffer.as_entire_binding(),
3440            // @group(0) @binding(2) var<storage,
3441            // read_write> preprocess_work_items:
3442            // array<PreprocessWorkItem>;
3443            work_item_buffer.as_entire_binding(),
3444            // @group(0) @binding(3) var<storage> bin_metadata:
3445            // array<BinMetadata>;
3446            job.bin_metadata_buffer.as_entire_binding(),
3447            // @group(0) @binding(4) var<storage>
3448            // bin_index_to_bin_metadata_index: array<u32>;
3449            job.bin_index_to_bin_metadata_index_buffer
3450                .as_entire_binding(),
3451        )),
3452    );
3453    ViewPhaseBinUnpackingBindGroup {
3454        metadata_index: job.bin_unpacking_metadata_index,
3455        bind_group,
3456        mesh_instance_count: job.mesh_instance_count,
3457    }
3458}
3459
3460/// Creates all bind groups needed to run the `allocate_uniforms` shader for all
3461/// the phases for a single view.
3462fn create_uniform_allocation_bind_groups(
3463    uniform_allocation_bind_groups: &mut UniformAllocationBindGroups,
3464    render_device: &RenderDevice,
3465    pipeline_cache: &PipelineCache,
3466    preprocess_pipelines: &PreprocessPipelines,
3467    indirect_parameters_buffers: &IndirectParametersBuffers,
3468    scene_unpacking_buffers: &SceneUnpackingBuffers,
3469    view_entity: &RetainedViewEntity,
3470) {
3471    let Some(uniform_allocation_metadata_buffer) =
3472        scene_unpacking_buffers.uniform_allocation_metadata.buffer()
3473    else {
3474        return;
3475    };
3476
3477    for (phase_type_id, phase_indirect_parameters_buffers) in indirect_parameters_buffers.iter() {
3478        let Some(view_phase_bin_unpacking_buffers) = scene_unpacking_buffers
3479            .view_phase_buffers
3480            .get(&SceneUnpackingBuffersKey {
3481                phase: *phase_type_id,
3482                view: *view_entity,
3483            })
3484        else {
3485            continue;
3486        };
3487
3488        // Create the actual bind groups.
3489        uniform_allocation_bind_groups.insert(
3490            SceneUnpackingBuffersKey {
3491                phase: *phase_type_id,
3492                view: *view_entity,
3493            },
3494            ViewPhaseUniformAllocationBindGroups {
3495                indexed: match phase_indirect_parameters_buffers.indexed.metadata_buffer() {
3496                    None => ::alloc::vec::Vec::new()vec![],
3497                    Some(indexed_indirect_parameters_metadata_buffer) => {
3498                        view_phase_bin_unpacking_buffers
3499                            .indexed_unpacking_jobs
3500                            .iter()
3501                            .map(|job| {
3502                                create_uniform_allocation_bind_group(
3503                                    render_device,
3504                                    preprocess_pipelines,
3505                                    pipeline_cache,
3506                                    job,
3507                                    uniform_allocation_metadata_buffer,
3508                                    indexed_indirect_parameters_metadata_buffer,
3509                                    true,
3510                                )
3511                            })
3512                            .collect()
3513                    }
3514                },
3515                non_indexed: match phase_indirect_parameters_buffers
3516                    .non_indexed
3517                    .metadata_buffer()
3518                {
3519                    None => ::alloc::vec::Vec::new()vec![],
3520                    Some(non_indexed_indirect_parameters_metadata_buffer) => {
3521                        view_phase_bin_unpacking_buffers
3522                            .non_indexed_unpacking_jobs
3523                            .iter()
3524                            .map(|job| {
3525                                create_uniform_allocation_bind_group(
3526                                    render_device,
3527                                    preprocess_pipelines,
3528                                    pipeline_cache,
3529                                    job,
3530                                    uniform_allocation_metadata_buffer,
3531                                    non_indexed_indirect_parameters_metadata_buffer,
3532                                    false,
3533                                )
3534                            })
3535                            .collect()
3536                    }
3537                },
3538            },
3539        );
3540    }
3541}
3542
3543/// Creates a bind group for the uniform allocation shader for a single (view,
3544/// phase, mesh indexed-ness) combination.
3545fn create_uniform_allocation_bind_group(
3546    render_device: &RenderDevice,
3547    preprocess_pipelines: &PreprocessPipelines,
3548    pipeline_cache: &PipelineCache,
3549    job: &SceneUnpackingJob,
3550    uniform_allocation_metadata_buffer: &Buffer,
3551    indirect_parameters_metadata_buffer: &Buffer,
3552    indexed: bool,
3553) -> ViewPhaseUniformAllocationBindGroup {
3554    let bind_group = render_device.create_bind_group(
3555        if indexed {
3556            "uniform allocation indexed bind group"
3557        } else {
3558            "uniform allocation non-indexed bind group"
3559        },
3560        &pipeline_cache.get_bind_group_layout(
3561            // All the pipelines' bind group layouts should be identical.
3562            &preprocess_pipelines
3563                .uniform_allocation
3564                .local_scan
3565                .bind_group_layout,
3566        ),
3567        &BindGroupEntries::sequential((
3568            // @group(0) @binding(0) var<uniform> allocate_uniforms_metadata:
3569            // AllocateUniformsMetadata;
3570            BindingResource::Buffer(BufferBinding {
3571                buffer: uniform_allocation_metadata_buffer,
3572                offset: job.uniform_allocation_metadata_index.uniform_offset() as u64,
3573                size: NonZeroU64::new(size_of::<GpuUniformAllocationMetadata>() as u64),
3574            }),
3575            // @group(0) @binding(1) var<storage> bin_metadata:
3576            // array<BinMetadata>;
3577            job.bin_metadata_buffer.as_entire_binding(),
3578            // @group(0) @binding(2) var<storage, read_write>
3579            // indirect_parameters_metadata: array<IndirectParametersMetadata>;
3580            indirect_parameters_metadata_buffer.as_entire_binding(),
3581            // @group(0) @binding(3) var<storage, read_write> fan_buffer:
3582            // array<u32>;
3583            job.fan_buffer.as_entire_binding(),
3584        )),
3585    );
3586    ViewPhaseUniformAllocationBindGroup {
3587        metadata_index: job.uniform_allocation_metadata_index,
3588        bind_group,
3589        bin_count: job.bin_count,
3590    }
3591}
3592
3593/// Writes the information needed to do GPU mesh culling to the GPU.
3594pub fn write_mesh_culling_data_buffer(
3595    render_device: Res<RenderDevice>,
3596    render_queue: Res<RenderQueue>,
3597    mut mesh_culling_data_buffer: ResMut<MeshCullingDataBuffer>,
3598    pipeline_cache: Res<PipelineCache>,
3599    mut sparse_buffer_update_jobs: ResMut<SparseBufferUpdateJobs>,
3600    mut sparse_buffer_update_bind_groups: ResMut<SparseBufferUpdateBindGroups>,
3601    sparse_buffer_update_pipelines: Res<SparseBufferUpdatePipelines>,
3602) {
3603    mesh_culling_data_buffer.write_buffers(&render_device, &render_queue);
3604    mesh_culling_data_buffer.prepare_to_populate_buffers(
3605        &render_device,
3606        &pipeline_cache,
3607        &mut sparse_buffer_update_jobs,
3608        &mut sparse_buffer_update_bind_groups,
3609        &sparse_buffer_update_pipelines,
3610    );
3611}