//! Builds GPU indirect draw parameters from metadata.
//!
//! This only runs when indirect drawing is enabled. It takes the output of
//! `mesh_preprocess.wesl` and creates indirect parameters for the GPU.
//!
//! This shader runs separately for indexed and non-indexed meshes. Unlike
//! `mesh_preprocess.wesl`, which runs one instance per mesh *instance*, one
//! instance of this shader corresponds to a single *batch* which could contain
//! arbitrarily many instances of a single mesh.
import bevy_render::occlusion_culling::mesh_preprocess_types::{
IndirectBatchSet,
IndirectParametersIndexed,
IndirectParametersNonIndexed,
IndirectParametersMetadata,
MeshInput
};
/// Specifies the batches that this shader invocation is to process.
struct IndirectParametersBuildJob {
/// The first batch index that this shader invocation should process
/// (inclusive).
first_batch_index: u32,
/// The last batch index that this shader invocation should process
/// (exclusive).
last_batch_index: u32,
}
/// The data for each mesh that the CPU supplied to the GPU.
@group(0) @binding(0) var<storage> current_input: array<MeshInput>;
/// Data that we use to generate the indirect parameters.
///
/// The `mesh_preprocess.wesl` shader emits these.
@group(0) @binding(1) var<storage> indirect_parameters_metadata:
array<IndirectParametersMetadata>;
/// Information about each batch set.
///
/// A *batch set* is a set of meshes that might be multi-drawn together.
@group(0) @binding(3) var<storage, read_write> indirect_batch_sets: array<IndirectBatchSet>;
/// Specifies the batches that this shader invocation is to process.
@group(0) @binding(4) var<uniform> indirect_parameters_build_job: IndirectParametersBuildJob;
@if(INDEXED)
/// The buffer of indirect draw parameters that we generate, and that the GPU
/// reads to issue the draws.
///
/// This buffer is for indexed meshes.
@group(0) @binding(5) var<storage, read_write> indirect_parameters:
array<IndirectParametersIndexed>;
@else // INDEXED
/// The buffer of indirect draw parameters that we generate, and that the GPU
/// reads to issue the draws.
///
/// This buffer is for non-indexed meshes.
@group(0) @binding(5) var<storage, read_write> indirect_parameters:
array<IndirectParametersNonIndexed>;
@compute
@workgroup_size(64)
fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
// Figure out our instance index (i.e. batch index). If this thread doesn't
// correspond to a valid index in the range.
let instance_index = global_invocation_id.x + indirect_parameters_build_job.first_batch_index;
if (instance_index >= indirect_parameters_build_job.last_batch_index) {
return;
}
// Unpack the metadata for this batch.
let base_output_index = indirect_parameters_metadata[instance_index].base_output_index;
let batch_set_index = indirect_parameters_metadata[instance_index].batch_set_index;
let mesh_index = indirect_parameters_metadata[instance_index].mesh_index;
// If we aren't using `multi_draw_indirect_count`, we have a 1:1 fixed
// assignment of batches to slots in the indirect parameters buffer, so we
// can just use the instance index as the index of our indirect parameters.
let early_instance_count =
indirect_parameters_metadata[instance_index].early_instance_count;
let late_instance_count = indirect_parameters_metadata[instance_index].late_instance_count;
// If in the early phase, we draw only the early meshes. If in the late
// phase, we draw only the late meshes. If in the main phase, draw all the
// meshes.
@if(EARLY_PHASE)
let instance_count = early_instance_count;
@elif(LATE_PHASE)
let instance_count = late_instance_count;
@else // LATE_PHASE
let instance_count = early_instance_count + late_instance_count;
var indirect_parameters_index = instance_index;
// If the current hardware and driver support `multi_draw_indirect_count`,
// dynamically reserve an index for the indirect parameters we're to
// generate.
@if(MULTI_DRAW_INDIRECT_COUNT_SUPPORTED)
// If this batch belongs to a batch set, then allocate space for the
// indirect commands in that batch set.
if (batch_set_index != 0xffffffffu) {
// Bail out now if there are no instances. Note that we can only bail if
// we're in a batch set. That's because only batch sets are drawn using
// `multi_draw_indirect_count`. If we aren't using
// `multi_draw_indirect_count`, then we need to continue in order to
// zero out the instance count; otherwise, it'll have garbage data in
// it.
if (instance_count == 0u) {
return;
}
let indirect_parameters_base =
indirect_batch_sets[batch_set_index].indirect_parameters_base;
let indirect_parameters_offset =
atomicAdd(&indirect_batch_sets[batch_set_index].indirect_parameters_count, 1u);
indirect_parameters_index = indirect_parameters_base + indirect_parameters_offset;
}
// Build up the indirect parameters. The structures for indexed and
// non-indexed meshes are slightly different.
indirect_parameters[indirect_parameters_index].instance_count = instance_count;
@if(LATE_PHASE)
// The late mesh instances are stored after the early mesh instances, so we
// offset the output index by the number of early mesh instances.
indirect_parameters[indirect_parameters_index].first_instance =
base_output_index + early_instance_count;
@else // LATE_PHASE
indirect_parameters[indirect_parameters_index].first_instance = base_output_index;
indirect_parameters[indirect_parameters_index].base_vertex =
current_input[mesh_index].first_vertex_index;
@if(INDEXED)
indirect_parameters[indirect_parameters_index].index_count =
current_input[mesh_index].index_count;
@if(INDEXED)
indirect_parameters[indirect_parameters_index].first_index =
current_input[mesh_index].first_index_index;
@if(!INDEXED)
indirect_parameters[indirect_parameters_index].vertex_count =
current_input[mesh_index].index_count;
}