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`].
89use core::num::{NonZero, NonZeroU64};
1011use bevy_app::{App, Plugin};
12use bevy_asset::{embedded_asset, load_embedded_asset, Handle};
13use bevy_core_pipeline::{
14deferred::node::late_deferred_prepass,
15 mip_generation::experimental::depth::{early_downsample_depth, ViewDepthPyramid},
16 prepass::{
17 node::{early_prepass, late_prepass},
18DeferredPrepass, DepthPrepass, MotionVectorPrepass, NormalPrepass, PreviousViewData,
19PreviousViewUniformOffset, PreviousViewUniforms,
20 },
21 schedule::{Core3d, Core3dSystems},
22};
23use bevy_derive::{Deref, DerefMut};
24use bevy_ecs::{
25component::Component,
26entity::Entity,
27prelude::resource_exists,
28 query::{Has, Or, With, Without},
29resource::Resource,
30 schedule::{common_conditions::any_match_filter, IntoScheduleConfigsas _},
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::{
39clear_scene_unpacking_buffers, BatchedInstanceBuffers, BinUnpackingMetadataIndex,
40BuildIndirectParametersMetadata, GpuBinMetadata, GpuBinUnpackingMetadata,
41GpuOcclusionCullingWorkItemBuffers, GpuPreprocessingMode, GpuPreprocessingSupport,
42GpuUniformAllocationMetadata, IndirectBatchSet, IndirectParametersBuffers,
43IndirectParametersBuildJob, IndirectParametersBuildJobs, IndirectParametersIndexed,
44IndirectParametersMetadata, IndirectParametersNonIndexed,
45LatePreprocessWorkItemIndirectParameters, PreprocessWorkItem, PreprocessWorkItemBuffers,
46SceneUnpackingBuffers, SceneUnpackingBuffersKey, SceneUnpackingJob,
47UniformAllocationMetadataIndex, UntypedPhaseBatchedInstanceBuffers,
48UntypedPhaseIndirectParametersBuffers,
49 },
50diagnostic::RecordDiagnosticsas _,
51occlusion_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},
55BindGroup, BindGroupEntries, BindGroupLayoutDescriptor, BindGroupLayoutEntries,
56BindingResource, Buffer, BufferBinding, BufferVec, CachedComputePipelineId,
57ComputePassDescriptor, ComputePipelineDescriptor, DynamicBindGroupLayoutEntries,
58PartialBufferVec, PipelineCache, RawBufferVec, ShaderStages, ShaderType,
59SparseBufferUpdateBindGroups, SparseBufferUpdateJobs, SparseBufferUpdatePipelines,
60SpecializedComputePipeline, SpecializedComputePipelines, TextureSampleType,
61UninitBufferVec,
62 },
63 renderer::{RenderContext, RenderDevice, RenderQueue, ViewQuery},
64settings::WgpuFeatures,
65 view::{
66ExtractedView, NoIndirectDrawing, RenderVisibilityRanges, RetainedViewEntity, ViewUniform,
67ViewUniformOffset, ViewUniforms,
68 },
69GpuResourceAppExt, Render, RenderApp, RenderSystems,
70};
71use bevy_shader::Shader;
72use bevy_utils::{default, TypeIdHashMap};
73use bitflags::bitflags;
74use smallvec::{smallvec, SmallVec};
75use tracing::warn;
7677use crate::{
78LightEntity, MeshCullingData, MeshCullingDataBuffer, MeshInputUniform, MeshUniform,
79PreviousMeshInputUniform,
80};
8182use super::{ShadowView, ViewLightEntities};
8384/// The GPU workgroup size.
85const WORKGROUP_SIZE: usize = 64;
8687/// 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.
96pub use_gpu_instance_buffer_builder: bool,
97}
9899/// 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.
105pub 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.
110pub 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.
115pub 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.
120pub 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.
123pub 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.
127pub 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.
130pub 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.
134pub late_phase: PreprocessPhasePipelines,
135/// Compute shader pipelines for the main color phase.
136pub main_phase: PreprocessPhasePipelines,
137/// Compute shader pipelines for the bin unpacking step.
138pub bin_unpacking: BinUnpackingPipeline,
139/// Compute shader pipelines for the uniform allocation step.
140pub uniform_allocation: UniformAllocationPipelines,
141}
142143/// 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.
150pub 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.
155pub 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.
160pub gpu_occlusion_culling_build_non_indexed_indirect_params: BuildIndirectParametersPipeline,
161}
162163/// The pipeline for the GPU mesh preprocessing shader.
164pub struct PreprocessPipeline {
165/// The bind group layout for the compute shader.
166pub bind_group_layout: BindGroupLayoutDescriptor,
167/// The shader asset handle.
168pub shader: Handle<Shader>,
169/// The pipeline ID for the compute shader.
170 ///
171 /// This gets filled in `prepare_preprocess_pipelines`.
172pub pipeline_id: Option<CachedComputePipelineId>,
173}
174175/// 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.
182pub bind_group_layout: BindGroupLayoutDescriptor,
183/// The shader asset handle.
184pub shader: Handle<Shader>,
185/// The pipeline ID for the compute shader.
186 ///
187 /// This gets filled in `prepare_preprocess_pipelines`.
188pub pipeline_id: Option<CachedComputePipelineId>,
189}
190191/// 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.
195pub bind_group_layout: BindGroupLayoutDescriptor,
196/// The shader asset handle.
197pub shader: Handle<Shader>,
198/// The pipeline ID for the compute shader.
199 ///
200 /// This gets filled in `prepare_preprocess_pipelines`.
201pub pipeline_id: Option<CachedComputePipelineId>,
202}
203204/// 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.
208pub bind_group_layout: BindGroupLayoutDescriptor,
209/// The shader asset handle.
210pub shader: Handle<Shader>,
211/// The pipeline ID for the compute shader.
212 ///
213 /// This gets filled in in the [`prepare_preprocess_pipelines`] system.
214pub pipeline_id: Option<CachedComputePipelineId>,
215}
216217/// 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.
227pub local_scan: UniformAllocationLocalScanPipeline,
228/// The pipeline for step 2: global scan.
229pub global_scan: UniformAllocationGlobalScanPipeline,
230/// The pipeline for step 3: fan.
231pub fan: UniformAllocationFanPipeline,
232}
233234/// 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.
239pub bind_group_layout: BindGroupLayoutDescriptor,
240/// The shader, also shared among all uniform allocation pipelines.
241pub shader: Handle<Shader>,
242/// The pipeline ID for the first step of the `allocate_uniforms` shader.
243pub pipeline_id_local_scan: Option<CachedComputePipelineId>,
244}
245246/// 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.
253pub bind_group_layout: BindGroupLayoutDescriptor,
254/// The shader, also shared among all uniform allocation pipelines.
255pub shader: Handle<Shader>,
256/// The pipeline ID for the second step of the `allocate_uniforms` shader.
257pub pipeline_id_global_scan: Option<CachedComputePipelineId>,
258}
259260/// 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.
267pub bind_group_layout: BindGroupLayoutDescriptor,
268/// The shader, also shared among all uniform allocation pipelines.
269pub shader: Handle<Shader>,
270/// The pipeline ID for the third step of the `allocate_uniforms` shader.
271pub pipeline_id_fan: Option<CachedComputePipelineId>,
272}
273274#[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)]
277pub struct PreprocessPipelineKey: u8 {
278/// Whether GPU frustum culling is in use.
279 ///
280 /// This `#define`'s `FRUSTUM_CULLING` in the shader.
281const FRUSTUM_CULLING = 1;
282/// Whether GPU two-phase occlusion culling is in use.
283 ///
284 /// This `#define`'s `OCCLUSION_CULLING` in the shader.
285const 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.
289const EARLY_PHASE = 4;
290 }
291292/// Specifies variants of the indirect parameter building shader.
293#[derive(Clone, Copy, PartialEq, Eq, Hash)]
294pub 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.
299const 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.
303const 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.
307const 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.
311const 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.
315const 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.
321const MAIN_PHASE = 32;
322 }
323}324325/// 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>);
333334/// 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.
343Direct(BindGroup),
344345/// 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.
350IndirectFrustumCulling {
351/// The bind group for indexed meshes.
352indexed: Option<BindGroup>,
353/// The bind group for non-indexed meshes.
354non_indexed: Option<BindGroup>,
355 },
356357/// 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.
364IndirectOcclusionCulling {
365/// The bind group for indexed meshes during the early mesh
366 /// preprocessing phase.
367early_indexed: Option<BindGroup>,
368/// The bind group for non-indexed meshes during the early mesh
369 /// preprocessing phase.
370early_non_indexed: Option<BindGroup>,
371/// The bind group for indexed meshes during the late mesh preprocessing
372 /// phase.
373late_indexed: Option<BindGroup>,
374/// The bind group for non-indexed meshes during the late mesh
375 /// preprocessing phase.
376late_non_indexed: Option<BindGroup>,
377 },
378}
379380/// 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(
387pub TypeIdHashMap<PhaseBuildIndirectParametersBindGroups>,
388);
389390impl BuildIndirectParametersBindGroups {
391/// Creates a new, empty [`BuildIndirectParametersBindGroups`] table.
392pub fn new() -> BuildIndirectParametersBindGroups {
393Self::default()
394 }
395}
396397/// 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.
402reset_indexed_indirect_batch_sets: Option<BindGroup>,
403/// The bind group for the `reset_indirect_batch_sets.wesl` shader, for
404 /// non-indexed meshes.
405reset_non_indexed_indirect_batch_sets: Option<BindGroup>,
406/// The bind group for the `build_indirect_params.wesl` shader, for indexed
407 /// meshes.
408build_indexed_indirect: Option<BindGroup>,
409/// The bind group for the `build_indirect_params.wesl` shader, for
410 /// non-indexed meshes.
411build_non_indexed_indirect: Option<BindGroup>,
412}
413414/// 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(
421pub HashMap<SceneUnpackingBuffersKey, ViewPhaseBinUnpackingBindGroups>,
422);
423424/// 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.
429indexed: Vec<ViewPhaseBinUnpackingBindGroup>,
430/// The bind groups for the non-indexed meshes, one for each batch set.
431non_indexed: Vec<ViewPhaseBinUnpackingBindGroup>,
432}
433434/// 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.
440pub metadata_index: BinUnpackingMetadataIndex,
441/// The actual shader bind group.
442pub bind_group: BindGroup,
443/// The number of mesh instances of the appropriate type (indexed or
444 /// non-indexed) for this batch set.
445pub mesh_instance_count: u32,
446}
447448/// 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(
455pub HashMap<SceneUnpackingBuffersKey, ViewPhaseUniformAllocationBindGroups>,
456);
457458/// 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.
463indexed: Vec<ViewPhaseUniformAllocationBindGroup>,
464/// The bind groups for the non-indexed meshes, one for each batch set.
465non_indexed: Vec<ViewPhaseUniformAllocationBindGroup>,
466}
467468/// 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.
474pub metadata_index: UniformAllocationMetadataIndex,
475/// The actual shader bind group.
476pub bind_group: BindGroup,
477/// The total number of bins in this batch set.
478pub bin_count: u32,
479}
480481/// 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;
485486type WithAnyPrepass = Or<(
487With<DepthPrepass>,
488With<NormalPrepass>,
489With<MotionVectorPrepass>,
490With<DeferredPrepass>,
491)>;
492493impl Pluginfor GpuMeshPreprocessPlugin {
494fn 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 }
501502fn finish(&self, app: &mut App) {
503let Some(render_app) = app.get_sub_app_mut(RenderApp) else {
504return;
505 };
506507// This plugin does nothing if GPU instance buffer building isn't in
508 // use.
509let gpu_preprocessing_support = render_app.world().resource::<GpuPreprocessingSupport>();
510if !self.use_gpu_instance_buffer_builder || !gpu_preprocessing_support.is_available() {
511return;
512 }
513514render_app515 .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(
526Render,
527 (
528clear_scene_unpacking_buffers.in_set(RenderSystems::PrepareResources),
529prepare_preprocess_pipelines.in_set(RenderSystems::Prepare),
530prepare_preprocess_bind_groups531 .run_if(resource_exists::<BatchedInstanceBuffers<
532MeshUniform,
533MeshInputUniform534 >>)
535 .in_set(RenderSystems::PrepareBindGroups)
536 .after(prepare_preprocess_pipelines),
537write_mesh_culling_data_buffer.in_set(RenderSystems::PrepareResourcesFlush),
538 ),
539 )
540 .add_systems(
541Core3d,
542 (
543 (
544allocate_uniforms,
545unpack_bins,
546early_gpu_preprocess,
547early_prepass_build_indirect_parameters.run_if(any_match_filter::<(
548With<PreprocessBindGroups>,
549Without<SkipGpuPreprocess>,
550Without<NoIndirectDrawing>,
551Or<(WithAnyPrepass, With<ShadowView>)>,
552 )>),
553 )
554 .chain()
555 .before(early_prepass),
556 (
557late_gpu_preprocess,
558late_prepass_build_indirect_parameters.run_if(any_match_filter::<(
559With<PreprocessBindGroups>,
560Without<SkipGpuPreprocess>,
561Without<NoIndirectDrawing>,
562Or<(WithAnyPrepass, With<ShadowView>)>,
563With<OcclusionCulling>,
564 )>),
565 )
566 .chain()
567 .after(early_downsample_depth)
568 .before(late_prepass),
569main_build_indirect_parameters570 .run_if(any_match_filter::<(
571With<PreprocessBindGroups>,
572Without<SkipGpuPreprocess>,
573Without<NoIndirectDrawing>,
574 )>)
575 .after(late_prepass_build_indirect_parameters)
576 .after(late_deferred_prepass)
577 .before(Core3dSystems::MainPass),
578 ),
579 );
580 }
581}
582583/// 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>,
597mut render_context: RenderContext,
598) {
599let diagnostics = render_context.diagnostic_recorder();
600let diagnostics = diagnostics.as_deref();
601602// Don't run if the shaders haven't been compiled yet.
603604let (
605Some(uniform_allocation_local_scan_pipeline_id),
606Some(uniform_allocation_global_scan_pipeline_id),
607Some(uniform_allocation_fan_pipeline_id),
608 ) = (
609preprocess_pipelines610 .uniform_allocation
611 .local_scan
612 .pipeline_id_local_scan,
613preprocess_pipelines614 .uniform_allocation
615 .global_scan
616 .pipeline_id_global_scan,
617preprocess_pipelines.uniform_allocation.fan.pipeline_id_fan,
618 )
619else {
620return;
621 };
622623let (
624Some(uniform_allocation_local_scan_pipeline),
625Some(uniform_allocation_global_scan_pipeline),
626Some(uniform_allocation_fan_pipeline),
627 ) = (
628pipeline_cache.get_compute_pipeline(uniform_allocation_local_scan_pipeline_id),
629pipeline_cache.get_compute_pipeline(uniform_allocation_global_scan_pipeline_id),
630pipeline_cache.get_compute_pipeline(uniform_allocation_fan_pipeline_id),
631 )
632else {
633return;
634 };
635636let command_encoder = render_context.command_encoder();
637let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
638 label: Some("uniform allocation"),
639 timestamp_writes: None,
640 });
641642let pass_span = diagnostics.pass_span(&mut compute_pass, "uniform_allocation");
643644// Gather up all views.
645let view_entity = current_view.entity();
646let shadow_cascade_views = current_view.into_inner();
647let all_views =
648gather_shadow_cascades_for_view(view_entity, shadow_cascade_views, &light_query);
649650// Loop over each view…
651for view_entity in all_views {
652let Ok(view) = view_query.get(view_entity) else {
653continue;
654 };
655656// …and each phase within each view.
657for phase_type_id in batched_instance_buffers.phase_instance_buffers.keys() {
658let uniform_allocation_buffers_key = SceneUnpackingBuffersKey {
659 phase: *phase_type_id,
660 view: view.retained_view_entity,
661 };
662663// Fetch the bind groups for this (view, phase) combination.
664let Some(phase_uniform_allocation_bind_groups) =
665 uniform_allocation_bind_groups.get(&uniform_allocation_buffers_key)
666else {
667continue;
668 };
669670// Invoke the shader for all batch sets corresponding to indexed
671 // meshes and then for all batch sets corresponding to
672 // non-indexed meshes.
673for 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).
679compute_pass.set_pipeline(uniform_allocation_local_scan_pipeline);
680 compute_pass.set_bind_group(0, &uniform_allocation_bind_group.bind_group, &[]);
681let local_scan_workgroup_count = uniform_allocation_bind_group
682 .bin_count
683 .div_ceil(UNIFORM_ALLOCATION_WORKGROUP_SIZE);
684if local_scan_workgroup_count > 0 {
685 compute_pass.dispatch_workgroups(local_scan_workgroup_count, 1, 1);
686 }
687688// If there are 256 or fewer draws in this batch, we're
689 // done. Otherwise, perform the other two steps.
690if local_scan_workgroup_count > 1 {
691// Invoke the global scan (step 2).
692compute_pass.set_pipeline(uniform_allocation_global_scan_pipeline);
693 compute_pass.dispatch_workgroups(1, 1, 1);
694695// Perform the fan operation (step 3).
696compute_pass.set_pipeline(uniform_allocation_fan_pipeline);
697let fan_workgroup_count = local_scan_workgroup_count - 1;
698 compute_pass.dispatch_workgroups(fan_workgroup_count, 1, 1);
699 }
700 }
701 }
702 }
703704pass_span.end(&mut compute_pass);
705}
706707/// 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>,
722mut render_context: RenderContext,
723) {
724let diagnostics = render_context.diagnostic_recorder();
725let diagnostics = diagnostics.as_deref();
726727let command_encoder = render_context.command_encoder();
728let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
729 label: Some("bin unpacking"),
730 timestamp_writes: None,
731 });
732733let pass_span = diagnostics.pass_span(&mut compute_pass, "bin_unpacking");
734735// Gather up all views.
736let view_entity = current_view.entity();
737let shadow_cascade_views = current_view.into_inner();
738let all_views =
739gather_shadow_cascades_for_view(view_entity, shadow_cascade_views, &light_query);
740741// Don't run if the shaders haven't been compiled yet.
742if let Some(bin_unpacking_pipeline_id) = preprocess_pipelines.bin_unpacking.pipeline_id
743 && let Some(bin_unpacking_pipeline) =
744pipeline_cache.get_compute_pipeline(bin_unpacking_pipeline_id)
745 {
746compute_pass.set_pipeline(bin_unpacking_pipeline);
747748// Loop over each view…
749for view_entity in all_views {
750let Ok(view) = view_query.get(view_entity) else {
751continue;
752 };
753754// …and each phase within each view.
755for phase_type_id in batched_instance_buffers.phase_instance_buffers.keys() {
756let scene_unpacking_buffers_key = SceneUnpackingBuffersKey {
757 phase: *phase_type_id,
758 view: view.retained_view_entity,
759 };
760761// Fetch the bind groups for this (view, phase) combination.
762let Some(phase_bin_unpacking_bind_groups) =
763 bin_unpacking_bind_groups.get(&scene_unpacking_buffers_key)
764else {
765continue;
766 };
767768// Invoke the shader for all batch sets corresponding to indexed
769 // meshes and then for all batch sets corresponding to
770 // non-indexed meshes.
771for 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, &[]);
777let workgroup_count = (bin_unpacking_bind_group.mesh_instance_count as usize)
778 .div_ceil(WORKGROUP_SIZE);
779if workgroup_count > 0 {
780 compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
781 }
782 }
783 }
784 }
785 }
786787pass_span.end(&mut compute_pass);
788}
789790pub fn early_gpu_preprocess(
791 current_view: ViewQuery<Option<&ViewLightEntities>, Without<SkipGpuPreprocess>>,
792 view_query: Query<
793 (
794&ExtractedView,
795Option<&PreprocessBindGroups>,
796Option<&ViewUniformOffset>,
797Has<NoIndirectDrawing>,
798Has<OcclusionCulling>,
799 ),
800Without<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>,
806mut ctx: RenderContext,
807) {
808let diagnostics = ctx.diagnostic_recorder();
809let diagnostics = diagnostics.as_deref();
810811let command_encoder = ctx.command_encoder();
812813let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
814 label: Some("early_mesh_preprocessing"),
815 timestamp_writes: None,
816 });
817818let pass_span = diagnostics.pass_span(&mut compute_pass, "early_mesh_preprocessing");
819820let view_entity = current_view.entity();
821let shadow_cascade_views = current_view.into_inner();
822let all_views =
823gather_shadow_cascades_for_view(view_entity, shadow_cascade_views, &light_query);
824825// Run the compute passes.
826for view_entity in all_views {
827let Ok((view, bind_groups, view_uniform_offset, no_indirect_drawing, occlusion_culling)) =
828 view_query.get(view_entity)
829else {
830continue;
831 };
832833let Some(bind_groups) = bind_groups else {
834continue;
835 };
836let Some(view_uniform_offset) = view_uniform_offset else {
837continue;
838 };
839840// Select the right pipeline, depending on whether GPU culling is in
841 // use.
842let 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 };
853854// Fetch the pipeline.
855let 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");
857continue;
858 };
859860let Some(preprocess_pipeline) = pipeline_cache.get_compute_pipeline(preprocess_pipeline_id)
861else {
862// This will happen while the pipeline is being compiled and is fine.
863continue;
864 };
865866 compute_pass.set_pipeline(preprocess_pipeline);
867868// Loop over each render phase.
869for (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.
873let Some(work_item_buffers) = batched_phase_instance_buffers
874 .work_item_buffers
875 .get(&view.retained_view_entity)
876else {
877continue;
878 };
879880// Fetch the bind group for the render phase.
881let Some(phase_bind_groups) = bind_groups.get(phase_type_id) else {
882continue;
883 };
884885// Make sure the mesh preprocessing shader has access to the
886 // view info it needs to do culling and motion vector
887 // computation.
888let dynamic_offsets = [view_uniform_offset.offset];
889890// Are we drawing directly or indirectly?
891match *phase_bind_groups {
892 PhasePreprocessBindGroups::Direct(ref bind_group) => {
893// Invoke the mesh preprocessing shader to transform
894 // meshes only, but not cull.
895let PreprocessWorkItemBuffers::Direct(work_item_buffer) = work_item_buffers
896else {
897continue;
898 };
899 compute_pass.set_bind_group(0, bind_group, &dynamic_offsets);
900let workgroup_count = work_item_buffer.len().div_ceil(WORKGROUP_SIZE);
901if workgroup_count > 0 {
902 compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
903 }
904 }
905906 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.
917let PreprocessWorkItemBuffers::Indirect {
918 indexed: indexed_buffer,
919 non_indexed: non_indexed_buffer,
920 ..
921 } = work_item_buffers
922else {
923continue;
924 };
925926// Transform and cull indexed meshes if there are any.
927if let Some(indexed_bind_group) = maybe_indexed_bind_group {
928if let PreprocessWorkItemBuffers::Indirect {
929 gpu_occlusion_culling:
930Some(GpuOcclusionCullingWorkItemBuffers {
931 late_indirect_parameters_indexed_offset,
932 ..
933 }),
934 ..
935 } = *work_item_buffers
936 {
937 compute_pass.set_immediates(
9380,
939 bytemuck::bytes_of(&late_indirect_parameters_indexed_offset),
940 );
941 }
942943 compute_pass.set_bind_group(0, indexed_bind_group, &dynamic_offsets);
944let workgroup_count = indexed_buffer.len().div_ceil(WORKGROUP_SIZE);
945if workgroup_count > 0 {
946 compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
947 }
948 }
949950// Transform and cull non-indexed meshes if there are any.
951if let Some(non_indexed_bind_group) = maybe_non_indexed_bind_group {
952if let PreprocessWorkItemBuffers::Indirect {
953 gpu_occlusion_culling:
954Some(GpuOcclusionCullingWorkItemBuffers {
955 late_indirect_parameters_non_indexed_offset,
956 ..
957 }),
958 ..
959 } = *work_item_buffers
960 {
961 compute_pass.set_immediates(
9620,
963 bytemuck::bytes_of(&late_indirect_parameters_non_indexed_offset),
964 );
965 }
966967 compute_pass.set_bind_group(0, non_indexed_bind_group, &dynamic_offsets);
968let workgroup_count = non_indexed_buffer.len().div_ceil(WORKGROUP_SIZE);
969if workgroup_count > 0 {
970 compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
971 }
972 }
973 }
974 }
975 }
976 }
977978pass_span.end(&mut compute_pass);
979}
980981/// 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]> {
988let mut all_views: SmallVec<[_; 8]> = SmallVec::new();
989all_views.push(view_entity);
990if let Some(shadow_cascade_views) = shadow_cascade_views {
991all_views.extend(
992shadow_cascade_views993 .lights
994 .iter()
995 .filter(|light_entity| {
996light_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 }
1003all_views1004}
10051006pub fn late_gpu_preprocess(
1007 current_view: ViewQuery<
1008 (&ExtractedView, &PreprocessBindGroups, &ViewUniformOffset),
1009 (
1010Without<SkipGpuPreprocess>,
1011Without<NoIndirectDrawing>,
1012With<OcclusionCulling>,
1013With<DepthPrepass>,
1014 ),
1015 >,
1016 batched_instance_buffers: Res<BatchedInstanceBuffers<MeshUniform, MeshInputUniform>>,
1017 pipeline_cache: Res<PipelineCache>,
1018 preprocess_pipelines: Res<PreprocessPipelines>,
1019mut ctx: RenderContext,
1020) {
1021let (view, bind_groups, view_uniform_offset) = current_view.into_inner();
10221023// Fetch the pipeline BEFORE starting diagnostic spans to avoid panic on early return
1024let maybe_pipeline_id = preprocess_pipelines1025 .late_gpu_occlusion_culling_preprocess
1026 .pipeline_id;
10271028let Some(preprocess_pipeline_id) = maybe_pipeline_idelse {
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");
1030return;
1031 };
10321033let Some(preprocess_pipeline) = pipeline_cache.get_compute_pipeline(preprocess_pipeline_id)
1034else {
1035// This will happen while the pipeline is being compiled and is fine.
1036return;
1037 };
10381039let diagnostics = ctx.diagnostic_recorder();
1040let diagnostics = diagnostics.as_deref();
10411042let command_encoder = ctx.command_encoder();
10431044let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
1045 label: Some("late_mesh_preprocessing"),
1046 timestamp_writes: None,
1047 });
10481049let pass_span = diagnostics.pass_span(&mut compute_pass, "late_mesh_preprocessing");
10501051compute_pass.set_pipeline(preprocess_pipeline);
10521053// Loop over each phase. Because we built the phases in parallel,
1054 // each phase has a separate set of instance buffers.
1055for (phase_type_id, batched_phase_instance_buffers) in
1056&batched_instance_buffers.phase_instance_buffers
1057 {
1058let UntypedPhaseBatchedInstanceBuffers {
1059ref work_item_buffers,
1060ref late_indexed_indirect_parameters_buffer,
1061ref late_non_indexed_indirect_parameters_buffer,
1062 ..
1063 } = *batched_phase_instance_buffers;
10641065// Grab the work item buffers for this view.
1066let Some(phase_work_item_buffers) = work_item_buffers.get(&view.retained_view_entity)
1067else {
1068continue;
1069 };
10701071let (
1072 PreprocessWorkItemBuffers::Indirect {
1073 gpu_occlusion_culling:
1074Some(GpuOcclusionCullingWorkItemBuffers {
1075 late_indirect_parameters_indexed_offset,
1076 late_indirect_parameters_non_indexed_offset,
1077 ..
1078 }),
1079 ..
1080 },
1081Some(PhasePreprocessBindGroups::IndirectOcclusionCulling {
1082 late_indexed: maybe_late_indexed_bind_group,
1083 late_non_indexed: maybe_late_non_indexed_bind_group,
1084 ..
1085 }),
1086Some(late_indexed_indirect_parameters_buffer),
1087Some(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 )
1094else {
1095continue;
1096 };
10971098let mut dynamic_offsets: SmallVec<[u32; 1]> = ::smallvec::SmallVec::new()smallvec![];
1099 dynamic_offsets.push(view_uniform_offset.offset);
11001101// 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.
11051106 // Transform and cull indexed meshes if there are any.
1107if let Some(late_indexed_bind_group) = maybe_late_indexed_bind_group {
1108 compute_pass.set_immediates(
11090,
1110 bytemuck::bytes_of(late_indirect_parameters_indexed_offset),
1111 );
11121113 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 }
11201121// Transform and cull non-indexed meshes if there are any.
1122if let Some(late_non_indexed_bind_group) = maybe_late_non_indexed_bind_group {
1123 compute_pass.set_immediates(
11240,
1125 bytemuck::bytes_of(late_indirect_parameters_non_indexed_offset),
1126 );
11271128 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 }
11361137pass_span.end(&mut compute_pass);
1138}
11391140/// 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>,
1152mut ctx: RenderContext,
1153) {
1154run_build_indirect_parameters(
1155&mut ctx,
1156current_view.into_inner().retained_view_entity,
1157build_indirect_params_bind_groups.as_deref(),
1158&pipeline_cache,
1159indirect_parameters_buffers.as_deref(),
1160&build_indirect_parameters_uniform_indices,
1161&preprocess_pipelines.early_phase,
1162"early_prepass_indirect_parameters_building",
1163 );
1164}
11651166/// 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>,
1179mut ctx: RenderContext,
1180) {
1181run_build_indirect_parameters(
1182&mut ctx,
1183current_view.into_inner().retained_view_entity,
1184build_indirect_params_bind_groups.as_deref(),
1185&pipeline_cache,
1186indirect_parameters_buffers.as_deref(),
1187&build_indirect_parameters_uniform_indices,
1188&preprocess_pipelines.late_phase,
1189"late_prepass_indirect_parameters_building",
1190 );
1191}
11921193/// 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>,
1203mut ctx: RenderContext,
1204) {
1205run_build_indirect_parameters(
1206&mut ctx,
1207current_view.into_inner().retained_view_entity,
1208build_indirect_params_bind_groups.as_deref(),
1209&pipeline_cache,
1210indirect_parameters_buffers.as_deref(),
1211&build_indirect_parameters_uniform_indices,
1212&preprocess_pipelines.main_phase,
1213"main_indirect_parameters_building",
1214 );
1215}
12161217/// 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) {
1229let Some(build_indirect_params_bind_groups) = build_indirect_params_bind_groupselse {
1230return;
1231 };
1232let Some(indirect_parameters_buffers) = indirect_parameters_bufferselse {
1233return;
1234 };
1235let Some(view_build_indirect_parameters_uniform_indices) =
1236build_indirect_parameters_uniform_indices.get(&retained_view_entity)
1237else {
1238return;
1239 };
12401241let command_encoder = ctx.command_encoder();
12421243let mut compute_pass = command_encoder.begin_compute_pass(&ComputePassDescriptor {
1244 label: Some(label),
1245 timestamp_writes: None,
1246 });
12471248// Fetch the pipeline.
1249let (
1250Some(reset_indirect_batch_sets_pipeline_id),
1251Some(build_indexed_indirect_params_pipeline_id),
1252Some(build_non_indexed_indirect_params_pipeline_id),
1253 ) = (
1254preprocess_phase_pipelines1255 .reset_indirect_batch_sets
1256 .pipeline_id,
1257preprocess_phase_pipelines1258 .gpu_occlusion_culling_build_indexed_indirect_params
1259 .pipeline_id,
1260preprocess_phase_pipelines1261 .gpu_occlusion_culling_build_non_indexed_indirect_params
1262 .pipeline_id,
1263 )
1264else {
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");
1266return;
1267 };
12681269let (
1270Some(reset_indirect_batch_sets_pipeline),
1271Some(build_indexed_indirect_params_pipeline),
1272Some(build_non_indexed_indirect_params_pipeline),
1273 ) = (
1274pipeline_cache.get_compute_pipeline(reset_indirect_batch_sets_pipeline_id),
1275pipeline_cache.get_compute_pipeline(build_indexed_indirect_params_pipeline_id),
1276pipeline_cache.get_compute_pipeline(build_non_indexed_indirect_params_pipeline_id),
1277 )
1278else {
1279// This will happen while the pipeline is being compiled and is fine.
1280return;
1281 };
12821283// Loop over each phase. As each has as separate set of buffers, we need to
1284 // build indirect parameters individually for each phase.
1285for (phase_type_id, phase_build_indirect_params_bind_groups) in
1286build_indirect_params_bind_groups.iter()
1287 {
1288let Some(phase_indirect_parameters_buffers) =
1289 indirect_parameters_buffers.get(phase_type_id)
1290else {
1291continue;
1292 };
1293let Some(build_indirect_parameters_uniform_index) =
1294 view_build_indirect_parameters_uniform_indices.get(phase_type_id)
1295else {
1296continue;
1297 };
12981299// Build indexed indirect parameters.
1300if let (
1301Some(reset_indexed_indirect_batch_sets_bind_group),
1302Some(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, &[]);
1309let workgroup_count = phase_indirect_parameters_buffers
1310 .batch_set_count(true)
1311 .div_ceil(WORKGROUP_SIZE);
1312if workgroup_count > 0 {
1313 compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
1314 }
13151316 compute_pass.set_pipeline(build_indexed_indirect_params_pipeline);
13171318for indexed_build_indirect_parameters_metadata in
1319&build_indirect_parameters_uniform_index.indexed
1320 {
1321 compute_pass.set_bind_group(
13220,
1323 build_indirect_indexed_params_bind_group,
1324&[indexed_build_indirect_parameters_metadata.uniform_offset],
1325 );
1326let workgroup_count = indexed_build_indirect_parameters_metadata
1327 .batch_count
1328 .div_ceil(WORKGROUP_SIZE as u32);
1329if workgroup_count > 0 {
1330 compute_pass.dispatch_workgroups(workgroup_count, 1, 1);
1331 }
1332 }
1333 }
13341335// Build non-indexed indirect parameters.
1336if let (
1337Some(reset_non_indexed_indirect_batch_sets_bind_group),
1338Some(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, &[]);
1345let workgroup_count = phase_indirect_parameters_buffers
1346 .batch_set_count(false)
1347 .div_ceil(WORKGROUP_SIZE);
1348if workgroup_count > 0 {
1349 compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
1350 }
13511352 compute_pass.set_pipeline(build_non_indexed_indirect_params_pipeline);
13531354for non_indexed_build_indirect_parameters_metadata in
1355&build_indirect_parameters_uniform_index.non_indexed
1356 {
1357 compute_pass.set_bind_group(
13580,
1359 build_indirect_non_indexed_params_bind_group,
1360&[non_indexed_build_indirect_parameters_metadata.uniform_offset],
1361 );
1362let workgroup_count = non_indexed_build_indirect_parameters_metadata
1363 .batch_count
1364 .div_ceil(WORKGROUP_SIZE as u32);
1365if workgroup_count > 0 {
1366 compute_pass.dispatch_workgroups(workgroup_count, 1, 1);
1367 }
1368 }
1369 }
1370 }
1371}
13721373impl PreprocessPipelines {
1374/// Returns true if the preprocessing and indirect parameters pipelines have
1375 /// been loaded or false otherwise.
1376pub(crate) fn pipelines_are_loaded(
1377&self,
1378 pipeline_cache: &PipelineCache,
1379 preprocessing_support: &GpuPreprocessingSupport,
1380 ) -> bool {
1381match preprocessing_support.max_supported_mode {
1382 GpuPreprocessingMode::None => false,
1383 GpuPreprocessingMode::PreprocessingOnly => {
1384self.direct_preprocess.is_loaded(pipeline_cache)
1385 && self1386 .gpu_frustum_culling_preprocess
1387 .is_loaded(pipeline_cache)
1388 }
1389 GpuPreprocessingMode::Culling => {
1390self.direct_preprocess.is_loaded(pipeline_cache)
1391 && self1392 .gpu_frustum_culling_preprocess
1393 .is_loaded(pipeline_cache)
1394 && self1395 .early_gpu_occlusion_culling_preprocess
1396 .is_loaded(pipeline_cache)
1397 && self1398 .late_gpu_occlusion_culling_preprocess
1399 .is_loaded(pipeline_cache)
1400 && self1401 .gpu_frustum_culling_build_indexed_indirect_params
1402 .is_loaded(pipeline_cache)
1403 && self1404 .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}
14131414impl PreprocessPhasePipelines {
1415fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1416self.reset_indirect_batch_sets.is_loaded(pipeline_cache)
1417 && self1418 .gpu_occlusion_culling_build_indexed_indirect_params
1419 .is_loaded(pipeline_cache)
1420 && self1421 .gpu_occlusion_culling_build_non_indexed_indirect_params
1422 .is_loaded(pipeline_cache)
1423 }
1424}
14251426impl PreprocessPipeline {
1427fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1428self.pipeline_id
1429 .is_some_and(|pipeline_id| pipeline_cache.get_compute_pipeline(pipeline_id).is_some())
1430 }
1431}
14321433impl ResetIndirectBatchSetsPipeline {
1434fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1435self.pipeline_id
1436 .is_some_and(|pipeline_id| pipeline_cache.get_compute_pipeline(pipeline_id).is_some())
1437 }
1438}
14391440impl BuildIndirectParametersPipeline {
1441/// Returns true if this pipeline has been loaded into the pipeline cache or
1442 /// false otherwise.
1443fn is_loaded(&self, pipeline_cache: &PipelineCache) -> bool {
1444self.pipeline_id
1445 .is_some_and(|pipeline_id| pipeline_cache.get_compute_pipeline(pipeline_id).is_some())
1446 }
1447}
14481449impl SpecializedComputePipelinefor PreprocessPipeline {
1450type Key = PreprocessPipelineKey;
14511452fn specialize(&self, key: Self::Key) -> ComputePipelineDescriptor {
1453let 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()];
1454if key.contains(PreprocessPipelineKey::FRUSTUM_CULLING) {
1455shader_defs.push("INDIRECT".into());
1456shader_defs.push("FRUSTUM_CULLING".into());
1457 }
1458if key.contains(PreprocessPipelineKey::OCCLUSION_CULLING) {
1459shader_defs.push("OCCLUSION_CULLING".into());
1460if key.contains(PreprocessPipelineKey::EARLY_PHASE) {
1461shader_defs.push("EARLY_PHASE".into());
1462 } else {
1463shader_defs.push("LATE_PHASE".into());
1464 }
1465 }
14661467ComputePipelineDescriptor {
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 ({})",
1471if 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) {
14884
1489} else {
14900
1491},
1492 shader: self.shader.clone(),
1493shader_defs,
1494 ..default()
1495 }
1496 }
1497}
14981499impl FromWorldfor PreprocessPipelines {
1500fn from_world(world: &mut World) -> Self {
1501// GPU culling bind group parameters are a superset of those in the CPU
1502 // culling (direct) shader.
1503let direct_bind_group_layout_entries = preprocess_direct_bind_group_layout_entries();
1504let gpu_frustum_culling_bind_group_layout_entries = gpu_culling_bind_group_layout_entries();
1505let gpu_early_occlusion_culling_bind_group_layout_entries =
1506gpu_occlusion_culling_bind_group_layout_entries().extend_with_indices((
1507 (
150812,
1509storage_buffer::<PreprocessWorkItem>(/*has_dynamic_offset=*/ false),
1510 ),
1511 (
151213,
1513storage_buffer::<LatePreprocessWorkItemIndirectParameters>(
1514/*has_dynamic_offset=*/ false,
1515 ),
1516 ),
1517 ));
1518let gpu_late_occlusion_culling_bind_group_layout_entries =
1519gpu_occlusion_culling_bind_group_layout_entries().extend_with_indices(((
152013,
1521storage_buffer_read_only::<LatePreprocessWorkItemIndirectParameters>(
1522/*has_dynamic_offset=*/ false,
1523 ),
1524 ),));
15251526let reset_indirect_batch_sets_bind_group_layout_entries =
1527DynamicBindGroupLayoutEntries::sequential(
1528ShaderStages::COMPUTE,
1529 (storage_buffer::<IndirectBatchSet>(false),),
1530 );
15311532// Indexed and non-indexed bind group parameters share all the bind
1533 // group layout entries except the final one.
1534let build_indexed_indirect_params_bind_group_layout_entries =
1535build_indirect_params_bind_group_layout_entries()
1536 .extend_sequential((storage_buffer::<IndirectParametersIndexed>(false),));
1537let build_non_indexed_indirect_params_bind_group_layout_entries =
1538build_indirect_params_bind_group_layout_entries()
1539 .extend_sequential((storage_buffer::<IndirectParametersNonIndexed>(false),));
15401541let bin_unpacking_bind_group_layout_entries = bin_unpacking_bind_group_layout_entries();
1542let uniform_allocation_bind_group_layout_entries =
1543uniform_allocation_bind_group_layout_entries();
15441545// Create the bind group layouts.
1546let direct_bind_group_layout = BindGroupLayoutDescriptor::new(
1547"build mesh uniforms direct bind group layout",
1548&direct_bind_group_layout_entries,
1549 );
1550let 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 );
1554let 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 );
1558let 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 );
1562let 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 );
1566let 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 );
1570let 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 );
1574let bin_unpacking_bind_group_layout = BindGroupLayoutDescriptor::new(
1575"bin unpacking bind group layout",
1576&bin_unpacking_bind_group_layout_entries,
1577 );
1578let uniform_allocation_bind_group_layout = BindGroupLayoutDescriptor::new(
1579"uniform allocation bind group layout",
1580&uniform_allocation_bind_group_layout_entries,
1581 );
15821583let 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");
1584let 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");
1586let 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");
1588let 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");
1589let 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");
15901591let 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:
1603BuildIndirectParametersPipeline {
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 };
16091610PreprocessPipelines {
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:
1637BuildIndirectParametersPipeline {
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}
16701671fn preprocess_direct_bind_group_layout_entries() -> DynamicBindGroupLayoutEntries {
1672DynamicBindGroupLayoutEntries::new_with_indices(
1673ShaderStages::COMPUTE,
1674 (
1675// `view`
1676(
16770,
1678uniform_buffer::<ViewUniform>(/* has_dynamic_offset= */ true),
1679 ),
1680// `current_input`
1681(3, storage_buffer_read_only::<MeshInputUniform>(false)),
1682// `previous_input`
1683(
16844,
1685storage_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}
16941695// 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 {
1698DynamicBindGroupLayoutEntries::new_with_indices(
1699ShaderStages::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(
17071,
1708storage_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}
17191720/// 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.
1725preprocess_direct_bind_group_layout_entries().extend_with_indices((
1726// @group(0) @binding(7) var<storage> indirect_parameters_metadata:
1727 // array<IndirectParametersMetadata>;
1728(
17297,
1730storage_buffer::<IndirectParametersMetadata>(/* has_dynamic_offset= */ false),
1731 ),
1732// `mesh_culling_data`
1733(
17349,
1735storage_buffer_read_only::<MeshCullingData>(/* has_dynamic_offset= */ false),
1736 ),
1737// `visibility_ranges`
1738(
173910,
1740storage_buffer_read_only::<Vec4>(/* has_dynamic_offset= */ false),
1741 ),
1742 ))
1743}
17441745fn gpu_occlusion_culling_bind_group_layout_entries() -> DynamicBindGroupLayoutEntries {
1746gpu_culling_bind_group_layout_entries().extend_with_indices((
1747 (
17482,
1749uniform_buffer::<PreviousViewData>(/*has_dynamic_offset=*/ false),
1750 ),
1751 (
175211,
1753texture_2d(TextureSampleType::Float { filterable: true }),
1754 ),
1755 ))
1756}
17571758/// 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> {
1761BindGroupLayoutEntries::sequential(
1762ShaderStages::COMPUTE,
1763 (
1764// @group(0) @binding(0) var<uniform> bin_unpacking_metadata:
1765 // BinUnpackingMetadata;
1766uniform_buffer::<GpuBinUnpackingMetadata>(false),
1767// @group(0) @binding(1) var<storage> binned_mesh_instances:
1768 // array<BinnedMeshInstance>;
1769storage_buffer_read_only::<GpuRenderBinnedMeshInstance>(false),
1770// @group(0) @binding(2) var<storage, read_write>
1771 // preprocess_work_items: array<PreprocessWorkItem>;
1772storage_buffer::<PreprocessWorkItem>(false),
1773// @group(0) @binding(3) var<storage> bin_metadata:
1774 // array<GpuBinMetadata>;
1775storage_buffer_read_only::<GpuBinMetadata>(false),
1776// @group(0) @binding(4) var<storage>
1777 // bin_index_to_bin_metadata_index: array<u32>;
1778storage_buffer_read_only::<u32>(false),
1779 ),
1780 )
1781}
17821783/// 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> {
1786BindGroupLayoutEntries::sequential(
1787ShaderStages::COMPUTE,
1788 (
1789// @group(0) @binding(0) var<uniform> allocate_uniforms_metadata:
1790 // AllocateUniformsMetadata;
1791uniform_buffer::<GpuUniformAllocationMetadata>(false),
1792// @group(0) @binding(1) var<storage> bin_metadata: array<BinMetadata>;
1793storage_buffer_read_only::<GpuBinMetadata>(false),
1794// @group(0) @binding(2) var<storage, read_write>
1795 // indirect_parameters_metadata: array<IndirectParametersMetadata>;
1796storage_buffer::<IndirectParametersMetadata>(false),
1797// @group(0) @binding(3) var<storage, read_write> fan_buffer:
1798 // array<u32>;
1799storage_buffer::<u32>(false),
1800 ),
1801 )
1802}
18031804/// 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>,
1814mut specialized_preprocess_pipelines: ResMut<SpecializedComputePipelines<PreprocessPipeline>>,
1815mut specialized_reset_indirect_batch_sets_pipelines: ResMut<
1816SpecializedComputePipelines<ResetIndirectBatchSetsPipeline>,
1817 >,
1818mut specialized_build_indirect_parameters_pipelines: ResMut<
1819SpecializedComputePipelines<BuildIndirectParametersPipeline>,
1820 >,
1821mut specialized_bin_unpacking_pipelines: ResMut<
1822SpecializedComputePipelines<BinUnpackingPipeline>,
1823 >,
1824mut specialized_uniform_allocation_local_scan_pipelines: ResMut<
1825SpecializedComputePipelines<UniformAllocationLocalScanPipeline>,
1826 >,
1827mut specialized_uniform_allocation_global_scan_pipelines: ResMut<
1828SpecializedComputePipelines<UniformAllocationGlobalScanPipeline>,
1829 >,
1830mut specialized_uniform_allocation_fan_pipelines: ResMut<
1831SpecializedComputePipelines<UniformAllocationFanPipeline>,
1832 >,
1833 preprocess_pipelines: ResMut<PreprocessPipelines>,
1834 gpu_preprocessing_support: Res<GpuPreprocessingSupport>,
1835) {
1836let preprocess_pipelines = preprocess_pipelines.into_inner();
18371838preprocess_pipelines.direct_preprocess.prepare(
1839&pipeline_cache,
1840&mut specialized_preprocess_pipelines,
1841PreprocessPipelineKey::empty(),
1842 );
1843preprocess_pipelines.gpu_frustum_culling_preprocess.prepare(
1844&pipeline_cache,
1845&mut specialized_preprocess_pipelines,
1846PreprocessPipelineKey::FRUSTUM_CULLING,
1847 );
18481849if gpu_preprocessing_support.is_culling_supported() {
1850preprocess_pipelines1851 .early_gpu_occlusion_culling_preprocess
1852 .prepare(
1853&pipeline_cache,
1854&mut specialized_preprocess_pipelines,
1855PreprocessPipelineKey::FRUSTUM_CULLING1856 | PreprocessPipelineKey::OCCLUSION_CULLING1857 | PreprocessPipelineKey::EARLY_PHASE,
1858 );
1859preprocess_pipelines1860 .late_gpu_occlusion_culling_preprocess
1861 .prepare(
1862&pipeline_cache,
1863&mut specialized_preprocess_pipelines,
1864PreprocessPipelineKey::FRUSTUM_CULLING | PreprocessPipelineKey::OCCLUSION_CULLING,
1865 );
1866 }
18671868let mut build_indirect_parameters_pipeline_key = BuildIndirectParametersPipelineKey::empty();
18691870// If the GPU and driver support `multi_draw_indirect_count`, tell the
1871 // shader that.
1872if render_device1873 .wgpu_device()
1874 .features()
1875 .contains(WgpuFeatures::MULTI_DRAW_INDIRECT_COUNT)
1876 {
1877build_indirect_parameters_pipeline_key1878 .insert(BuildIndirectParametersPipelineKey::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED);
1879 }
18801881preprocess_pipelines1882 .gpu_frustum_culling_build_indexed_indirect_params
1883 .prepare(
1884&pipeline_cache,
1885&mut specialized_build_indirect_parameters_pipelines,
1886build_indirect_parameters_pipeline_key | BuildIndirectParametersPipelineKey::INDEXED,
1887 );
1888preprocess_pipelines1889 .gpu_frustum_culling_build_non_indexed_indirect_params
1890 .prepare(
1891&pipeline_cache,
1892&mut specialized_build_indirect_parameters_pipelines,
1893build_indirect_parameters_pipeline_key,
1894 );
18951896if !gpu_preprocessing_support.is_culling_supported() {
1897return;
1898 }
18991900for (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 }
19401941// Prepare the bin unpacking compute pipeline.
1942preprocess_pipelines1943 .bin_unpacking
1944 .prepare(&pipeline_cache, &mut specialized_bin_unpacking_pipelines);
19451946// Prepare the uniform allocation compute pipeline.
1947preprocess_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}
19541955impl PreprocessPipeline {
1956fn prepare(
1957&mut self,
1958 pipeline_cache: &PipelineCache,
1959 pipelines: &mut SpecializedComputePipelines<PreprocessPipeline>,
1960 key: PreprocessPipelineKey,
1961 ) {
1962if self.pipeline_id.is_some() {
1963return;
1964 }
19651966let preprocess_pipeline_id = pipelines.specialize(pipeline_cache, self, key);
1967self.pipeline_id = Some(preprocess_pipeline_id);
1968 }
1969}
19701971impl SpecializedComputePipelinefor ResetIndirectBatchSetsPipeline {
1972type Key = ();
19731974fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
1975ComputePipelineDescriptor {
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}
19831984impl SpecializedComputePipelinefor BuildIndirectParametersPipeline {
1985type Key = BuildIndirectParametersPipelineKey;
19861987fn specialize(&self, key: Self::Key) -> ComputePipelineDescriptor {
1988let mut shader_defs = ::alloc::vec::Vec::new()vec![];
1989if key.contains(BuildIndirectParametersPipelineKey::INDEXED) {
1990shader_defs.push("INDEXED".into());
1991 }
1992if key.contains(BuildIndirectParametersPipelineKey::MULTI_DRAW_INDIRECT_COUNT_SUPPORTED) {
1993shader_defs.push("MULTI_DRAW_INDIRECT_COUNT_SUPPORTED".into());
1994 }
1995if key.contains(BuildIndirectParametersPipelineKey::OCCLUSION_CULLING) {
1996shader_defs.push("OCCLUSION_CULLING".into());
1997 }
1998if key.contains(BuildIndirectParametersPipelineKey::EARLY_PHASE) {
1999shader_defs.push("EARLY_PHASE".into());
2000 }
2001if key.contains(BuildIndirectParametersPipelineKey::LATE_PHASE) {
2002shader_defs.push("LATE_PHASE".into());
2003 }
2004if key.contains(BuildIndirectParametersPipelineKey::MAIN_PHASE) {
2005shader_defs.push("MAIN_PHASE".into());
2006 }
20072008let 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",
2010if !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},
2019if key.contains(BuildIndirectParametersPipelineKey::INDEXED) {
2020""
2021} else {
2022"non-"
2023}
2024 );
20252026ComputePipelineDescriptor {
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(),
2030shader_defs,
2031 ..default()
2032 }
2033 }
2034}
20352036impl SpecializedComputePipelinefor BinUnpackingPipeline {
2037type Key = ();
20382039fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2040ComputePipelineDescriptor {
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}
20492050impl SpecializedComputePipelinefor UniformAllocationLocalScanPipeline {
2051type Key = ();
20522053fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2054ComputePipelineDescriptor {
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}
20642065impl SpecializedComputePipelinefor UniformAllocationGlobalScanPipeline {
2066type Key = ();
20672068fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2069ComputePipelineDescriptor {
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}
20792080impl SpecializedComputePipelinefor UniformAllocationFanPipeline {
2081type Key = ();
20822083fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
2084ComputePipelineDescriptor {
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}
20942095impl ResetIndirectBatchSetsPipeline {
2096fn prepare(
2097&mut self,
2098 pipeline_cache: &PipelineCache,
2099 pipelines: &mut SpecializedComputePipelines<ResetIndirectBatchSetsPipeline>,
2100 ) {
2101if self.pipeline_id.is_some() {
2102return;
2103 }
21042105let reset_indirect_batch_sets_pipeline_id = pipelines.specialize(pipeline_cache, self, ());
2106self.pipeline_id = Some(reset_indirect_batch_sets_pipeline_id);
2107 }
2108}
21092110impl BuildIndirectParametersPipeline {
2111fn prepare(
2112&mut self,
2113 pipeline_cache: &PipelineCache,
2114 pipelines: &mut SpecializedComputePipelines<BuildIndirectParametersPipeline>,
2115 key: BuildIndirectParametersPipelineKey,
2116 ) {
2117if self.pipeline_id.is_some() {
2118return;
2119 }
21202121let build_indirect_parameters_pipeline_id = pipelines.specialize(pipeline_cache, self, key);
2122self.pipeline_id = Some(build_indirect_parameters_pipeline_id);
2123 }
2124}
21252126impl BinUnpackingPipeline {
2127/// Specializes a single pipeline for the bin unpacking shader.
2128fn prepare(
2129&mut self,
2130 pipeline_cache: &PipelineCache,
2131 pipelines: &mut SpecializedComputePipelines<BinUnpackingPipeline>,
2132 ) {
2133if self.pipeline_id.is_some() {
2134return;
2135 }
21362137let bin_unpacking_pipeline_id = pipelines.specialize(pipeline_cache, self, ());
2138self.pipeline_id = Some(bin_unpacking_pipeline_id);
2139 }
2140}
21412142impl UniformAllocationPipelines {
2143/// Specializes all three pipelines that use the uniform allocation shader.
2144fn prepare(
2145&mut self,
2146 pipeline_cache: &PipelineCache,
2147 uniform_allocation_local_scan_pipelines: &mut SpecializedComputePipelines<
2148UniformAllocationLocalScanPipeline,
2149 >,
2150 uniform_allocation_global_scan_pipelines: &mut SpecializedComputePipelines<
2151UniformAllocationGlobalScanPipeline,
2152 >,
2153 uniform_allocation_fan_pipelines: &mut SpecializedComputePipelines<
2154UniformAllocationFanPipeline,
2155 >,
2156 ) {
2157if self.local_scan.pipeline_id_local_scan.is_none() {
2158self.local_scan.pipeline_id_local_scan =
2159Some(uniform_allocation_local_scan_pipelines.specialize(
2160pipeline_cache,
2161&self.local_scan,
2162 (),
2163 ));
2164 }
21652166if self.global_scan.pipeline_id_global_scan.is_none() {
2167self.global_scan.pipeline_id_global_scan =
2168Some(uniform_allocation_global_scan_pipelines.specialize(
2169pipeline_cache,
2170&self.global_scan,
2171 (),
2172 ));
2173 }
21742175if self.fan.pipeline_id_fan.is_none() {
2176self.fan.pipeline_id_fan =
2177Some(uniform_allocation_fan_pipelines.specialize(pipeline_cache, &self.fan, ()));
2178 }
2179 }
2180}
21812182/// 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(
2189mut 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>,
2203mut bin_unpacking_bind_groups: ResMut<BinUnpackingBindGroups>,
2204mut uniform_allocation_bind_groups: ResMut<UniformAllocationBindGroups>,
2205) {
2206// Grab the `BatchedInstanceBuffers`.
2207let 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();
22122213let (Some(current_input_buffer), Some(previous_input_buffer)) = (
2214current_input_buffer_vec.buffer().buffer(),
2215previous_input_buffer_vec.buffer(),
2216 ) else {
2217return;
2218 };
22192220// Record whether we have any meshes that are to be drawn indirectly. If we
2221 // don't, then we can skip building indirect parameters.
2222let mut any_indirect = false;
22232224// Loop over each view.
2225for (view_entity, view) in &views {
2226let mut bind_groups = TypeIdHashMap::default();
22272228// Loop over each phase.
2229for (phase_type_id, phase_instance_buffers) in phase_instance_buffers {
2230let UntypedPhaseBatchedInstanceBuffers {
2231 data_buffer: ref data_buffer_vec,
2232ref work_item_buffers,
2233ref late_indexed_indirect_parameters_buffer,
2234ref late_non_indexed_indirect_parameters_buffer,
2235 } = *phase_instance_buffers;
22362237let Some(data_buffer) = data_buffer_vec.buffer() else {
2238continue;
2239 };
22402241// Grab the indirect parameters buffers for this phase.
2242let Some(phase_indirect_parameters_buffers) =
2243 indirect_parameters_buffers.get(phase_type_id)
2244else {
2245continue;
2246 };
22472248let Some(work_item_buffers) = work_item_buffers.get(&view.retained_view_entity) else {
2249continue;
2250 };
22512252// Create the `PreprocessBindGroupBuilder`.
2253let 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 };
22692270// Depending on the type of work items we have, construct the
2271 // appropriate bind groups.
2272let (was_indirect, bind_group) = match *work_item_buffers {
2273 PreprocessWorkItemBuffers::Direct(ref work_item_buffer) => (
2274false,
2275 preprocess_bind_group_builder
2276 .create_direct_preprocess_bind_groups(work_item_buffer),
2277 ),
22782279 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 } => (
2284true,
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 ),
22932294 PreprocessWorkItemBuffers::Indirect {
2295 indexed: ref indexed_work_item_buffer,
2296 non_indexed: ref non_indexed_work_item_buffer,
2297 gpu_occlusion_culling: None,
2298 } => (
2299true,
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 };
23072308// Write that bind group in.
2309if let Some(bind_group) = bind_group {
2310 any_indirect = any_indirect || was_indirect;
2311 bind_groups.insert(*phase_type_id, bind_group);
2312 }
2313 }
23142315// Save the bind groups.
2316commands
2317 .entity(view_entity)
2318 .insert(PreprocessBindGroups(bind_groups));
2319 }
23202321// Now, if there were any indirect draw commands, create the bind groups for
2322 // the indirect parameters building shader.
2323if any_indirect {
2324create_build_indirect_parameters_bind_groups(
2325&mut commands,
2326&render_device,
2327&pipeline_cache,
2328&pipelines,
2329current_input_buffer,
2330&indirect_parameters_buffers,
2331&indirect_parameters_build_jobs,
2332 );
2333 }
23342335// Create the bind groups we'll need for each dispatch of the bin unpacking
2336 // (`unpack_bins`) and uniform allocation (`allocate_uniforms`) shaders.
2337for (_, 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}
23592360/// 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.
2364view: Entity,
2365/// The indirect compute dispatch parameters buffer for indexed meshes in
2366 /// the late prepass.
2367late_indexed_indirect_parameters_buffer:
2368&'a RawBufferVec<LatePreprocessWorkItemIndirectParameters>,
2369/// The indirect compute dispatch parameters buffer for non-indexed meshes
2370 /// in the late prepass.
2371late_non_indexed_indirect_parameters_buffer:
2372&'a RawBufferVec<LatePreprocessWorkItemIndirectParameters>,
2373/// The device.
2374render_device: &'a RenderDevice,
2375/// The pipeline cache
2376pipeline_cache: &'a PipelineCache,
2377/// The buffers that store indirect draw parameters.
2378phase_indirect_parameters_buffers: &'a UntypedPhaseIndirectParametersBuffers,
2379/// The GPU buffer that stores the information needed to cull each mesh.
2380mesh_culling_data_buffer: &'a MeshCullingDataBuffer,
2381/// The device buffer that stores the information needed to process
2382 /// visibility ranges on the GPU.
2383visibility_range_data_buffer: &'a BufferVec<Vec4>,
2384/// The GPU buffer that stores information about the view.
2385view_uniforms: &'a ViewUniforms,
2386/// The GPU buffer that stores information about the view from last frame.
2387previous_view_uniforms: &'a PreviousViewUniforms,
2388/// The pipelines for the mesh preprocessing shader.
2389pipelines: &'a PreprocessPipelines,
2390/// The GPU buffer containing the list of [`MeshInputUniform`]s for the
2391 /// current frame.
2392current_input_buffer: &'a Buffer,
2393/// The GPU buffer containing the list of [`MeshInputUniform`]s for the
2394 /// previous frame.
2395previous_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.
2401data_buffer: &'a Buffer,
2402}
24032404impl<'a> PreprocessBindGroupBuilder<'a> {
2405/// Creates the bind groups for mesh preprocessing when GPU frustum culling
2406 /// and GPU occlusion culling are both disabled.
2407fn 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.
2414let work_item_buffer_size = NonZero::<u64>::try_from(
2415work_item_buffer.len() as u64 * u64::from(PreprocessWorkItem::min_size()),
2416 )
2417 .ok();
24182419Some(PhasePreprocessBindGroups::Direct(
2420self.render_device.create_bind_group(
2421"preprocess_direct_bind_group",
2422&self2423 .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 (
24305,
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 }
24422443/// Creates the bind groups for mesh preprocessing when GPU occlusion
2444 /// culling is enabled.
2445fn 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> {
2452let 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;
24572458let (view_depth_pyramid, previous_view_uniform_offset) =
2459 view_depth_pyramids.get(self.view).ok()?;
24602461Some(PhasePreprocessBindGroups::IndirectOcclusionCulling {
2462 early_indexed: self.create_indirect_occlusion_culling_early_indexed_bind_group(
2463view_depth_pyramid,
2464previous_view_uniform_offset,
2465indexed_work_item_buffer,
2466late_indexed_work_item_buffer,
2467 ),
24682469 early_non_indexed: self.create_indirect_occlusion_culling_early_non_indexed_bind_group(
2470view_depth_pyramid,
2471previous_view_uniform_offset,
2472non_indexed_work_item_buffer,
2473late_non_indexed_work_item_buffer,
2474 ),
24752476 late_indexed: self.create_indirect_occlusion_culling_late_indexed_bind_group(
2477view_depth_pyramid,
2478previous_view_uniform_offset,
2479late_indexed_work_item_buffer,
2480 ),
24812482 late_non_indexed: self.create_indirect_occlusion_culling_late_non_indexed_bind_group(
2483view_depth_pyramid,
2484previous_view_uniform_offset,
2485late_non_indexed_work_item_buffer,
2486 ),
2487 })
2488 }
24892490/// Creates the bind group for the first phase of mesh preprocessing of
2491 /// indexed meshes when GPU occlusion culling is enabled.
2492fn 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> {
2499let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2500let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2501let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2502let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
25032504match (
2505self.phase_indirect_parameters_buffers
2506 .indexed
2507 .metadata_buffer(),
2508indexed_work_item_buffer.buffer(),
2509late_indexed_work_item_buffer.buffer(),
2510self.late_indexed_indirect_parameters_buffer.buffer(),
2511 ) {
2512 (
2513Some(indexed_metadata_buffer),
2514Some(indexed_work_item_gpu_buffer),
2515Some(late_indexed_work_item_gpu_buffer),
2516Some(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.
2521let indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2522indexed_work_item_buffer.len() as u642523 * u64::from(PreprocessWorkItem::min_size()),
2524 )
2525 .ok();
25262527Some(
2528self.render_device.create_bind_group(
2529"preprocess_early_indexed_gpu_occlusion_culling_bind_group",
2530&self.pipeline_cache.get_bind_group_layout(
2531&self2532 .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(
25465,
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(
25742,
2575BufferBinding {
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(
258512,
2586BufferBinding {
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(
259613,
2597BufferBinding {
2598 buffer: late_indexed_indirect_parameters_buffer,
2599 offset: 0,
2600 size: NonZeroU64::new(
2601late_indexed_indirect_parameters_buffer.size(),
2602 ),
2603 },
2604 ),
2605 )),
2606 ),
2607 )
2608 }
2609_ => None,
2610 }
2611 }
26122613/// Creates the bind group for the first phase of mesh preprocessing of
2614 /// non-indexed meshes when GPU occlusion culling is enabled.
2615fn 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> {
2622let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2623let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2624let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2625let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
26262627match (
2628self.phase_indirect_parameters_buffers
2629 .non_indexed
2630 .metadata_buffer(),
2631non_indexed_work_item_buffer.buffer(),
2632late_non_indexed_work_item_buffer.buffer(),
2633self.late_non_indexed_indirect_parameters_buffer.buffer(),
2634 ) {
2635 (
2636Some(non_indexed_metadata_buffer),
2637Some(non_indexed_work_item_gpu_buffer),
2638Some(late_non_indexed_work_item_buffer),
2639Some(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.
2644let non_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2645non_indexed_work_item_buffer.len() as u642646 * u64::from(PreprocessWorkItem::min_size()),
2647 )
2648 .ok();
26492650Some(
2651self.render_device.create_bind_group(
2652"preprocess_early_non_indexed_gpu_occlusion_culling_bind_group",
2653&self.pipeline_cache.get_bind_group_layout(
2654&self2655 .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(
26695,
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(
26952,
2696BufferBinding {
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(
270612,
2707BufferBinding {
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(
271713,
2718BufferBinding {
2719 buffer: late_non_indexed_indirect_parameters_buffer,
2720 offset: 0,
2721 size: NonZeroU64::new(
2722late_non_indexed_indirect_parameters_buffer.size(),
2723 ),
2724 },
2725 ),
2726 )),
2727 ),
2728 )
2729 }
2730_ => None,
2731 }
2732 }
27332734/// Creates the bind group for the second phase of mesh preprocessing of
2735 /// indexed meshes when GPU occlusion culling is enabled.
2736fn 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> {
2742let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2743let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2744let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2745let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
27462747match (
2748self.phase_indirect_parameters_buffers
2749 .indexed
2750 .metadata_buffer(),
2751late_indexed_work_item_buffer.buffer(),
2752self.late_indexed_indirect_parameters_buffer.buffer(),
2753 ) {
2754 (
2755Some(indexed_metadata_buffer),
2756Some(late_indexed_work_item_gpu_buffer),
2757Some(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.
2762let late_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2763late_indexed_work_item_buffer.len() as u642764 * u64::from(PreprocessWorkItem::min_size()),
2765 )
2766 .ok();
27672768Some(
2769self.render_device.create_bind_group(
2770"preprocess_late_indexed_gpu_occlusion_culling_bind_group",
2771&self.pipeline_cache.get_bind_group_layout(
2772&self2773 .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(
27875,
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(
28152,
2816BufferBinding {
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(
282613,
2827BufferBinding {
2828 buffer: late_indexed_indirect_parameters_buffer,
2829 offset: 0,
2830 size: NonZeroU64::new(
2831late_indexed_indirect_parameters_buffer.size(),
2832 ),
2833 },
2834 ),
2835 )),
2836 ),
2837 )
2838 }
2839_ => None,
2840 }
2841 }
28422843/// Creates the bind group for the second phase of mesh preprocessing of
2844 /// non-indexed meshes when GPU occlusion culling is enabled.
2845fn 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> {
2851let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2852let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2853let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
2854let previous_view_buffer = self.previous_view_uniforms.uniforms.buffer()?;
28552856match (
2857self.phase_indirect_parameters_buffers
2858 .non_indexed
2859 .metadata_buffer(),
2860late_non_indexed_work_item_buffer.buffer(),
2861self.late_non_indexed_indirect_parameters_buffer.buffer(),
2862 ) {
2863 (
2864Some(non_indexed_metadata_buffer),
2865Some(non_indexed_work_item_gpu_buffer),
2866Some(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.
2871let non_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2872late_non_indexed_work_item_buffer.len() as u642873 * u64::from(PreprocessWorkItem::min_size()),
2874 )
2875 .ok();
28762877Some(
2878self.render_device.create_bind_group(
2879"preprocess_late_non_indexed_gpu_occlusion_culling_bind_group",
2880&self.pipeline_cache.get_bind_group_layout(
2881&self2882 .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(
28965,
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(
29242,
2925BufferBinding {
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(
293513,
2936BufferBinding {
2937 buffer: late_non_indexed_indirect_parameters_buffer,
2938 offset: 0,
2939 size: NonZeroU64::new(
2940late_non_indexed_indirect_parameters_buffer.size(),
2941 ),
2942 },
2943 ),
2944 )),
2945 ),
2946 )
2947 }
2948_ => None,
2949 }
2950 }
29512952/// Creates the bind groups for mesh preprocessing when GPU frustum culling
2953 /// is enabled, but GPU occlusion culling is disabled.
2954fn 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> {
2959Some(PhasePreprocessBindGroups::IndirectFrustumCulling {
2960 indexed: self2961 .create_indirect_frustum_culling_indexed_bind_group(indexed_work_item_buffer),
2962 non_indexed: self.create_indirect_frustum_culling_non_indexed_bind_group(
2963non_indexed_work_item_buffer,
2964 ),
2965 })
2966 }
29672968/// Creates the bind group for mesh preprocessing of indexed meshes when GPU
2969 /// frustum culling is enabled, but GPU occlusion culling is disabled.
2970fn create_indirect_frustum_culling_indexed_bind_group(
2971&self,
2972 indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
2973 ) -> Option<BindGroup> {
2974let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
2975let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
2976let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
29772978match (
2979self.phase_indirect_parameters_buffers
2980 .indexed
2981 .metadata_buffer(),
2982indexed_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.
2988let indexed_work_item_buffer_size = NonZero::<u64>::try_from(
2989indexed_work_item_buffer.len() as u642990 * u64::from(PreprocessWorkItem::min_size()),
2991 )
2992 .ok();
29932994Some(
2995self.render_device.create_bind_group(
2996"preprocess_gpu_indexed_frustum_culling_bind_group",
2997&self.pipeline_cache.get_bind_group_layout(
2998&self2999 .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 (
30075,
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 }
30263027/// Creates the bind group for mesh preprocessing of non-indexed meshes when
3028 /// GPU frustum culling is enabled, but GPU occlusion culling is disabled.
3029fn create_indirect_frustum_culling_non_indexed_bind_group(
3030&self,
3031 non_indexed_work_item_buffer: &PartialBufferVec<PreprocessWorkItem>,
3032 ) -> Option<BindGroup> {
3033let mesh_culling_data_buffer = self.mesh_culling_data_buffer.buffer()?;
3034let visibility_range_binding = self.visibility_range_data_buffer.binding()?;
3035let view_uniforms_binding = self.view_uniforms.uniforms.binding()?;
30363037match (
3038self.phase_indirect_parameters_buffers
3039 .non_indexed
3040 .metadata_buffer(),
3041non_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.
3047let non_indexed_work_item_buffer_size = NonZero::<u64>::try_from(
3048non_indexed_work_item_buffer.len() as u643049 * u64::from(PreprocessWorkItem::min_size()),
3050 )
3051 .ok();
30523053Some(
3054self.render_device.create_bind_group(
3055"preprocess_gpu_non_indexed_frustum_culling_bind_group",
3056&self.pipeline_cache.get_bind_group_layout(
3057&self3058 .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(
30725,
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}
31023103/// 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) {
3115let mut build_indirect_parameters_bind_groups = BuildIndirectParametersBindGroups::new();
31163117for (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 }),
31403141 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 }),
31603161 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 (
3168Some(indexed_indirect_parameters_metadata_buffer),
3169Some(indexed_indirect_parameters_data_buffer),
3170Some(indexed_batch_sets_buffer),
3171Some(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(
31901,
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(
32014,
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(
32145,
3215 indexed_indirect_parameters_data_buffer.as_entire_binding(),
3216 ),
3217 )),
3218 ),
3219 ),
3220_ => None,
3221 },
32223223 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 (
3234Some(non_indexed_indirect_parameters_metadata_buffer),
3235Some(non_indexed_indirect_parameters_data_buffer),
3236Some(non_indexed_batch_sets_buffer),
3237Some(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(
32591,
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()
3267as 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(
32804,
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(
32935,
3294 non_indexed_indirect_parameters_data_buffer.as_entire_binding(),
3295 ),
3296 )),
3297 ),
3298 ),
3299_ => None,
3300 },
3301 },
3302 );
3303 }
33043305commands.insert_resource(build_indirect_parameters_bind_groups);
3306}
33073308/// 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) {
3320let Some(bin_unpacking_metadata_buffer) =
3321scene_unpacking_buffers.bin_unpacking_metadata.buffer()
3322else {
3323return;
3324 };
33253326// We run the bin unpacking shader once per phase, so loop over all phases.
3327for phase_type_id in indirect_parameters_buffers.keys() {
3328// Fetch the buffers we need.
3329let Some(phase_batched_instance_buffers) = phase_instance_buffers.get(phase_type_id) else {
3330continue;
3331 };
3332let Some(work_item_buffers) = phase_batched_instance_buffers
3333 .work_item_buffers
3334 .get(view_entity)
3335else {
3336continue;
3337 };
3338let 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 })
3344else {
3345continue;
3346 };
33473348// Fetch the work item buffers.
3349let 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 };
3353let maybe_non_indexed_work_item_buffer = match *work_item_buffers {
3354 PreprocessWorkItemBuffers::Direct(ref raw_buffer_vec) => raw_buffer_vec.buffer(),
3355 PreprocessWorkItemBuffers::Indirect {
3356ref non_indexed, ..
3357 } => non_indexed.buffer(),
3358 };
33593360// Create the actual bind groups.
3361bin_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 {
3368Some(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,
3379true,
3380 )
3381 })
3382 .collect(),
3383None => ::alloc::vec::Vec::new()vec![],
3384 },
3385 non_indexed: match maybe_non_indexed_work_item_buffer {
3386Some(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,
3397false,
3398 )
3399 })
3400 .collect(),
3401None => ::alloc::vec::Vec::new()vec![],
3402 },
3403 },
3404 );
3405 }
3406}
34073408/// 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 {
3419let bind_group = render_device.create_bind_group(
3420if indexed {
3421"bin unpacking indexed bind group"
3422} else {
3423"bin unpacking non-indexed bind group"
3424},
3425&pipeline_cache3426 .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;
3431BindingResource::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>;
3439job.render_binned_mesh_instance_buffer.as_entire_binding(),
3440// @group(0) @binding(2) var<storage,
3441 // read_write> preprocess_work_items:
3442 // array<PreprocessWorkItem>;
3443work_item_buffer.as_entire_binding(),
3444// @group(0) @binding(3) var<storage> bin_metadata:
3445 // array<BinMetadata>;
3446job.bin_metadata_buffer.as_entire_binding(),
3447// @group(0) @binding(4) var<storage>
3448 // bin_index_to_bin_metadata_index: array<u32>;
3449job.bin_index_to_bin_metadata_index_buffer
3450 .as_entire_binding(),
3451 )),
3452 );
3453ViewPhaseBinUnpackingBindGroup {
3454 metadata_index: job.bin_unpacking_metadata_index,
3455bind_group,
3456 mesh_instance_count: job.mesh_instance_count,
3457 }
3458}
34593460/// 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) {
3471let Some(uniform_allocation_metadata_buffer) =
3472scene_unpacking_buffers.uniform_allocation_metadata.buffer()
3473else {
3474return;
3475 };
34763477for (phase_type_id, phase_indirect_parameters_buffers) in indirect_parameters_buffers.iter() {
3478let 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 })
3484else {
3485continue;
3486 };
34873488// Create the actual bind groups.
3489uniform_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() {
3496None => ::alloc::vec::Vec::new()vec![],
3497Some(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,
3509true,
3510 )
3511 })
3512 .collect()
3513 }
3514 },
3515 non_indexed: match phase_indirect_parameters_buffers
3516 .non_indexed
3517 .metadata_buffer()
3518 {
3519None => ::alloc::vec::Vec::new()vec![],
3520Some(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,
3532false,
3533 )
3534 })
3535 .collect()
3536 }
3537 },
3538 },
3539 );
3540 }
3541}
35423543/// 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 {
3554let bind_group = render_device.create_bind_group(
3555if 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_pipelines3563 .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;
3570BindingResource::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>;
3577job.bin_metadata_buffer.as_entire_binding(),
3578// @group(0) @binding(2) var<storage, read_write>
3579 // indirect_parameters_metadata: array<IndirectParametersMetadata>;
3580indirect_parameters_metadata_buffer.as_entire_binding(),
3581// @group(0) @binding(3) var<storage, read_write> fan_buffer:
3582 // array<u32>;
3583job.fan_buffer.as_entire_binding(),
3584 )),
3585 );
3586ViewPhaseUniformAllocationBindGroup {
3587 metadata_index: job.uniform_allocation_metadata_index,
3588bind_group,
3589 bin_count: job.bin_count,
3590 }
3591}
35923593/// 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>,
3597mut mesh_culling_data_buffer: ResMut<MeshCullingDataBuffer>,
3598 pipeline_cache: Res<PipelineCache>,
3599mut sparse_buffer_update_jobs: ResMut<SparseBufferUpdateJobs>,
3600mut sparse_buffer_update_bind_groups: ResMut<SparseBufferUpdateBindGroups>,
3601 sparse_buffer_update_pipelines: Res<SparseBufferUpdatePipelines>,
3602) {
3603mesh_culling_data_buffer.write_buffers(&render_device, &render_queue);
3604mesh_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}