Book a Demo!
CoCalc Logo Icon
StoreFeaturesDocsShareSupportNewsAboutPoliciesSign UpSign In
bevyengine
GitHub Repository: bevyengine/bevy
Path: blob/main/crates/bevy_pbr/src/render/unpack_bins.wgsl
30636 views
// A compute shader that unpacks bins.
//
// This shader runs before mesh preprocessing in order to generate
// `PreprocessWorkItem`s from cached bins. A single dispatch of this shader
// corresponds to a single batch set, and one invocation of this shader
// corresponds to one binned entity. Each shader invocation builds the work item
// for its entity by copying over the mesh input uniform index and calculating
// the position of the command in the indirect parameters buffer that will draw
// that entity.

#import bevy_pbr::mesh_preprocess_types::PreprocessWorkItem

// Information needed to unpack bins belonging to a single batch set.
struct BinUnpackingMetadata {
    // The index of the first `PreprocessWorkItem` that this compute shader
    // dispatch is to write to.
    base_output_work_item_index: u32,
    // The index of the first GPU indirect parameters command for this batch
    // set.
    base_indirect_parameters_index: u32,
    // The number of binned mesh instances in the `binned_mesh_instances`
    // array.
    binned_mesh_instance_count: u32,
    // Padding.
    pad_a: u32,
    // Padding.
    pad_b: array<vec4<u32>, 15>,
};

// One mesh instance in a bin.
//
// This corresponds to the CPU-side `GpuRenderBinnedMeshInstance` structure.
//
// Note that this structure isn't sorted within the
// `binned_mesh_instances` buffer. Instances from the same bin aren't
// guaranteed to be adjacent to one another.
struct BinnedMeshInstance {
    // The index of the `MeshInputUniform` corresponding to this mesh instance
    // in the `MeshInputUniform` buffer.
    input_uniform_index: u32,
    // The index of the bin that this mesh instance belongs to.
    bin_index: u32,
};

// Metadata for the entire batch set.
@group(0) @binding(0) var<uniform> bin_unpacking_metadata: BinUnpackingMetadata;

// The input array of `BinnedMeshInstance`s.
//
// Note that this array isn't sorted.
@group(0) @binding(1) var<storage> binned_mesh_instances: array<BinnedMeshInstance>;

// The output list of `PreprocessWorkItem`s.
@group(0) @binding(2) var<storage, read_write> preprocess_work_items: array<PreprocessWorkItem>;

// A mapping from each `bin_index` to the index of the GPU indirect parameters
// for this bin, relative to the start of the indirect parameters for this batch
// set.
@group(0) @binding(3) var<storage> bin_index_to_indirect_parameters_offset: array<u32>;

@compute
@workgroup_size(64)
fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
    // Figure out which instance we're looking at.
    let global_id = global_invocation_id.x;
    if (global_id >= bin_unpacking_metadata.binned_mesh_instance_count) {
        return;
    }

    // Unpack the `BinnedMeshInstance`.
    let input_uniform_index = binned_mesh_instances[global_id].input_uniform_index;
    let bin_index = binned_mesh_instances[global_id].bin_index;

    // Look up the indirect parameters index for this bin, relative to the first
    // indirect parameters offset for this batch set.
    let indirect_parameters_offset = bin_index_to_indirect_parameters_offset[bin_index];

    // Determine the location we should write the work item to.
    let output_index = bin_unpacking_metadata.base_output_work_item_index + global_id;

    // Write out the resulting work item.
    preprocess_work_items[output_index].input_index = input_uniform_index;
    preprocess_work_items[output_index].output_or_indirect_parameters_index =
        bin_unpacking_metadata.base_indirect_parameters_index + indirect_parameters_offset;
}