Search across packages, versions, files, declarations, and comments. Use multiple terms to narrow results (AND), and add @vertex, @fragment, or @compute to filter by shader stage.
1//! A compute shader that unpacks bins.
2//!
3//! This shader runs before mesh preprocessing in order to generate
4//! `PreprocessWorkItem`s from cached bins. A single dispatch of this shader
5//! corresponds to a single batch set, and one invocation of this shader
6//! corresponds to one binned entity. Each shader invocation builds the work item
7//! for its entity by copying over the mesh input uniform index and calculating
8//! the position of the command in the indirect parameters buffer that will draw
9//! that entity.
11import bevy_render::occlusion_culling::mesh_preprocess_types::{BinMetadata, PreprocessWorkItem};
13/// Information needed to unpack bins belonging to a single batch set.
14struct BinUnpackingMetadata {
15 /// The index of the first `PreprocessWorkItem` that this compute shader
16 /// dispatch is to write to.
17 base_output_work_item_index: u32,
18 /// The index of the first GPU indirect parameters command for this batch
19 /// set.
20 base_indirect_parameters_index: u32,
21 /// The number of binned mesh instances in the `binned_mesh_instances`
22 /// array.
23 binned_mesh_instance_count: u32,
24 /// Padding.
25 pad_a: u32,
26 /// Padding.
27 pad_b: array<vec4<u32>, 15>,
28};
30/// One mesh instance in a bin.
31///
32/// This corresponds to the CPU-side `GpuRenderBinnedMeshInstance` structure.
33///
34/// Note that this structure isn't sorted within the
35/// `binned_mesh_instances` buffer. Instances from the same bin aren't
36/// guaranteed to be adjacent to one another.
37struct BinnedMeshInstance {
38 /// The index of the `MeshInputUniform` corresponding to this mesh instance
39 /// in the `MeshInputUniform` buffer.
40 input_uniform_index: u32,
41 /// The index of the bin that this mesh instance belongs to.
42 bin_index: u32,
43};
45/// Metadata for the entire batch set.
46@group(0) @binding(0) var<uniform> bin_unpacking_metadata: BinUnpackingMetadata;
48/// The input array of `BinnedMeshInstance`s.
49///
50/// Note that this array isn't sorted.
51@group(0) @binding(1) var<storage> binned_mesh_instances: array<BinnedMeshInstance>;
53/// The output list of `PreprocessWorkItem`s.
54@group(0) @binding(2) var<storage, read_write> preprocess_work_items: array<PreprocessWorkItem>;
56/// The bin metadata, which contains the index of the GPU indirect parameters for
57/// this bin, relative to the start of the indirect parameters for this batch
58/// set.
59///
60/// This is indexed by the bin metadata index below.
61@group(0) @binding(3) var<storage> bin_metadata: array<BinMetadata>;
63/// A mapping from the bin index to the metadata index within the `bin_metadata`
64/// buffer.
65@group(0) @binding(4) var<storage> bin_index_to_bin_metadata_index: array<u32>;
67@compute
68@workgroup_size(64)
69fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
70 // Figure out which instance we're looking at.
71 let global_id = global_invocation_id.x;
72 if (global_id >= bin_unpacking_metadata.binned_mesh_instance_count) {
73 return;
74 }
76 // Unpack the `BinnedMeshInstance`.
77 let input_uniform_index = binned_mesh_instances[global_id].input_uniform_index;
78 let bin_index = binned_mesh_instances[global_id].bin_index;
80 // Look up the indirect parameters index for this bin, relative to the first
81 // indirect parameters offset for this batch set.
82 let bin_metadata_index = bin_index_to_bin_metadata_index[bin_index];
83 let indirect_parameters_offset = bin_metadata[bin_metadata_index].indirect_parameters_offset;
85 // Determine the location we should write the work item to.
86 let output_index = bin_unpacking_metadata.base_output_work_item_index + global_id;
88 // Write out the resulting work item.
89 preprocess_work_items[output_index].input_index = input_uniform_index;
90 preprocess_work_items[output_index].output_or_indirect_parameters_index =
91 bin_unpacking_metadata.base_indirect_parameters_index + indirect_parameters_offset;
Contents