Skip to main content

bevy_pbr/meshlet/
mod.rs

1//! Render high-poly 3d meshes using an efficient GPU-driven method. See [`MeshletPlugin`] and [`MeshletMesh`] for details.
2
3mod asset;
4#[cfg(feature = "meshlet_processor")]
5mod from_mesh;
6mod instance_manager;
7mod material_pipeline_prepare;
8mod material_shade_nodes;
9mod meshlet_mesh_manager;
10mod persistent_buffer;
11mod persistent_buffer_impls;
12mod pipelines;
13mod resource_manager;
14mod visibility_buffer_raster_node;
15
16pub(crate) use self::{
17    instance_manager::{queue_material_meshlet_meshes, InstanceManager},
18    material_pipeline_prepare::{
19        prepare_material_meshlet_meshes_main_opaque_pass, prepare_material_meshlet_meshes_prepass,
20    },
21};
22
23pub use self::asset::{
24    MeshletMesh, MeshletMeshLoader, MeshletMeshSaver, MESHLET_MESH_ASSET_VERSION,
25};
26#[cfg(feature = "meshlet_processor")]
27pub use self::from_mesh::{
28    MeshToMeshletMeshConversionError, MESHLET_DEFAULT_VERTEX_POSITION_QUANTIZATION_FACTOR,
29};
30use self::{
31    instance_manager::extract_meshlet_mesh_entities,
32    material_pipeline_prepare::{
33        MeshletViewMaterialsDeferredGBufferPrepass, MeshletViewMaterialsMainOpaquePass,
34        MeshletViewMaterialsPrepass,
35    },
36    material_shade_nodes::{
37        meshlet_deferred_gbuffer_prepass, meshlet_main_opaque_pass, meshlet_prepass,
38    },
39    meshlet_mesh_manager::perform_pending_meshlet_mesh_writes,
40    pipelines::*,
41    resource_manager::{
42        prepare_meshlet_per_frame_resources, prepare_meshlet_view_bind_groups, ResourceManager,
43    },
44    visibility_buffer_raster_node::meshlet_visibility_buffer_raster,
45};
46use crate::render::{per_view_shadow_pass, EARLY_SHADOW_PASS};
47use crate::{meshlet::meshlet_mesh_manager::init_meshlet_mesh_manager, PreviousGlobalTransform};
48use bevy_app::{App, Plugin};
49use bevy_asset::{embedded_asset, AssetApp, AssetId, Handle};
50use bevy_camera::visibility::{self, Visibility, VisibilityClass};
51use bevy_core_pipeline::{
52    core_3d::main_opaque_pass_3d,
53    prepass::{DeferredPrepass, MotionVectorPrepass, NormalPrepass},
54    schedule::{Core3d, Core3dSystems},
55};
56use bevy_derive::{Deref, DerefMut};
57use bevy_ecs::{
58    component::Component,
59    entity::Entity,
60    query::Has,
61    reflect::ReflectComponent,
62    schedule::IntoScheduleConfigs,
63    system::{Commands, Query, Res},
64    template::FromTemplate,
65};
66use bevy_reflect::{std_traits::ReflectDefault, Reflect};
67use bevy_render::{
68    renderer::RenderDevice,
69    settings::WgpuFeatures,
70    view::{prepare_view_targets, Msaa},
71    ExtractSchedule, Render, RenderApp, RenderStartup, RenderSystems,
72};
73use bevy_shader::load_shader_library;
74use bevy_transform::components::Transform;
75use derive_more::From;
76use tracing::error;
77
78/// Provides a plugin for rendering large amounts of high-poly 3d meshes using an efficient GPU-driven method. See also [`MeshletMesh`].
79///
80/// Rendering dense scenes made of high-poly meshes with thousands or millions of triangles is extremely expensive in Bevy's standard renderer.
81/// Once meshes are pre-processed into a [`MeshletMesh`], this plugin can render these kinds of scenes very efficiently.
82///
83/// In comparison to Bevy's standard renderer:
84/// * Much more efficient culling. Meshlets can be culled individually, instead of all or nothing culling for entire meshes at a time.
85///   Additionally, occlusion culling can eliminate meshlets that would cause overdraw.
86/// * Much more efficient batching. All geometry can be rasterized in a single draw.
87/// * Scales better with large amounts of dense geometry and overdraw. Bevy's standard renderer will bottleneck sooner.
88/// * Near-seamless level of detail (LOD).
89/// * Much greater base overhead. Rendering will be slower and use more memory than Bevy's standard renderer
90///   with small amounts of geometry and overdraw.
91/// * Requires preprocessing meshes. See [`MeshletMesh`] for details.
92/// * Limitations on the kinds of materials you can use. See [`MeshletMesh`] for details.
93///
94/// This plugin requires a fairly recent GPU that supports [`WgpuFeatures::TEXTURE_INT64_ATOMIC`].
95///
96/// This plugin currently works only on the Vulkan and Metal backends.
97///
98/// This plugin is not compatible with [`Msaa`]. Any camera rendering a [`MeshletMesh`] must have
99/// [`Msaa`] set to [`Msaa::Off`].
100///
101/// Mixing forward+prepass and deferred rendering for opaque materials is not currently supported when using this plugin.
102/// You must use one or the other by setting [`crate::DefaultOpaqueRendererMethod`].
103/// Do not override [`crate::Material::opaque_render_method`] for any material when using this plugin.
104///
105/// ![A render of the Stanford dragon as a `MeshletMesh`](https://raw.githubusercontent.com/bevyengine/bevy/main/crates/bevy_pbr/src/meshlet/meshlet_preview.png)
106pub struct MeshletPlugin {
107    /// The maximum amount of clusters that can be processed at once,
108    /// used to control the size of a pre-allocated GPU buffer.
109    ///
110    /// If this number is too low, you'll see rendering artifacts like missing or blinking meshes.
111    ///
112    /// Each cluster slot costs 4 bytes of VRAM.
113    ///
114    /// Must not be greater than 2^25.
115    pub cluster_buffer_slots: u32,
116}
117
118impl MeshletPlugin {
119    /// [`WgpuFeatures`] required for this plugin to function.
120    pub fn required_wgpu_features() -> WgpuFeatures {
121        WgpuFeatures::TEXTURE_INT64_ATOMIC
122            | WgpuFeatures::TEXTURE_ATOMIC
123            | WgpuFeatures::SHADER_INT64
124            | WgpuFeatures::SUBGROUP
125            | WgpuFeatures::DEPTH_CLIP_CONTROL
126            | WgpuFeatures::IMMEDIATES
127    }
128}
129
130impl Plugin for MeshletPlugin {
131    fn build(&self, app: &mut App) {
132        #[cfg(target_endian = "big")]
133        compile_error!("MeshletPlugin is only supported on little-endian processors.");
134
135        if self.cluster_buffer_slots > 2_u32.pow(25) {
136            {
    use ::tracing::__macro_support::Callsite as _;
    static __CALLSITE: ::tracing::callsite::DefaultCallsite =
        {
            static META: ::tracing::Metadata<'static> =
                {
                    ::tracing_core::metadata::Metadata::new("event src/meshlet/mod.rs:136",
                        "bevy_pbr::meshlet", ::tracing::Level::ERROR,
                        ::tracing_core::__macro_support::Option::Some("src/meshlet/mod.rs"),
                        ::tracing_core::__macro_support::Option::Some(136u32),
                        ::tracing_core::__macro_support::Option::Some("bevy_pbr::meshlet"),
                        ::tracing_core::field::FieldSet::new(&["message"],
                            ::tracing_core::callsite::Identifier(&__CALLSITE)),
                        ::tracing::metadata::Kind::EVENT)
                };
            ::tracing::callsite::DefaultCallsite::new(&META)
        };
    let enabled =
        ::tracing::Level::ERROR <= ::tracing::level_filters::STATIC_MAX_LEVEL
                &&
                ::tracing::Level::ERROR <=
                    ::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!("MeshletPlugin::cluster_buffer_slots must not be greater than 2^25.")
                                            as &dyn ::tracing::field::Value))])
            });
    } else { ; }
};error!("MeshletPlugin::cluster_buffer_slots must not be greater than 2^25.");
137            std::process::exit(1);
138        }
139
140        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "bindings.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "bindings.wesl");
        embedded.insert_asset(watched_path, &path,
            b"import package::render::mesh_types::Mesh;\nimport bevy_render::view::View;\nimport package::prepass::bindings::PreviousViewUniforms;\nimport bevy_render::utils::octahedral_decode_signed;\n\nstruct BvhNode {\n    aabbs: array<MeshletAabbErrorOffset, 8>,\n    lod_bounds: array<vec4<f32>, 8>,\n    child_counts: array<u32, 2>,\n    _padding: vec2<u32>,\n}\n\nstruct Meshlet {\n    start_vertex_position_bit: u32,\n    start_vertex_attribute_id: u32,\n    start_index_id: u32,\n    packed_a: u32,\n    packed_b: u32,\n    min_vertex_position_channel_x: f32,\n    min_vertex_position_channel_y: f32,\n    min_vertex_position_channel_z: f32,\n}\n\nfn get_meshlet_vertex_count(meshlet: ptr<function, Meshlet>) -> u32 {\n    return extractBits((*meshlet).packed_a, 0u, 8u) + 1u;\n}\n\nfn get_meshlet_triangle_count(meshlet: ptr<function, Meshlet>) -> u32 {\n    return extractBits((*meshlet).packed_a, 8u, 8u);\n}\n\nstruct MeshletCullData {\n    aabb: MeshletAabbErrorOffset,\n    lod_group_sphere: vec4<f32>,\n}\n\nstruct MeshletAabb {\n    center: vec3<f32>,\n    half_extent: vec3<f32>,\n}\n\nstruct MeshletAabbErrorOffset {\n    center_and_error: vec4<f32>,\n    half_extent_and_child_offset: vec4<f32>,\n}\n\nfn get_aabb(aabb: ptr<function, MeshletAabbErrorOffset>) -> MeshletAabb {\n    return MeshletAabb(\n        (*aabb).center_and_error.xyz,\n        (*aabb).half_extent_and_child_offset.xyz,\n    );\n}\n\nfn get_aabb_error(aabb: ptr<function, MeshletAabbErrorOffset>) -> f32 {\n    return (*aabb).center_and_error.w;\n}\n\nfn get_aabb_child_offset(aabb: ptr<function, MeshletAabbErrorOffset>) -> u32 {\n    return bitcast<u32>((*aabb).half_extent_and_child_offset.w);\n}\n\nstruct DispatchIndirectArgs {\n    x: atomic<u32>,\n    y: u32,\n    z: u32,\n}\n\nstruct DrawIndirectArgs {\n    vertex_count: u32,\n    instance_count: atomic<u32>,\n    first_vertex: u32,\n    first_instance: u32,\n}\n\n/// Either a BVH node or a meshlet, along with the instance it is associated with.\n/// Refers to BVH nodes in `meshlet_bvh_cull_queue` and `meshlet_second_pass_bvh_queue`, where `offset` is the index into `meshlet_bvh_nodes`.\n/// Refers to meshlets in `meshlet_meshlet_cull_queue` and `meshlet_raster_clusters`.\n/// In `meshlet_meshlet_cull_queue`, `offset` is the index into `meshlet_cull_data`.\n/// In `meshlet_raster_clusters`, `offset` is the index into `meshlets`.\nstruct InstancedOffset {\n    instance_id: u32,\n    offset: u32,\n}\n\nconst CENTIMETERS_PER_METER = 100.0;\n\n@if(MESHLET_INSTANCE_CULLING_PASS)\nstruct Constants { scene_instance_count: u32 }\n@if(MESHLET_INSTANCE_CULLING_PASS) {\nvar<immediate> constants: Constants;\n\n// Cull data\n@group(0) @binding(0) var depth_pyramid: texture_2d<f32>;\n@group(0) @binding(1) var<uniform> view: View;\n@group(0) @binding(2) var<uniform> previous_view: PreviousViewUniforms;\n\n// Per entity instance data\n@group(0) @binding(3) var<storage, read> meshlet_instance_uniforms: array<Mesh>;\n@group(0) @binding(4) var<storage, read> meshlet_view_instance_visibility: array<u32>; // 1 bit per entity instance, packed as a bitmask\n@group(0) @binding(5) var<storage, read> meshlet_instance_aabbs: array<MeshletAabb>;\n@group(0) @binding(6) var<storage, read> meshlet_instance_bvh_root_nodes: array<u32>;\n\n// BVH cull queue data\n@group(0) @binding(7) var<storage, read_write> meshlet_bvh_cull_count_write: atomic<u32>;\n@group(0) @binding(8) var<storage, read_write> meshlet_bvh_cull_dispatch: DispatchIndirectArgs;\n@group(0) @binding(9) var<storage, read_write> meshlet_bvh_cull_queue: array<InstancedOffset>;\n}\n\n// Second pass queue data\n@if(MESHLET_INSTANCE_CULLING_PASS && MESHLET_FIRST_CULLING_PASS) {\n@group(0) @binding(10) var<storage, read_write> meshlet_second_pass_instance_count: atomic<u32>;\n@group(0) @binding(11) var<storage, read_write> meshlet_second_pass_instance_dispatch: DispatchIndirectArgs;\n@group(0) @binding(12) var<storage, read_write> meshlet_second_pass_instance_candidates: array<u32>;\n}\n@if(MESHLET_INSTANCE_CULLING_PASS && !MESHLET_FIRST_CULLING_PASS) {\n@group(0) @binding(10) var<storage, read> meshlet_second_pass_instance_count: u32;\n@group(0) @binding(11) var<storage, read> meshlet_second_pass_instance_candidates: array<u32>;\n}\n\n@if(MESHLET_BVH_CULLING_PASS)\nstruct Constants { read_from_front: u32, rightmost_slot: u32 }\n@if(MESHLET_BVH_CULLING_PASS) {\nvar<immediate> constants: Constants;\n\n// Cull data\n@group(0) @binding(0) var depth_pyramid: texture_2d<f32>; // From the end of the last frame for the first culling pass, and from the first raster pass for the second culling pass\n@group(0) @binding(1) var<uniform> view: View;\n@group(0) @binding(2) var<uniform> previous_view: PreviousViewUniforms;\n\n// Global mesh data\n@group(0) @binding(3) var<storage, read> meshlet_bvh_nodes: array<BvhNode>;\n\n// Per entity instance data\n@group(0) @binding(4) var<storage, read> meshlet_instance_uniforms: array<Mesh>;\n\n// BVH cull queue data\n@group(0) @binding(5) var<storage, read> meshlet_bvh_cull_count_read: u32;\n@group(0) @binding(6) var<storage, read_write> meshlet_bvh_cull_count_write: atomic<u32>;\n@group(0) @binding(7) var<storage, read_write> meshlet_bvh_cull_dispatch: DispatchIndirectArgs;\n@group(0) @binding(8) var<storage, read_write> meshlet_bvh_cull_queue: array<InstancedOffset>;\n\n// Meshlet cull queue data\n@group(0) @binding(9) var<storage, read_write> meshlet_meshlet_cull_count_early: atomic<u32>;\n@group(0) @binding(10) var<storage, read_write> meshlet_meshlet_cull_count_late: atomic<u32>;\n@group(0) @binding(11) var<storage, read_write> meshlet_meshlet_cull_dispatch_early: DispatchIndirectArgs;\n@group(0) @binding(12) var<storage, read_write> meshlet_meshlet_cull_dispatch_late: DispatchIndirectArgs;\n@group(0) @binding(13) var<storage, read_write> meshlet_meshlet_cull_queue: array<InstancedOffset>;\n}\n\n// Second pass queue data\n@if(MESHLET_BVH_CULLING_PASS && MESHLET_FIRST_CULLING_PASS) {\n@group(0) @binding(14) var<storage, read_write> meshlet_second_pass_bvh_count: atomic<u32>;\n@group(0) @binding(15) var<storage, read_write> meshlet_second_pass_bvh_dispatch: DispatchIndirectArgs;\n@group(0) @binding(16) var<storage, read_write> meshlet_second_pass_bvh_queue: array<InstancedOffset>;\n}\n\n@if(MESHLET_CLUSTER_CULLING_PASS)\nstruct Constants { rightmost_slot: u32 }\n@if(MESHLET_CLUSTER_CULLING_PASS) {\nvar<immediate> constants: Constants;\n\n// Cull data\n@group(0) @binding(0) var depth_pyramid: texture_2d<f32>; // From the end of the last frame for the first culling pass, and from the first raster pass for the second culling pass\n@group(0) @binding(1) var<uniform> view: View;\n@group(0) @binding(2) var<uniform> previous_view: PreviousViewUniforms;\n\n// Global mesh data\n@group(0) @binding(3) var<storage, read> meshlet_cull_data: array<MeshletCullData>;\n\n// Per entity instance data\n@group(0) @binding(4) var<storage, read> meshlet_instance_uniforms: array<Mesh>;\n\n// Raster queue data\n@group(0) @binding(5) var<storage, read_write> meshlet_software_raster_indirect_args: DispatchIndirectArgs;\n@group(0) @binding(6) var<storage, read_write> meshlet_hardware_raster_indirect_args: DrawIndirectArgs;\n@group(0) @binding(7) var<storage, read> meshlet_previous_raster_counts: array<u32>;\n@group(0) @binding(8) var<storage, read_write> meshlet_raster_clusters: array<InstancedOffset>;\n\n// Meshlet cull queue data\n@group(0) @binding(9) var<storage, read> meshlet_meshlet_cull_count_read: u32;\n}\n\n// Second pass queue data\n@if(MESHLET_CLUSTER_CULLING_PASS && MESHLET_FIRST_CULLING_PASS) {\n@group(0) @binding(10) var<storage, read_write> meshlet_meshlet_cull_count_write: atomic<u32>;\n@group(0) @binding(11) var<storage, read_write> meshlet_meshlet_cull_dispatch: DispatchIndirectArgs;\n@group(0) @binding(12) var<storage, read_write> meshlet_meshlet_cull_queue: array<InstancedOffset>;\n}\n@if(MESHLET_CLUSTER_CULLING_PASS && !MESHLET_FIRST_CULLING_PASS)\n@group(0) @binding(10) var<storage, read> meshlet_meshlet_cull_queue: array<InstancedOffset>;\n\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS) {\n@group(0) @binding(0) var<storage, read> meshlet_raster_clusters: array<InstancedOffset>; // Per cluster\n@group(0) @binding(1) var<storage, read> meshlets: array<Meshlet>; // Per meshlet\n@group(0) @binding(2) var<storage, read> meshlet_indices: array<u32>; // Many per meshlet\n@group(0) @binding(3) var<storage, read> meshlet_vertex_positions: array<u32>; // Many per meshlet\n@group(0) @binding(4) var<storage, read> meshlet_instance_uniforms: array<Mesh>; // Per entity instance\n@group(0) @binding(5) var<storage, read> meshlet_previous_raster_counts: array<u32>;\n@group(0) @binding(6) var<storage, read> meshlet_software_raster_cluster_count: u32;\n}\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS && MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n@group(0) @binding(7) var meshlet_visibility_buffer: texture_storage_2d<r64uint, atomic>;\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS && !MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n@group(0) @binding(7) var meshlet_visibility_buffer: texture_storage_2d<r32uint, atomic>;\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS)\n@group(0) @binding(8) var<uniform> view: View;\n\n// TODO: Load only twice, instead of 3x in cases where you load 3 indices per thread?\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS)\nfn get_meshlet_vertex_id(index_id: u32) -> u32 {\n    let packed_index = meshlet_indices[index_id / 4u];\n    let bit_offset = (index_id % 4u) * 8u;\n    return extractBits(packed_index, bit_offset, 8u);\n}\n\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS)\nfn get_meshlet_vertex_position(meshlet: ptr<function, Meshlet>, vertex_id: u32) -> vec3<f32> {\n    // Get bitstream start for the vertex\n    let unpacked = unpack4xU8((*meshlet).packed_b);\n    let bits_per_channel = unpacked.xyz;\n    let bits_per_vertex = bits_per_channel.x + bits_per_channel.y + bits_per_channel.z;\n    var start_bit = (*meshlet).start_vertex_position_bit + (vertex_id * bits_per_vertex);\n\n    // Read each vertex channel from the bitstream\n    var vertex_position_packed = vec3(0u);\n    for (var i = 0u; i < 3u; i++) {\n        let lower_word_index = start_bit / 32u;\n        let lower_word_bit_offset = start_bit & 31u;\n        var next_32_bits = meshlet_vertex_positions[lower_word_index] >> lower_word_bit_offset;\n        if lower_word_bit_offset + bits_per_channel[i] > 32u {\n            next_32_bits |= meshlet_vertex_positions[lower_word_index + 1u] << (32u - lower_word_bit_offset);\n        }\n        vertex_position_packed[i] = extractBits(next_32_bits, 0u, bits_per_channel[i]);\n        start_bit += bits_per_channel[i];\n    }\n\n    // Remap [0, range_max - range_min] vec3<u32> to [range_min, range_max] vec3<f32>\n    var vertex_position = vec3<f32>(vertex_position_packed) + vec3(\n        (*meshlet).min_vertex_position_channel_x,\n        (*meshlet).min_vertex_position_channel_y,\n        (*meshlet).min_vertex_position_channel_z,\n    );\n\n    // Reverse vertex quantization\n    let vertex_position_quantization_factor = unpacked.w;\n    vertex_position /= f32(1u << vertex_position_quantization_factor) * CENTIMETERS_PER_METER;\n\n    return vertex_position;\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS) {\n@group(2) @binding(0) var meshlet_visibility_buffer: texture_storage_2d<r64uint, read>;\n@group(2) @binding(1) var<storage, read> meshlet_raster_clusters: array<InstancedOffset>; // Per cluster\n@group(2) @binding(2) var<storage, read> meshlets: array<Meshlet>; // Per meshlet\n@group(2) @binding(3) var<storage, read> meshlet_indices: array<u32>; // Many per meshlet\n@group(2) @binding(4) var<storage, read> meshlet_vertex_positions: array<u32>; // Many per meshlet\n@group(2) @binding(5) var<storage, read> meshlet_vertex_normals: array<u32>; // Many per meshlet\n@group(2) @binding(6) var<storage, read> meshlet_vertex_uvs: array<vec2<f32>>; // Many per meshlet\n@group(2) @binding(7) var<storage, read> meshlet_instance_uniforms: array<Mesh>; // Per entity instance\n}\n\n// TODO: Load only twice, instead of 3x in cases where you load 3 indices per thread?\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn get_meshlet_vertex_id(index_id: u32) -> u32 {\n    let packed_index = meshlet_indices[index_id / 4u];\n    let bit_offset = (index_id % 4u) * 8u;\n    return extractBits(packed_index, bit_offset, 8u);\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn get_meshlet_vertex_position(meshlet: ptr<function, Meshlet>, vertex_id: u32) -> vec3<f32> {\n    // Get bitstream start for the vertex\n    let unpacked = unpack4xU8((*meshlet).packed_b);\n    let bits_per_channel = unpacked.xyz;\n    let bits_per_vertex = bits_per_channel.x + bits_per_channel.y + bits_per_channel.z;\n    var start_bit = (*meshlet).start_vertex_position_bit + (vertex_id * bits_per_vertex);\n\n    // Read each vertex channel from the bitstream\n    var vertex_position_packed = vec3(0u);\n    for (var i = 0u; i < 3u; i++) {\n        let lower_word_index = start_bit / 32u;\n        let lower_word_bit_offset = start_bit & 31u;\n        var next_32_bits = meshlet_vertex_positions[lower_word_index] >> lower_word_bit_offset;\n        if lower_word_bit_offset + bits_per_channel[i] > 32u {\n            next_32_bits |= meshlet_vertex_positions[lower_word_index + 1u] << (32u - lower_word_bit_offset);\n        }\n        vertex_position_packed[i] = extractBits(next_32_bits, 0u, bits_per_channel[i]);\n        start_bit += bits_per_channel[i];\n    }\n\n    // Remap [0, range_max - range_min] vec3<u32> to [range_min, range_max] vec3<f32>\n    var vertex_position = vec3<f32>(vertex_position_packed) + vec3(\n        (*meshlet).min_vertex_position_channel_x,\n        (*meshlet).min_vertex_position_channel_y,\n        (*meshlet).min_vertex_position_channel_z,\n    );\n\n    // Reverse vertex quantization\n    let vertex_position_quantization_factor = unpacked.w;\n    vertex_position /= f32(1u << vertex_position_quantization_factor) * CENTIMETERS_PER_METER;\n\n    return vertex_position;\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn get_meshlet_vertex_normal(meshlet: ptr<function, Meshlet>, vertex_id: u32) -> vec3<f32> {\n    let packed_normal = meshlet_vertex_normals[(*meshlet).start_vertex_attribute_id + vertex_id];\n    return octahedral_decode_signed(unpack2x16snorm(packed_normal));\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn get_meshlet_vertex_uv(meshlet: ptr<function, Meshlet>, vertex_id: u32) -> vec2<f32> {\n    return meshlet_vertex_uvs[(*meshlet).start_vertex_attribute_id + vertex_id];\n}\n");
    }
};
let handle:
        ::bevy_shader::_macro::bevy_asset::prelude::Handle<::bevy_shader::prelude::Shader> =
    {
        let (path, asset_server) =
            {
                let path =
                    {
                        {
                            let crate_name =
                                "bevy_pbr::meshlet".split(':').next().unwrap();
                            ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                                "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                                "bindings.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(app);
                (path, asset_server)
            };
        asset_server.load(path)
    };
::core::mem::forget(handle);load_shader_library!(app, "bindings.wesl");
141        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "visibility_buffer_resolve.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "visibility_buffer_resolve.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! Functions to be used by materials for reading from a meshlet visibility buffer texture.\n\nimport package::{\n    meshlet::bindings::{\n        Meshlet,\n        meshlet_visibility_buffer,\n        meshlet_raster_clusters,\n        meshlets,\n        meshlet_instance_uniforms,\n        get_meshlet_vertex_id,\n        get_meshlet_vertex_position,\n        get_meshlet_vertex_normal,\n        get_meshlet_vertex_uv,\n    },\n    render::{\n        mesh_view_bindings::view,\n        mesh_functions::mesh_position_local_to_world,\n        mesh_types::Mesh,\n        view_transformations::{position_world_to_clip, frag_coord_to_ndc},\n    },\n};\nimport bevy_render::maths::{affine3_to_square, mat2x4_f32_to_mat3x3_unpack};\n\n@if(PREPASS_FRAGMENT && MOTION_VECTOR_PREPASS)\nimport package::{\n    prepass::bindings::previous_view_uniforms,\n    render::pbr_prepass_functions::calculate_motion_vector,\n};\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nstruct PartialDerivatives {\n    barycentrics: vec3<f32>,\n    ddx: vec3<f32>,\n    ddy: vec3<f32>,\n}\n\n/// https://github.com/ConfettiFX/The-Forge/blob/9d43e69141a9cd0ce2ce2d2db5122234d3a2d5b5/Common_3/Renderer/VisibilityBuffer2/Shaders/FSL/vb_shading_utilities.h.fsl#L90-L150\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn compute_partial_derivatives(vertex_world_positions: array<vec4<f32>, 3>, ndc_uv: vec2<f32>, two_over_screen_size: vec2<f32>) -> PartialDerivatives {\n    var result: PartialDerivatives;\n\n    let vertex_clip_position_0 = position_world_to_clip(vertex_world_positions[0].xyz);\n    let vertex_clip_position_1 = position_world_to_clip(vertex_world_positions[1].xyz);\n    let vertex_clip_position_2 = position_world_to_clip(vertex_world_positions[2].xyz);\n\n    let inv_w = 1.0 / vec3(vertex_clip_position_0.w, vertex_clip_position_1.w, vertex_clip_position_2.w);\n    let ndc_0 = vertex_clip_position_0.xy * inv_w[0];\n    let ndc_1 = vertex_clip_position_1.xy * inv_w[1];\n    let ndc_2 = vertex_clip_position_2.xy * inv_w[2];\n\n    let inv_det = 1.0 / determinant(mat2x2(ndc_2 - ndc_1, ndc_0 - ndc_1));\n    result.ddx = vec3(ndc_1.y - ndc_2.y, ndc_2.y - ndc_0.y, ndc_0.y - ndc_1.y) * inv_det * inv_w;\n    result.ddy = vec3(ndc_2.x - ndc_1.x, ndc_0.x - ndc_2.x, ndc_1.x - ndc_0.x) * inv_det * inv_w;\n\n    var ddx_sum = dot(result.ddx, vec3(1.0));\n    var ddy_sum = dot(result.ddy, vec3(1.0));\n\n    let delta_v = ndc_uv - ndc_0;\n    let interp_inv_w = inv_w.x + delta_v.x * ddx_sum + delta_v.y * ddy_sum;\n    let interp_w = 1.0 / interp_inv_w;\n\n    result.barycentrics = vec3(\n        interp_w * (inv_w[0] + delta_v.x * result.ddx.x + delta_v.y * result.ddy.x),\n        interp_w * (delta_v.x * result.ddx.y + delta_v.y * result.ddy.y),\n        interp_w * (delta_v.x * result.ddx.z + delta_v.y * result.ddy.z),\n    );\n\n    result.ddx *= two_over_screen_size.x;\n    result.ddy *= two_over_screen_size.y;\n    ddx_sum *= two_over_screen_size.x;\n    ddy_sum *= two_over_screen_size.y;\n\n    result.ddy *= -1.0;\n    ddy_sum *= -1.0;\n\n    let interp_ddx_w = 1.0 / (interp_inv_w + ddx_sum);\n    let interp_ddy_w = 1.0 / (interp_inv_w + ddy_sum);\n\n    result.ddx = interp_ddx_w * (result.barycentrics * interp_inv_w + result.ddx) - result.barycentrics;\n    result.ddy = interp_ddy_w * (result.barycentrics * interp_inv_w + result.ddy) - result.barycentrics;\n    return result;\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nstruct VertexOutput {\n    position: vec4<f32>,\n    world_position: vec4<f32>,\n    world_normal: vec3<f32>,\n    uv: vec2<f32>,\n    ddx_uv: vec2<f32>,\n    ddy_uv: vec2<f32>,\n    world_tangent: vec4<f32>,\n    mesh_flags: u32,\n    cluster_id: u32,\n    material_bind_group_slot: u32,\n    @if(PREPASS_FRAGMENT && MOTION_VECTOR_PREPASS)\n    motion_vector: vec2<f32>,\n}\n\n/// Load the visibility buffer texture and resolve it into a VertexOutput.\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn resolve_vertex_output(frag_coord: vec4<f32>) -> VertexOutput {\n    let packed_ids = u32(textureLoad(meshlet_visibility_buffer, vec2<u32>(frag_coord.xy)).r);\n    let cluster_id = packed_ids >> 7u;\n    let instanced_offset = meshlet_raster_clusters[cluster_id];\n    let meshlet_id = instanced_offset.offset;\n    var meshlet = meshlets[meshlet_id];\n\n    let triangle_id = extractBits(packed_ids, 0u, 7u);\n    let index_ids = meshlet.start_index_id + (triangle_id * 3u) + vec3(0u, 1u, 2u);\n    let vertex_ids = vec3(get_meshlet_vertex_id(index_ids[0]), get_meshlet_vertex_id(index_ids[1]), get_meshlet_vertex_id(index_ids[2]));\n    let vertex_0 = load_vertex(&meshlet, vertex_ids[0]);\n    let vertex_1 = load_vertex(&meshlet, vertex_ids[1]);\n    let vertex_2 = load_vertex(&meshlet, vertex_ids[2]);\n\n    let instance_id = instanced_offset.instance_id;\n    var instance_uniform = meshlet_instance_uniforms[instance_id];\n\n    let world_from_local = affine3_to_square(instance_uniform.world_from_local);\n    let world_position_0 = mesh_position_local_to_world(world_from_local, vec4(vertex_0.position, 1.0));\n    let world_position_1 = mesh_position_local_to_world(world_from_local, vec4(vertex_1.position, 1.0));\n    let world_position_2 = mesh_position_local_to_world(world_from_local, vec4(vertex_2.position, 1.0));\n\n    let frag_coord_ndc = frag_coord_to_ndc(frag_coord).xy;\n    let partial_derivatives = compute_partial_derivatives(\n        array(world_position_0, world_position_1, world_position_2),\n        frag_coord_ndc,\n        2.0 / view.viewport.zw,\n    );\n\n    let world_position = mat3x4(world_position_0, world_position_1, world_position_2) * partial_derivatives.barycentrics;\n    let world_positions_camera_relative = mat3x3(\n        world_position_0.xyz - view.world_position,\n        world_position_1.xyz - view.world_position,\n        world_position_2.xyz - view.world_position,\n    );\n    let ddx_world_position = world_positions_camera_relative * partial_derivatives.ddx;\n    let ddy_world_position = world_positions_camera_relative * partial_derivatives.ddy;\n\n    let world_normal = mat3x3(\n        normal_local_to_world(vertex_0.normal, &instance_uniform),\n        normal_local_to_world(vertex_1.normal, &instance_uniform),\n        normal_local_to_world(vertex_2.normal, &instance_uniform),\n    ) * partial_derivatives.barycentrics;\n\n    let uv = mat3x2(vertex_0.uv, vertex_1.uv, vertex_2.uv) * partial_derivatives.barycentrics;\n    let ddx_uv = mat3x2(vertex_0.uv, vertex_1.uv, vertex_2.uv) * partial_derivatives.ddx;\n    let ddy_uv = mat3x2(vertex_0.uv, vertex_1.uv, vertex_2.uv) * partial_derivatives.ddy;\n\n    let world_tangent = calculate_world_tangent(world_normal, ddx_world_position, ddy_world_position, ddx_uv, ddy_uv);\n\n    var motion_vector: vec2<f32>;\n@if(PREPASS_FRAGMENT && MOTION_VECTOR_PREPASS) {\n    let previous_world_from_local = affine3_to_square(instance_uniform.previous_world_from_local);\n    let previous_world_position_0 = mesh_position_local_to_world(previous_world_from_local, vec4(vertex_0.position, 1.0));\n    let previous_world_position_1 = mesh_position_local_to_world(previous_world_from_local, vec4(vertex_1.position, 1.0));\n    let previous_world_position_2 = mesh_position_local_to_world(previous_world_from_local, vec4(vertex_2.position, 1.0));\n    let previous_world_position = mat3x4(previous_world_position_0, previous_world_position_1, previous_world_position_2) * partial_derivatives.barycentrics;\n    motion_vector = calculate_motion_vector(world_position, previous_world_position);\n}\n\n    var out: VertexOutput;\n    out.position = frag_coord;\n    out.world_position = world_position;\n    out.world_normal = world_normal;\n    out.uv = uv;\n    out.ddx_uv = ddx_uv;\n    out.ddy_uv = ddy_uv;\n    out.world_tangent = world_tangent;\n    out.mesh_flags = instance_uniform.flags;\n    out.cluster_id = instance_id ^ meshlet_id;\n    out.material_bind_group_slot = instance_uniform.material_and_lightmap_bind_group_slot & 0xffffu;\n    @if(PREPASS_FRAGMENT && MOTION_VECTOR_PREPASS)\n    out.motion_vector = motion_vector;\n    return out;\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nstruct MeshletVertex {\n    position: vec3<f32>,\n    normal: vec3<f32>,\n    uv: vec2<f32>,\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn load_vertex(meshlet: ptr<function, Meshlet>, vertex_id: u32) -> MeshletVertex {\n    return MeshletVertex(\n        get_meshlet_vertex_position(meshlet, vertex_id),\n        get_meshlet_vertex_normal(meshlet, vertex_id),\n        get_meshlet_vertex_uv(meshlet, vertex_id),\n    );\n}\n\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn normal_local_to_world(vertex_normal: vec3<f32>, instance_uniform: ptr<function, Mesh>) -> vec3<f32> {\n    if any(vertex_normal != vec3<f32>(0.0)) {\n        return normalize(\n            mat2x4_f32_to_mat3x3_unpack(\n                (*instance_uniform).local_from_world_transpose_a,\n                (*instance_uniform).local_from_world_transpose_b,\n            ) * vertex_normal\n        );\n    } else {\n        return vertex_normal;\n    }\n}\n\n/// https://www.jeremyong.com/graphics/2023/12/16/surface-gradient-bump-mapping/#surface-gradient-from-a-tangent-space-normal-vector-without-an-explicit-tangent-basis\n@if(MESHLET_MESH_MATERIAL_PASS)\nfn calculate_world_tangent(\n    world_normal: vec3<f32>,\n    ddx_world_position: vec3<f32>,\n    ddy_world_position: vec3<f32>,\n    ddx_uv: vec2<f32>,\n    ddy_uv: vec2<f32>,\n) -> vec4<f32> {\n    // Project the position gradients onto the tangent plane\n    let ddx_world_position_s = ddx_world_position - dot(ddx_world_position, world_normal) * world_normal;\n    let ddy_world_position_s = ddy_world_position - dot(ddy_world_position, world_normal) * world_normal;\n\n    // Compute the jacobian matrix to leverage the chain rule\n    let jacobian_sign = sign(ddx_uv.x * ddy_uv.y - ddx_uv.y * ddy_uv.x);\n\n    var world_tangent = jacobian_sign * (ddy_uv.y * ddx_world_position_s - ddx_uv.y * ddy_world_position_s);\n\n    // The sign intrinsic returns 0 if the argument is 0\n    if jacobian_sign != 0.0 {\n        world_tangent = normalize(world_tangent);\n    }\n\n    // The second factor here ensures a consistent handedness between\n    // the tangent frame and surface basis w.r.t. screenspace.\n    let w = jacobian_sign * sign(dot(ddy_world_position, cross(world_normal, ddx_world_position)));\n\n    return vec4(world_tangent, -w); // TODO: Unclear why we need to negate this to match mikktspace generated tangents\n}\n");
    }
};
let handle:
        ::bevy_shader::_macro::bevy_asset::prelude::Handle<::bevy_shader::prelude::Shader> =
    {
        let (path, asset_server) =
            {
                let path =
                    {
                        {
                            let crate_name =
                                "bevy_pbr::meshlet".split(':').next().unwrap();
                            ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                                "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                                "visibility_buffer_resolve.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(app);
                (path, asset_server)
            };
        asset_server.load(path)
    };
::core::mem::forget(handle);load_shader_library!(app, "visibility_buffer_resolve.wesl");
142        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "cull_shared.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "cull_shared.wesl");
        embedded.insert_asset(watched_path, &path,
            b"import package::meshlet::bindings::{\n    MeshletAabb,\n    DispatchIndirectArgs,\n    InstancedOffset,\n    depth_pyramid,\n    view,\n    previous_view,\n    meshlet_instance_uniforms,\n};\nimport bevy_render::maths::affine3_to_square;\n\n/// https://github.com/zeux/meshoptimizer/blob/1e48e96c7e8059321de492865165e9ef071bffba/demo/nanite.cpp#L115\nfn lod_error_is_imperceptible(lod_sphere: vec4<f32>, simplification_error: f32, instance_id: u32) -> bool {\n    let world_from_local = affine3_to_square(meshlet_instance_uniforms[instance_id].world_from_local);\n    let world_scale = max(length(world_from_local[0]), max(length(world_from_local[1]), length(world_from_local[2])));\n    let camera_pos = view.world_position;\n\n    let projection = view.clip_from_view;\n    if projection[3][3] == 1.0 {\n        // Orthographic\n        let world_error = simplification_error * world_scale;\n        let proj = projection[1][1];\n        let height = 2.0 / proj;\n        let norm_error = world_error / height;\n        return norm_error * view.viewport.w < 1.0;\n    } else {\n        // Perspective\n        var near = projection[3][2];\n        let world_sphere_center = (world_from_local * vec4<f32>(lod_sphere.xyz, 1.0)).xyz;\n        let world_sphere_radius = lod_sphere.w * world_scale;\n        let d_pos = world_sphere_center - camera_pos;\n        let d = sqrt(dot(d_pos, d_pos)) - world_sphere_radius;\n        let norm_error = simplification_error / max(d, near) * projection[1][1] * 0.5;\n        return norm_error * view.viewport.w < 1.0;\n    }\n}\n\nfn normalize_plane(p: vec4<f32>) -> vec4<f32> {\n    return p / length(p.xyz);\n}\n\n/// https://fgiesen.wordpress.com/2012/08/31/frustum-planes-from-the-projection-matrix/\n/// https://fgiesen.wordpress.com/2010/10/17/view-frustum-culling/\nfn aabb_in_frustum(aabb: MeshletAabb, instance_id: u32) -> bool {\n    let world_from_local = affine3_to_square(meshlet_instance_uniforms[instance_id].world_from_local);\n    let clip_from_local = view.clip_from_world * world_from_local;\n    let row_major = transpose(clip_from_local);\n    let planes = array(\n        row_major[3] + row_major[0],\n        row_major[3] - row_major[0],\n        row_major[3] + row_major[1],\n        row_major[3] - row_major[1],\n        row_major[2],\n    );\n\n    for (var i = 0; i < 5; i++) {\n        let plane = normalize_plane(planes[i]);\n        let flipped = aabb.half_extent * sign(plane.xyz);\n        if dot(aabb.center + flipped, plane.xyz) <= -plane.w {\n            return false;\n        }\n    }\n    return true;\n}\n\nstruct ScreenAabb {\n    min: vec3<f32>,\n    max: vec3<f32>,\n}\n\nfn min8(a: vec3<f32>, b: vec3<f32>, c: vec3<f32>, d: vec3<f32>, e: vec3<f32>, f: vec3<f32>, g: vec3<f32>, h: vec3<f32>) -> vec3<f32> {\n    return min(min(min(a, b), min(c, d)), min(min(e, f), min(g, h)));\n}\n\nfn max8(a: vec3<f32>, b: vec3<f32>, c: vec3<f32>, d: vec3<f32>, e: vec3<f32>, f: vec3<f32>, g: vec3<f32>, h: vec3<f32>) -> vec3<f32> {\n    return max(max(max(a, b), max(c, d)), max(max(e, f), max(g, h)));\n}\n\nfn min8_4(a: vec4<f32>, b: vec4<f32>, c: vec4<f32>, d: vec4<f32>, e: vec4<f32>, f: vec4<f32>, g: vec4<f32>, h: vec4<f32>) -> vec4<f32> {\n    return min(min(min(a, b), min(c, d)), min(min(e, f), min(g, h)));\n}\n\n/// https://zeux.io/2023/01/12/approximate-projected-bounds/\nfn project_aabb(clip_from_local: mat4x4<f32>, near: f32, aabb: MeshletAabb, out: ptr<function, ScreenAabb>) -> bool {\n    let extent = aabb.half_extent * 2.0;\n    let sx = clip_from_local * vec4<f32>(extent.x, 0.0, 0.0, 0.0);\n    let sy = clip_from_local * vec4<f32>(0.0, extent.y, 0.0, 0.0);\n    let sz = clip_from_local * vec4<f32>(0.0, 0.0, extent.z, 0.0);\n\n    let p0 = clip_from_local * vec4<f32>(aabb.center - aabb.half_extent, 1.0);\n    let p1 = p0 + sz;\n    let p2 = p0 + sy;\n    let p3 = p2 + sz;\n    let p4 = p0 + sx;\n    let p5 = p4 + sz;\n    let p6 = p4 + sy;\n    let p7 = p6 + sz;\n\n    let depth = min8_4(p0, p1, p2, p3, p4, p5, p6, p7).w;\n    // do not occlusion cull if we are inside the aabb\n    if depth < near {\n        return false;\n    }\n\n    let dp0 = p0.xyz / p0.w;\n    let dp1 = p1.xyz / p1.w;\n    let dp2 = p2.xyz / p2.w;\n    let dp3 = p3.xyz / p3.w;\n    let dp4 = p4.xyz / p4.w;\n    let dp5 = p5.xyz / p5.w;\n    let dp6 = p6.xyz / p6.w;\n    let dp7 = p7.xyz / p7.w;\n    let min = min8(dp0, dp1, dp2, dp3, dp4, dp5, dp6, dp7);\n    let max = max8(dp0, dp1, dp2, dp3, dp4, dp5, dp6, dp7);\n    var vaabb = vec4<f32>(min.xy, max.xy);\n    // convert ndc to texture coordinates by rescaling and flipping Y\n    vaabb = vaabb.xwzy * vec4<f32>(0.5, -0.5, 0.5, -0.5) + 0.5;\n    (*out).min = vec3<f32>(vaabb.xy, min.z);\n    (*out).max = vec3<f32>(vaabb.zw, max.z);\n    return true;\n}\n\nfn sample_hzb(smin: vec2<u32>, smax: vec2<u32>, mip: i32) -> f32 {\n    let texel = vec4<u32>(0, 1, 2, 3);\n    let sx = min(smin.x + texel, smax.xxxx);\n    let sy = min(smin.y + texel, smax.yyyy);\n    // TODO: switch to min samplers when wgpu has them\n    // sampling 16 times a finer mip is worth the extra cost for better culling\n    let a = sample_hzb_row(sx, sy.x, mip);\n    let b = sample_hzb_row(sx, sy.y, mip);\n    let c = sample_hzb_row(sx, sy.z, mip);\n    let d = sample_hzb_row(sx, sy.w, mip);\n    return min(min(a, b), min(c, d));\n}\n\nfn sample_hzb_row(sx: vec4<u32>, sy: u32, mip: i32) -> f32 {\n    let a = textureLoad(depth_pyramid, vec2(sx.x, sy), mip).x;\n    let b = textureLoad(depth_pyramid, vec2(sx.y, sy), mip).x;\n    let c = textureLoad(depth_pyramid, vec2(sx.z, sy), mip).x;\n    let d = textureLoad(depth_pyramid, vec2(sx.w, sy), mip).x;\n    return min(min(a, b), min(c, d));\n}\n\nfn occlusion_cull_screen_aabb(aabb: ScreenAabb, screen: vec2<f32>) -> bool {\n    let hzb_size = vec2<f32>(textureDimensions(depth_pyramid).xy);\n    let aabb_min = aabb.min.xy * hzb_size;\n    let aabb_max = aabb.max.xy * hzb_size;\n\n    let min_texel = vec2<u32>(max(aabb_min, vec2<f32>(0.0)));\n    let max_texel = vec2<u32>(min(aabb_max, hzb_size - 1.0));\n    let size = max_texel - min_texel;\n    let max_size = max(size.x, size.y);\n\n    // note: add 1 before max because the unsigned overflow behavior is intentional\n    // it wraps around firstLeadingBit(0) = ~0 to 0\n    // TODO: we actually sample a 4x4 block, so ideally this would be `max(..., 3u) - 3u`.\n    var mip = max(firstLeadingBit(max_size) + 1u, 2u) - 2u;\n\n    if any((max_texel >> vec2(mip)) > (min_texel >> vec2(mip)) + 3) {\n        mip += 1u;\n    }\n\n    let smin = min_texel >> vec2<u32>(mip);\n    let smax = max_texel >> vec2<u32>(mip);\n\n    let curr_depth = sample_hzb(smin, smax, i32(mip));\n    return aabb.max.z <= curr_depth;\n}\n\nfn occlusion_cull_projection() -> mat4x4<f32> {\n    @if(MESHLET_FIRST_CULLING_PASS)\n    return previous_view.clip_from_world;\n    @else\n    return view.clip_from_world;\n}\n\nfn occlusion_cull_clip_from_local(instance_id: u32) -> mat4x4<f32> {\n    @if(MESHLET_FIRST_CULLING_PASS)\n    let prev_world_from_local = affine3_to_square(meshlet_instance_uniforms[instance_id].previous_world_from_local);\n    @if(MESHLET_FIRST_CULLING_PASS)\n    return previous_view.clip_from_world * prev_world_from_local;\n    @if(!MESHLET_FIRST_CULLING_PASS)\n    let world_from_local = affine3_to_square(meshlet_instance_uniforms[instance_id].world_from_local);\n    @if(!MESHLET_FIRST_CULLING_PASS)\n    return view.clip_from_world * world_from_local;\n}\n\nfn should_occlusion_cull_aabb(aabb: MeshletAabb, instance_id: u32) -> bool {\n    let projection = occlusion_cull_projection();\n    var near: f32;\n    if projection[3][3] == 1.0 {\n        near = projection[3][2] / projection[2][2];\n    } else {\n        near = projection[3][2];\n    }\n\n    let clip_from_local = occlusion_cull_clip_from_local(instance_id);\n    var screen_aabb = ScreenAabb(vec3<f32>(0.0), vec3<f32>(0.0));\n    if project_aabb(clip_from_local, near, aabb, &screen_aabb) {\n        return occlusion_cull_screen_aabb(screen_aabb, view.viewport.zw);\n    }\n    return false;\n}\n");
    }
};
let handle:
        ::bevy_shader::_macro::bevy_asset::prelude::Handle<::bevy_shader::prelude::Shader> =
    {
        let (path, asset_server) =
            {
                let path =
                    {
                        {
                            let crate_name =
                                "bevy_pbr::meshlet".split(':').next().unwrap();
                            ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                                "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                                "cull_shared.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(app);
                (path, asset_server)
            };
        asset_server.load(path)
    };
::core::mem::forget(handle);load_shader_library!(app, "cull_shared.wesl");
143        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "clear_visibility_buffer.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "clear_visibility_buffer.wesl");
        embedded.insert_asset(watched_path, &path,
            b"@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n@group(0) @binding(0) var meshlet_visibility_buffer: texture_storage_2d<r64uint, write>;\n@else\n@group(0) @binding(0) var meshlet_visibility_buffer: texture_storage_2d<r32uint, write>;\nvar<immediate> view_size: vec2<u32>;\n\n@compute\n@workgroup_size(16, 16, 1)\nfn clear_visibility_buffer(@builtin(global_invocation_id) global_id: vec3<u32>) {\n    if any(global_id.xy >= view_size) { return; }\n\n    @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n    textureStore(meshlet_visibility_buffer, global_id.xy, vec4(0lu));\n    @else\n    textureStore(meshlet_visibility_buffer, global_id.xy, vec4(0u));\n}\n");
    }
};embedded_asset!(app, "clear_visibility_buffer.wesl");
144        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "cull_instances.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "cull_instances.wesl");
        embedded.insert_asset(watched_path, &path,
            b"import package::meshlet::bindings::{\n    InstancedOffset,\n    constants,\n    meshlet_view_instance_visibility,\n    meshlet_instance_aabbs,\n    meshlet_instance_bvh_root_nodes,\n    meshlet_bvh_cull_count_write,\n    meshlet_bvh_cull_dispatch,\n    meshlet_bvh_cull_queue,\n    meshlet_second_pass_instance_count,\n    meshlet_second_pass_instance_dispatch,\n    meshlet_second_pass_instance_candidates,\n};\nimport package::meshlet::cull_shared::{\n    aabb_in_frustum,\n    should_occlusion_cull_aabb,\n};\n\nfn instance_count() -> u32 {\n    @if(MESHLET_FIRST_CULLING_PASS)\n    return constants.scene_instance_count;\n    @else\n    return meshlet_second_pass_instance_count;\n}\n\nfn map_instance_id(id: u32) -> u32 {\n    @if(MESHLET_FIRST_CULLING_PASS)\n    return id;\n    @else\n    return meshlet_second_pass_instance_candidates[id];\n}\n\nfn should_cull_instance(instance_id: u32) -> bool {\n    let index = instance_id >> 5u;\n    let bit_offset = instance_id & 31u;\n    let packed_visibility = meshlet_view_instance_visibility[index];\n    return bool(extractBits(packed_visibility, bit_offset, 1u));\n}\n\n@compute\n@workgroup_size(128, 1, 1) // 1 instance per thread\nfn cull_instances(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {\n    // Calculate the instance ID for this thread\n    let dispatch_id = global_invocation_id.x;\n    if dispatch_id >= instance_count() { return; }\n\n    let instance_id = map_instance_id(dispatch_id);\n    let aabb = meshlet_instance_aabbs[instance_id];\n\n    // Visibility and frustum cull, but only in the first pass\n    @if(MESHLET_FIRST_CULLING_PASS)\n    if should_cull_instance(instance_id) || !aabb_in_frustum(aabb, instance_id) { return; }\n\n    // If we pass, try occlusion culling\n    // If this instance was occluded, push it to the second pass to check against this frame\'s HZB\n    if should_occlusion_cull_aabb(aabb, instance_id) {\n@if(MESHLET_FIRST_CULLING_PASS) {\n        let id = atomicAdd(&meshlet_second_pass_instance_count, 1u);\n        meshlet_second_pass_instance_candidates[id] = instance_id;\n        if ((id & 127u) == 0u) {\n            atomicAdd(&meshlet_second_pass_instance_dispatch.x, 1u);\n        }\n}\n        return;\n    }\n\n    // If we pass, push the instance\'s root node to BVH cull\n    let root_node = meshlet_instance_bvh_root_nodes[instance_id];\n    let id = atomicAdd(&meshlet_bvh_cull_count_write, 1u);\n    meshlet_bvh_cull_queue[id] = InstancedOffset(instance_id, root_node);\n    if ((id & 15u) == 0u) {\n        atomicAdd(&meshlet_bvh_cull_dispatch.x, 1u);\n    }\n}\n");
    }
};embedded_asset!(app, "cull_instances.wesl");
145        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "cull_bvh.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "cull_bvh.wesl");
        embedded.insert_asset(watched_path, &path,
            b"import package::meshlet::bindings::{\n    InstancedOffset,\n    get_aabb,\n    get_aabb_error,\n    get_aabb_child_offset,\n    constants,\n    meshlet_bvh_nodes,\n    meshlet_bvh_cull_count_read,\n    meshlet_bvh_cull_count_write,\n    meshlet_bvh_cull_dispatch,\n    meshlet_bvh_cull_queue,\n    meshlet_meshlet_cull_count_early,\n    meshlet_meshlet_cull_count_late,\n    meshlet_meshlet_cull_dispatch_early,\n    meshlet_meshlet_cull_dispatch_late,\n    meshlet_meshlet_cull_queue,\n    meshlet_second_pass_bvh_count,\n    meshlet_second_pass_bvh_dispatch,\n    meshlet_second_pass_bvh_queue,\n};\nimport package::meshlet::cull_shared::{\n    lod_error_is_imperceptible,\n    aabb_in_frustum,\n    should_occlusion_cull_aabb,\n};\n\n@compute\n@workgroup_size(128, 1, 1) // 8 threads per node, 16 nodes per workgroup\nfn cull_bvh(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {\n    // Calculate the queue ID for this thread\n    let dispatch_id = global_invocation_id.x;\n    var node = dispatch_id >> 3u;\n    let subnode = dispatch_id & 7u;\n    if node >= meshlet_bvh_cull_count_read { return; }\n\n    node = select(node, constants.rightmost_slot - node, constants.read_from_front == 0u);\n    let instanced_offset = meshlet_bvh_cull_queue[node];\n    let instance_id = instanced_offset.instance_id;\n    let bvh_node = &meshlet_bvh_nodes[instanced_offset.offset];\n\n    var aabb_error_offset = (*bvh_node).aabbs[subnode];\n    let aabb = get_aabb(&aabb_error_offset);\n    let parent_error = get_aabb_error(&aabb_error_offset);\n    let lod_sphere = (*bvh_node).lod_bounds[subnode];\n\n    let parent_is_imperceptible = lod_error_is_imperceptible(lod_sphere, parent_error, instance_id);\n    // Error and frustum cull, in both passes\n    if parent_is_imperceptible || !aabb_in_frustum(aabb, instance_id) { return; }\n\n    let child_offset = get_aabb_child_offset(&aabb_error_offset);\n    let index = subnode >> 2u;\n    let bit_offset = subnode & 3u;\n    let packed_child_count = (*bvh_node).child_counts[index];\n    let child_count = extractBits(packed_child_count, bit_offset * 8u, 8u);\n    var value = InstancedOffset(instance_id, child_offset);\n\n    // If we pass, try occlusion culling\n    // If this node was occluded, push it\'s children to the second pass to check against this frame\'s HZB\n    if should_occlusion_cull_aabb(aabb, instance_id) {\n        @if(MESHLET_FIRST_CULLING_PASS)\n        if child_count == 255u {\n            let id = atomicAdd(&meshlet_second_pass_bvh_count, 1u);\n            meshlet_second_pass_bvh_queue[id] = value;\n            if ((id & 15u) == 0u) {\n                atomicAdd(&meshlet_second_pass_bvh_dispatch.x, 1u);\n            }\n        } else {\n            let base = atomicAdd(&meshlet_meshlet_cull_count_late, child_count);\n            let start = constants.rightmost_slot - base;\n            for (var i = start; i > start - child_count; i--) {\n                meshlet_meshlet_cull_queue[i] = value;\n                value.offset += 1u;\n            }\n            let req = (base + child_count + 127u) >> 7u;\n            atomicMax(&meshlet_meshlet_cull_dispatch_late.x, req);\n        }\n        return;\n    }\n\n    // If we pass, push the children to the next BVH cull\n    if child_count == 255u {\n        let id = atomicAdd(&meshlet_bvh_cull_count_write, 1u);\n        let index = select(constants.rightmost_slot - id, id, constants.read_from_front == 0u);\n        meshlet_bvh_cull_queue[index] = value;\n        if ((id & 15u) == 0u) {\n            atomicAdd(&meshlet_bvh_cull_dispatch.x, 1u);\n        }\n    } else {\n@if(MESHLET_FIRST_CULLING_PASS) {\n        let base = atomicAdd(&meshlet_meshlet_cull_count_early, child_count);\n        let end = base + child_count;\n        for (var i = base; i < end; i++) {\n            meshlet_meshlet_cull_queue[i] = value;\n            value.offset += 1u;\n        }\n        let req = (end + 127u) >> 7u;\n        atomicMax(&meshlet_meshlet_cull_dispatch_early.x, req);\n} @else {\n        let base = atomicAdd(&meshlet_meshlet_cull_count_late, child_count);\n        let start = constants.rightmost_slot - base;\n        for (var i = start; i > start - child_count; i--) {\n            meshlet_meshlet_cull_queue[i] = value;\n            value.offset += 1u;\n        }\n        let req = (base + child_count + 127u) >> 7u;\n        atomicMax(&meshlet_meshlet_cull_dispatch_late.x, req);\n}\n    }\n}\n");
    }
};embedded_asset!(app, "cull_bvh.wesl");
146        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "cull_clusters.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "cull_clusters.wesl");
        embedded.insert_asset(watched_path, &path,
            b"import package::meshlet::bindings::{\n    InstancedOffset,\n    get_aabb,\n    get_aabb_error,\n    constants,\n    view,\n    meshlet_instance_uniforms,\n    meshlet_cull_data,\n    meshlet_software_raster_indirect_args,\n    meshlet_hardware_raster_indirect_args,\n    meshlet_previous_raster_counts,\n    meshlet_raster_clusters,\n    meshlet_meshlet_cull_count_read,\n    meshlet_meshlet_cull_count_write,\n    meshlet_meshlet_cull_dispatch,\n    meshlet_meshlet_cull_queue,\n};\nimport package::meshlet::cull_shared::{\n    ScreenAabb,\n    project_aabb,\n    lod_error_is_imperceptible,\n    aabb_in_frustum,\n    should_occlusion_cull_aabb,\n};\nimport bevy_render::maths::affine3_to_square;\n\n@compute\n@workgroup_size(128, 1, 1) // 1 cluster per thread\nfn cull_clusters(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {\n    if global_invocation_id.x >= meshlet_meshlet_cull_count_read { return; }\n\n    @if(MESHLET_FIRST_CULLING_PASS)\n    let meshlet_id = global_invocation_id.x;\n    @else\n    let meshlet_id = constants.rightmost_slot - global_invocation_id.x;\n    let instanced_offset = meshlet_meshlet_cull_queue[meshlet_id];\n    let instance_id = instanced_offset.instance_id;\n    let cull_data = &meshlet_cull_data[instanced_offset.offset];\n    var aabb_error_offset = (*cull_data).aabb;\n    let aabb = get_aabb(&aabb_error_offset);\n    let error = get_aabb_error(&aabb_error_offset);\n    let lod_sphere = (*cull_data).lod_group_sphere;\n\n    let is_imperceptible = lod_error_is_imperceptible(lod_sphere, error, instance_id);\n    // Error and frustum cull, in both passes\n    if !is_imperceptible || !aabb_in_frustum(aabb, instance_id) { return; }\n\n    // If we pass, try occlusion culling\n    // If this node was occluded, push it\'s children to the second pass to check against this frame\'s HZB\n    if should_occlusion_cull_aabb(aabb, instance_id) {\n@if(MESHLET_FIRST_CULLING_PASS) {\n        let id = atomicAdd(&meshlet_meshlet_cull_count_write, 1u);\n        let value = InstancedOffset(instance_id, instanced_offset.offset);\n        meshlet_meshlet_cull_queue[constants.rightmost_slot - id] = value;\n        if ((id & 127u) == 0) {\n            atomicAdd(&meshlet_meshlet_cull_dispatch.x, 1u);\n        }\n}\n        return;\n    }\n\n    // If we pass, rasterize the meshlet\n    // Check how big the cluster is in screen space\n    let world_from_local = affine3_to_square(meshlet_instance_uniforms[instance_id].world_from_local);\n    let clip_from_local  = view.clip_from_world * world_from_local;\n    let projection = view.clip_from_world;\n    var near: f32;\n    if projection[3][3] == 1.0 {\n        near = projection[3][2] / projection[2][2];\n    } else {\n        near = projection[3][2];\n    }\n    var screen_aabb = ScreenAabb(vec3<f32>(0.0), vec3<f32>(0.0));\n    var sw_raster = project_aabb(clip_from_local, near, aabb, &screen_aabb);\n    if sw_raster {\n        let aabb_size = (screen_aabb.max.xy - screen_aabb.min.xy) * view.viewport.zw;\n        sw_raster = all(aabb_size <= vec2<f32>(64.0));\n    }\n\n    var buffer_slot: u32;\n    if sw_raster {\n        // Append this cluster to the list for software rasterization\n        buffer_slot = atomicAdd(&meshlet_software_raster_indirect_args.x, 1u);\n        buffer_slot += meshlet_previous_raster_counts[0];\n    } else {\n        // Append this cluster to the list for hardware rasterization\n        buffer_slot = atomicAdd(&meshlet_hardware_raster_indirect_args.instance_count, 1u);\n        buffer_slot += meshlet_previous_raster_counts[1];\n        buffer_slot = constants.rightmost_slot - buffer_slot;\n    }\n    meshlet_raster_clusters[buffer_slot] = InstancedOffset(instance_id, instanced_offset.offset);\n}\n");
    }
};embedded_asset!(app, "cull_clusters.wesl");
147        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "visibility_buffer_software_raster.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "visibility_buffer_software_raster.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! Compute shader for rasterizing small clusters into a visibility buffer.\n\nimport package::{\n    meshlet::bindings::{\n        meshlet_cluster_meshlet_ids,\n        meshlets,\n        meshlet_cluster_instance_ids,\n        meshlet_instance_uniforms,\n        meshlet_raster_clusters,\n        meshlet_previous_raster_counts,\n        meshlet_software_raster_cluster_count,\n        meshlet_visibility_buffer,\n        view,\n        get_meshlet_vertex_count,\n        get_meshlet_triangle_count,\n        get_meshlet_vertex_id,\n        get_meshlet_vertex_position,\n    },\n    render::{\n        mesh_functions::mesh_position_local_to_world,\n        view_transformations::ndc_to_uv,\n    },\n};\nimport bevy_render::maths::affine3_to_square;\n\n// TODO: Fixed-point math and top-left rule\n\nvar<workgroup> viewport_vertices: array<vec3f, 256>;\n\n@compute\n@workgroup_size(128, 1, 1) // 128 threads per workgroup, 1-2 vertices per thread, 1 triangle per thread, 1 cluster per workgroup\nfn rasterize_cluster(\n    @builtin(workgroup_id) workgroup_id: vec3<u32>,\n    @builtin(local_invocation_index) local_invocation_index: u32,\n    @if(MESHLET_2D_DISPATCH)\n    @builtin(num_workgroups) num_workgroups: vec3<u32>,\n) {\n    var workgroup_id_1d = workgroup_id.x;\n\n@if(MESHLET_2D_DISPATCH) {\n    workgroup_id_1d += workgroup_id.y * num_workgroups.x;\n    if workgroup_id_1d >= meshlet_software_raster_cluster_count { return; }\n}\n\n    let cluster_id = workgroup_id_1d + meshlet_previous_raster_counts[0];\n    let instanced_offset = meshlet_raster_clusters[cluster_id];\n    var meshlet = meshlets[instanced_offset.offset];\n\n    let instance_uniform = meshlet_instance_uniforms[instanced_offset.instance_id];\n    let world_from_local = affine3_to_square(instance_uniform.world_from_local);\n\n    // Load and project 1 vertex per thread, and then again if there are more than 128 vertices in the meshlet\n    for (var i = 0u; i <= 128u; i += 128u) {\n        let vertex_id = local_invocation_index + i;\n        if vertex_id < get_meshlet_vertex_count(&meshlet) {\n            let vertex_position = get_meshlet_vertex_position(&meshlet, vertex_id);\n\n            // Project vertex to viewport space\n            let world_position = mesh_position_local_to_world(world_from_local, vec4(vertex_position, 1.0));\n            let clip_position = view.clip_from_world * vec4(world_position.xyz, 1.0);\n            let ndc_position = clip_position.xyz / clip_position.w;\n            let viewport_position_xy = ndc_to_uv(ndc_position.xy) * view.viewport.zw;\n\n            // Write vertex to workgroup shared memory\n            viewport_vertices[vertex_id] = vec3(viewport_position_xy, ndc_position.z);\n        }\n    }\n    workgroupBarrier();\n\n    // Load 1 triangle\'s worth of vertex data per thread\n    let triangle_id = local_invocation_index;\n    if triangle_id >= get_meshlet_triangle_count(&meshlet) { return; }\n    let index_ids = meshlet.start_index_id + (triangle_id * 3u) + vec3(0u, 1u, 2u);\n    let vertex_ids = vec3(get_meshlet_vertex_id(index_ids[0]), get_meshlet_vertex_id(index_ids[1]), get_meshlet_vertex_id(index_ids[2]));\n    let vertex_0 = viewport_vertices[vertex_ids[2]];\n    let vertex_1 = viewport_vertices[vertex_ids[1]];\n    let vertex_2 = viewport_vertices[vertex_ids[0]];\n    let packed_ids = (cluster_id << 7u) | triangle_id;\n\n    // Backface culling\n    let triangle_double_area = edge_function(vertex_0.xy, vertex_1.xy, vertex_2.xy);\n    if triangle_double_area <= 0.0 { return; }\n\n    // Setup triangle gradients\n    let w_x = vec3(vertex_1.y - vertex_2.y, vertex_2.y - vertex_0.y, vertex_0.y - vertex_1.y);\n    let w_y = vec3(vertex_2.x - vertex_1.x, vertex_0.x - vertex_2.x, vertex_1.x - vertex_0.x);\n    let vertices_z = vec3(vertex_0.z, vertex_1.z, vertex_2.z) / triangle_double_area;\n    let z_x = dot(vertices_z, w_x);\n    let z_y = dot(vertices_z, w_y);\n\n    // Compute triangle bounding box\n    var min_x = floor(min3(vertex_0.x, vertex_1.x, vertex_2.x));\n    var min_y = floor(min3(vertex_0.y, vertex_1.y, vertex_2.y));\n    var max_x = ceil(max3(vertex_0.x, vertex_1.x, vertex_2.x));\n    var max_y = ceil(max3(vertex_0.y, vertex_1.y, vertex_2.y));\n    min_x = max(min_x, 0.0);\n    min_y = max(min_y, 0.0);\n    max_x = min(max_x, view.viewport.z - 1.0);\n    max_y = min(max_y, view.viewport.w - 1.0);\n\n    // Setup initial triangle equations\n    let starting_pixel = vec2(min_x, min_y) + 0.5;\n    var w_row = vec3(\n        edge_function(vertex_1.xy, vertex_2.xy, starting_pixel),\n        edge_function(vertex_2.xy, vertex_0.xy, starting_pixel),\n        edge_function(vertex_0.xy, vertex_1.xy, starting_pixel),\n    );\n    var z_row = dot(vertices_z, w_row);\n\n    // Rasterize triangle\n    if subgroupAny(max_x - min_x > 4.0) {\n        // Scanline setup\n        let edge_012 = -w_x;\n        let open_edge = edge_012 < vec3(0.0);\n        let inverse_edge_012 = select(1.0 / edge_012, vec3(1e8), edge_012 == vec3(0.0));\n        let max_x_diff = vec3(max_x - min_x);\n        for (var y = min_y; y <= max_y; y += 1.0) {\n            // Calculate start and end X interval for pixels in this row within the triangle\n            let cross_x = w_row * inverse_edge_012;\n            let min_x2 = select(vec3(0.0), cross_x, open_edge);\n            let max_x2 = select(cross_x, max_x_diff, open_edge);\n            var x0 = ceil(max3(min_x2[0], min_x2[1], min_x2[2]));\n            var x1 = min3(max_x2[0], max_x2[1], max_x2[2]);\n\n            var w = w_row + w_x * x0;\n            var z = z_row + z_x * x0;\n            x0 += min_x;\n            x1 += min_x;\n\n            // Iterate scanline X interval\n            for (var x = x0; x <= x1; x += 1.0) {\n                // Check if point at pixel is within triangle (TODO: this shouldn\'t be needed, but there\'s bugs without it)\n                if min3(w[0], w[1], w[2]) >= 0.0 {\n                    write_visibility_buffer_pixel(x, y, z, packed_ids);\n                }\n\n                // Increment triangle equations along the X-axis\n                w += w_x;\n                z += z_x;\n            }\n\n            // Increment triangle equations along the Y-axis\n            w_row += w_y;\n            z_row += z_y;\n        }\n    } else {\n        // Iterate over every pixel in the triangle\'s bounding box\n        for (var y = min_y; y <= max_y; y += 1.0) {\n            var w = w_row;\n            var z = z_row;\n\n            for (var x = min_x; x <= max_x; x += 1.0) {\n                // Check if point at pixel is within triangle\n                if min3(w[0], w[1], w[2]) >= 0.0 {\n                    write_visibility_buffer_pixel(x, y, z, packed_ids);\n                }\n\n                // Increment triangle equations along the X-axis\n                w += w_x;\n                z += z_x;\n            }\n\n            // Increment triangle equations along the Y-axis\n            w_row += w_y;\n            z_row += z_y;\n        }\n    }\n}\n\nfn write_visibility_buffer_pixel(x: f32, y: f32, z: f32, packed_ids: u32) {\n    let depth = bitcast<u32>(z);\n    @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n    let visibility = (u64(depth) << 32u) | u64(packed_ids);\n    @else\n    let visibility = depth;\n    textureAtomicMax(meshlet_visibility_buffer, vec2(u32(x), u32(y)), visibility);\n}\n\nfn edge_function(a: vec2<f32>, b: vec2<f32>, c: vec2<f32>) -> f32 {\n    return (b.x - a.x) * (c.y - a.y) - (b.y - a.y) * (c.x - a.x);\n}\n\nfn min3(a: f32, b: f32, c: f32) -> f32 {\n    return min(a, min(b, c));\n}\n\nfn max3(a: f32, b: f32, c: f32) -> f32 {\n    return max(a, max(b, c));\n}\n");
    }
};embedded_asset!(app, "visibility_buffer_software_raster.wesl");
148        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "visibility_buffer_hardware_raster.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "visibility_buffer_hardware_raster.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! Vertex/fragment shader for rasterizing large clusters into a visibility buffer.\n\nimport package::{\n    meshlet::bindings::{\n        meshlet_cluster_meshlet_ids,\n        meshlets,\n        meshlet_cluster_instance_ids,\n        meshlet_instance_uniforms,\n        meshlet_raster_clusters,\n        meshlet_previous_raster_counts,\n        meshlet_visibility_buffer,\n        view,\n        get_meshlet_triangle_count,\n        get_meshlet_vertex_id,\n        get_meshlet_vertex_position,\n    },\n    render::mesh_functions::mesh_position_local_to_world,\n};\nimport bevy_render::maths::affine3_to_square;\nvar<immediate> meshlet_raster_cluster_rightmost_slot: u32;\n\nstruct VertexOutput {\n    @builtin(position) position: vec4<f32>,\n    @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n    @location(0) @interpolate(flat) packed_ids: u32,\n}\n\n@vertex\nfn vertex(@builtin(instance_index) instance_index: u32, @builtin(vertex_index) vertex_index: u32) -> VertexOutput {\n    let cluster_in_draw = meshlet_previous_raster_counts[1] + instance_index;\n    let cluster_id = meshlet_raster_cluster_rightmost_slot - cluster_in_draw;\n    let instanced_offset = meshlet_raster_clusters[cluster_id];\n    var meshlet = meshlets[instanced_offset.offset];\n\n    let triangle_id = vertex_index / 3u;\n    if triangle_id >= get_meshlet_triangle_count(&meshlet) { return dummy_vertex(); }\n    let index_id = vertex_index;\n    let vertex_id = get_meshlet_vertex_id(meshlet.start_index_id + index_id);\n\n    let instance_uniform = meshlet_instance_uniforms[instanced_offset.instance_id];\n\n    let vertex_position = get_meshlet_vertex_position(&meshlet, vertex_id);\n    let world_from_local = affine3_to_square(instance_uniform.world_from_local);\n    let world_position = mesh_position_local_to_world(world_from_local, vec4(vertex_position, 1.0));\n    let clip_position = view.clip_from_world * vec4(world_position.xyz, 1.0);\n\n    @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n    return VertexOutput(\n        clip_position,\n        (cluster_id << 7u) | triangle_id,\n    );\n    @else\n    return VertexOutput(\n        clip_position,\n    );\n}\n\n@fragment\nfn fragment(vertex_output: VertexOutput) {\n    let depth = bitcast<u32>(vertex_output.position.z);\n    @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n    let visibility = (u64(depth) << 32u) | u64(vertex_output.packed_ids);\n    @else\n    let visibility = depth;\n    textureAtomicMax(meshlet_visibility_buffer, vec2<u32>(vertex_output.position.xy), visibility);\n}\n\nfn dummy_vertex() -> VertexOutput {\n    @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n    return VertexOutput(\n        vec4(divide(0.0, 0.0)), // NaN vertex position\n        0u,\n    );\n    @else\n    return VertexOutput(\n        vec4(divide(0.0, 0.0)),\n    );\n}\n\n/// Naga doesn\'t allow divide by zero literals, but this lets us work around it\nfn divide(a: f32, b: f32) -> f32 {\n    return a / b;\n}\n");
    }
};embedded_asset!(app, "visibility_buffer_hardware_raster.wesl");
149        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "meshlet_mesh_material.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "meshlet_mesh_material.wesl");
        embedded.insert_asset(watched_path, &path,
            b"import package::{\n    meshlet::visibility_buffer_resolve::resolve_vertex_output,\n    prepass::io as prepass_io,\n    render::{\n        view_transformations::uv_to_ndc,\n        pbr_prepass_functions,\n        utils::rand_f,\n    },\n};\n\n@vertex\nfn vertex(@builtin(vertex_index) vertex_input: u32) -> @builtin(position) vec4<f32> {\n    let vertex_index = vertex_input % 3u;\n    let material_id = vertex_input / 3u;\n    let material_depth = f32(material_id) / 65535.0;\n    let uv = vec2<f32>(vec2(vertex_index >> 1u, vertex_index & 1u)) * 2.0;\n    return vec4(uv_to_ndc(uv), material_depth, 1.0);\n}\n\n@fragment\nfn fragment(@builtin(position) frag_coord: vec4<f32>) -> @location(0) vec4<f32> {\n    let vertex_output = resolve_vertex_output(frag_coord);\n    var rng = vertex_output.cluster_id;\n    let color = vec3(rand_f(&rng), rand_f(&rng), rand_f(&rng));\n    return vec4(color, 1.0);\n}\n\n@if(PREPASS_FRAGMENT)\n@fragment\nfn prepass_fragment(@builtin(position) frag_coord: vec4<f32>) -> prepass_io::FragmentOutput {\n    let vertex_output = resolve_vertex_output(frag_coord);\n\n    var out: prepass_io::FragmentOutput;\n\n    @if(NORMAL_PREPASS)\n    out.normal = vec4(vertex_output.world_normal * 0.5 + vec3(0.5), 1.0);\n\n    @if(MOTION_VECTOR_PREPASS)\n    out.motion_vector = vertex_output.motion_vector;\n\n    @if(DEFERRED_PREPASS)\n    // There isn\'t any material info available for this default prepass shader so we are just writing\xc2\xa0\n    // emissive magenta out to the deferred gbuffer to be rendered by the first deferred lighting pass layer.\n    // This is here so if the default prepass fragment is used for deferred magenta will be rendered, and also\n    // as an example to show that a user could write to the deferred gbuffer if they were to start from this shader.\n    out.deferred = vec4(0u, package::render::rgb9e5::vec3_to_rgb9e5_(vec3(1.0, 0.0, 1.0)), 0u, 0u);\n    @if(DEFERRED_PREPASS)\n    out.deferred_lighting_pass_id = 1u;\n\n    return out;\n}\n");
    }
};embedded_asset!(app, "meshlet_mesh_material.wesl");
150        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "resolve_render_targets.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "resolve_render_targets.wesl");
        embedded.insert_asset(watched_path, &path,
            b"import bevy_core_pipeline::fullscreen_vertex_shader::FullscreenVertexOutput;\nimport package::meshlet::bindings::InstancedOffset;\n\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n@group(0) @binding(0) var meshlet_visibility_buffer: texture_storage_2d<r64uint, read>;\n@else\n@group(0) @binding(0) var meshlet_visibility_buffer: texture_storage_2d<r32uint, read>;\n@group(0) @binding(1) var<storage, read> meshlet_raster_clusters: array<InstancedOffset>;  // Per cluster\n@group(0) @binding(2) var<storage, read> meshlet_instance_material_ids: array<u32>; // Per entity instance\n\n/// This pass writes out the depth texture.\n@fragment\nfn resolve_depth(in: FullscreenVertexOutput) -> @builtin(frag_depth) f32 {\n    let visibility = textureLoad(meshlet_visibility_buffer, vec2<u32>(in.position.xy)).r;\n    @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n    let depth = u32(visibility >> 32u);\n    @else\n    let depth = visibility;\n\n    if depth == 0u { discard; }\n\n    return bitcast<f32>(depth);\n}\n\n/// This pass writes out the material depth texture.\n@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)\n@fragment\nfn resolve_material_depth(in: FullscreenVertexOutput) -> @builtin(frag_depth) f32 {\n    let visibility = textureLoad(meshlet_visibility_buffer, vec2<u32>(in.position.xy)).r;\n\n    let depth = visibility >> 32u;\n    if depth == 0lu { discard; }\n\n    let cluster_id = u32(visibility) >> 7u;\n    let instance_id = meshlet_raster_clusters[cluster_id].instance_id;\n    let material_id = meshlet_instance_material_ids[instance_id];\n    return f32(material_id) / 65535.0;\n}\n");
    }
};embedded_asset!(app, "resolve_render_targets.wesl");
151        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "remap_1d_to_2d_dispatch.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "remap_1d_to_2d_dispatch.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! Remaps an indirect 1d to 2d dispatch for devices with low dispatch size limit.\n\nstruct DispatchIndirectArgs {\n    x: u32,\n    y: u32,\n    z: u32,\n}\n\n@group(0) @binding(0) var<storage, read_write> meshlet_software_raster_indirect_args: DispatchIndirectArgs;\n@group(0) @binding(1) var<storage, read_write> meshlet_software_raster_cluster_count: u32;\nvar<immediate> max_compute_workgroups_per_dimension: u32;\n\n@compute\n@workgroup_size(1, 1, 1)\nfn remap_dispatch() {\n    let cluster_count = meshlet_software_raster_indirect_args.x;\n\n    if cluster_count > max_compute_workgroups_per_dimension {\n        let n = u32(ceil(sqrt(f32(cluster_count))));\n        meshlet_software_raster_indirect_args.x = n;\n        meshlet_software_raster_indirect_args.y = n;\n        meshlet_software_raster_cluster_count = cluster_count;\n    }\n}\n");
    }
};embedded_asset!(app, "remap_1d_to_2d_dispatch.wesl");
152        {
    {
        let mut embedded =
            app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
        let path =
            {
                let crate_name =
                    "bevy_pbr::meshlet".split(':').next().unwrap();
                ::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
                    "src".as_ref(), "src/meshlet/mod.rs".as_ref(),
                    "fill_counts.wesl".as_ref())
            };
        let watched_path =
            ::bevy_asset::io::embedded::watched_path("src/meshlet/mod.rs",
                "fill_counts.wesl");
        embedded.insert_asset(watched_path, &path,
            b"//! Copies the counts of meshlets in the hardware and software buckets, resetting the counters in the process.\n\nstruct DispatchIndirectArgs {\n    x: u32,\n    y: u32,\n    z: u32,\n}\n\nstruct DrawIndirectArgs {\n    vertex_count: u32,\n    instance_count: u32,\n    first_vertex: u32,\n    first_instance: u32,\n}\n\n@group(0) @binding(0) var<storage, read_write> meshlet_software_raster_indirect_args: DispatchIndirectArgs;\n@group(0) @binding(1) var<storage, read_write> meshlet_hardware_raster_indirect_args: DrawIndirectArgs;\n@group(0) @binding(2) var<storage, read_write> meshlet_previous_raster_counts: array<u32>;\n@if(MESHLET_2D_DISPATCH)\n@group(0) @binding(3) var<storage, read_write> meshlet_software_raster_cluster_count: u32;\n\n@compute\n@workgroup_size(1, 1, 1)\nfn fill_counts() {\n    @if(MESHLET_2D_DISPATCH)\n    meshlet_previous_raster_counts[0] += meshlet_software_raster_cluster_count;\n    @else\n    meshlet_previous_raster_counts[0] += meshlet_software_raster_indirect_args.x;\n    meshlet_software_raster_indirect_args.x = 0;\n\n    meshlet_previous_raster_counts[1] += meshlet_hardware_raster_indirect_args.instance_count;\n    meshlet_hardware_raster_indirect_args.instance_count = 0;\n}\n");
    }
};embedded_asset!(app, "fill_counts.wesl");
153
154        app.init_asset::<MeshletMesh>()
155            .register_asset_loader(MeshletMeshLoader);
156
157        let Some(render_app) = app.get_sub_app_mut(RenderApp) else {
158            return;
159        };
160
161        // Create a variable here so we can move-capture it.
162        let cluster_buffer_slots = self.cluster_buffer_slots;
163        let init_resource_manager_system =
164            move |mut commands: Commands, render_device: Res<RenderDevice>| {
165                commands
166                    .insert_resource(ResourceManager::new(cluster_buffer_slots, &render_device));
167            };
168
169        render_app
170            .insert_resource(InstanceManager::new())
171            .add_systems(
172                RenderStartup,
173                (
174                    check_meshlet_features,
175                    (
176                        (init_resource_manager_system, init_meshlet_pipelines).chain(),
177                        init_meshlet_mesh_manager,
178                    ),
179                )
180                    .chain(),
181            )
182            .add_systems(ExtractSchedule, extract_meshlet_mesh_entities)
183            .add_systems(
184                Render,
185                (
186                    perform_pending_meshlet_mesh_writes.in_set(RenderSystems::PrepareAssets),
187                    configure_meshlet_views
188                        .after(prepare_view_targets)
189                        .in_set(RenderSystems::PrepareViews),
190                    prepare_meshlet_per_frame_resources.in_set(RenderSystems::PrepareResources),
191                    prepare_meshlet_view_bind_groups.in_set(RenderSystems::PrepareBindGroups),
192                    queue_material_meshlet_meshes.in_set(RenderSystems::QueueMeshes),
193                    prepare_material_meshlet_meshes_main_opaque_pass
194                        .in_set(RenderSystems::QueueMeshes)
195                        .before(queue_material_meshlet_meshes),
196                ),
197            )
198            .add_systems(
199                Core3d,
200                (
201                    meshlet_visibility_buffer_raster
202                        .before(per_view_shadow_pass::<EARLY_SHADOW_PASS>),
203                    meshlet_prepass
204                        .after(per_view_shadow_pass::<EARLY_SHADOW_PASS>)
205                        .in_set(Core3dSystems::Prepass),
206                    meshlet_deferred_gbuffer_prepass
207                        .after(meshlet_prepass)
208                        .in_set(Core3dSystems::Prepass),
209                    meshlet_main_opaque_pass
210                        .before(main_opaque_pass_3d)
211                        .in_set(Core3dSystems::MainPass),
212                ),
213            );
214    }
215}
216
217fn check_meshlet_features(render_device: Res<RenderDevice>) {
218    let features = render_device.features();
219    if !features.contains(MeshletPlugin::required_wgpu_features()) {
220        {
    use ::tracing::__macro_support::Callsite as _;
    static __CALLSITE: ::tracing::callsite::DefaultCallsite =
        {
            static META: ::tracing::Metadata<'static> =
                {
                    ::tracing_core::metadata::Metadata::new("event src/meshlet/mod.rs:220",
                        "bevy_pbr::meshlet", ::tracing::Level::ERROR,
                        ::tracing_core::__macro_support::Option::Some("src/meshlet/mod.rs"),
                        ::tracing_core::__macro_support::Option::Some(220u32),
                        ::tracing_core::__macro_support::Option::Some("bevy_pbr::meshlet"),
                        ::tracing_core::field::FieldSet::new(&["message"],
                            ::tracing_core::callsite::Identifier(&__CALLSITE)),
                        ::tracing::metadata::Kind::EVENT)
                };
            ::tracing::callsite::DefaultCallsite::new(&META)
        };
    let enabled =
        ::tracing::Level::ERROR <= ::tracing::level_filters::STATIC_MAX_LEVEL
                &&
                ::tracing::Level::ERROR <=
                    ::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!("MeshletPlugin can\'t be used. GPU lacks support for required features: {0:?}.",
                                                    MeshletPlugin::required_wgpu_features().difference(features))
                                            as &dyn ::tracing::field::Value))])
            });
    } else { ; }
};error!(
221            "MeshletPlugin can't be used. GPU lacks support for required features: {:?}.",
222            MeshletPlugin::required_wgpu_features().difference(features)
223        );
224        std::process::exit(1);
225    }
226}
227
228/// The meshlet mesh equivalent of [`bevy_mesh::Mesh3d`].
229#[derive(
230    #[doc =
"**Required Components**: [`Transform`], [`PreviousGlobalTransform`], [`Visibility`], [`VisibilityClass`]. \n\n A component's Required Components are inserted whenever it is inserted. Note that this will also insert the required components _of_ the required components, recursively, in depth-first order."]
impl bevy_ecs::component::Component for MeshletMesh3d 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) {
        required_components.register_required::<Transform>(<Transform as
                ::core::default::Default>::default);
        required_components.register_required::<PreviousGlobalTransform>(<PreviousGlobalTransform
                as ::core::default::Default>::default);
        required_components.register_required::<Visibility>(<Visibility as
                ::core::default::Default>::default);
        required_components.register_required::<VisibilityClass>(<VisibilityClass
                as ::core::default::Default>::default);
    }
    fn on_add()
        -> ::core::option::Option<bevy_ecs::lifecycle::ComponentHook> {
        ::core::option::Option::Some(visibility::add_visibility_class::<MeshletMesh3d>)
    }
    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, impl bevy_ecs::template::FromTemplate for MeshletMesh3d {
    type Template = MeshletMesh3dTemplate;
}
impl ::core::marker::Unpin for MeshletMesh3d where
    for<'a> [()]: bevy_ecs::template::SpecializeFromTemplate {}
#[allow(missing_docs)]
pub struct MeshletMesh3dTemplate(pub <Handle<MeshletMesh> as
    bevy_ecs::template::FromTemplate>::Template);
impl bevy_ecs::template::Template for MeshletMesh3dTemplate {
    type Output = MeshletMesh3d;
    fn build_template(&self,
        context: &mut bevy_ecs::template::TemplateContext)
        -> bevy_ecs::error::Result<Self::Output> {
        bevy_ecs::error::Result::Ok(MeshletMesh3d(self.0.build_template(context)?))
    }
    fn clone_template(&self) -> Self {
        Self(bevy_ecs::template::Template::clone_template(&self.0))
    }
}
impl ::core::default::Default for MeshletMesh3dTemplate {
    fn default() -> Self { Self(::core::default::Default::default()) }
}FromTemplate, #[automatically_derived]
impl ::core::clone::Clone for MeshletMesh3d {
    #[inline]
    fn clone(&self) -> Self { Self(::core::clone::Clone::clone(&self.0)) }
}Clone, #[automatically_derived]
impl ::core::fmt::Debug for MeshletMesh3d {
    #[inline]
    fn fmt(&self, f: &mut ::core::fmt::Formatter) -> ::core::fmt::Result {
        ::core::fmt::Formatter::debug_tuple_field1_finish(f, "MeshletMesh3d",
            &&self.0)
    }
}Debug, #[automatically_derived]
impl ::core::default::Default for MeshletMesh3d {
    #[inline]
    fn default() -> Self { Self(::core::default::Default::default()) }
}Default, impl ::core::ops::Deref for MeshletMesh3d {
    type Target = Handle<MeshletMesh>;
    fn deref(&self) -> &Self::Target { &self.0 }
}Deref, impl ::core::ops::DerefMut for MeshletMesh3d {
    fn deref_mut(&mut self) -> &mut Self::Target { &mut self.0 }
}DerefMut, const _: () =
    {
        impl bevy_reflect::GetTypeRegistration for MeshletMesh3d where  {
            fn get_type_registration() -> bevy_reflect::TypeRegistration {
                let mut registration =
                    bevy_reflect::TypeRegistration::of::<Self>();
                registration.insert(<bevy_reflect::ReflectFromPtr as
                            bevy_reflect::CreateTypeData<Self>>::create_type_data(()));
                registration.insert(<bevy_reflect::ReflectFromReflect as
                            bevy_reflect::CreateTypeData<Self>>::create_type_data(()));
                registration.register_type_data_with::<ReflectComponent, Self,
                    _>(());
                registration.register_type_data_with::<ReflectDefault, Self,
                    _>(());
                registration
            }
            #[inline(never)]
            fn register_type_dependencies(registry:
                    &mut bevy_reflect::TypeRegistry) {
                <Handle<MeshletMesh> as
                        bevy_reflect::__macro_exports::RegisterForReflection>::__register(registry);
            }
        }
        impl bevy_reflect::Typed for MeshletMesh3d where  {
            #[inline(never)]
            fn type_info() -> &'static bevy_reflect::TypeInfo {
                static CELL: bevy_reflect::utility::NonGenericTypeInfoCell =
                    bevy_reflect::utility::NonGenericTypeInfoCell::new();
                CELL.get_or_set(#[inline(never)] ||
                        {
                            bevy_reflect::TypeInfo::TupleStruct(bevy_reflect::tuple_struct::TupleStructInfo::new::<Self>(&[bevy_reflect::UnnamedField::new::<Handle<MeshletMesh>>(0usize)]))
                        })
            }
        }
        #[allow(deprecated, reason =
        "derives on a deprecated type shouldn't be considered a usage")]
        impl bevy_reflect::TypePath for MeshletMesh3d where  {
            fn type_path() -> &'static str {
                "bevy_pbr::meshlet::MeshletMesh3d"
            }
            fn short_type_path() -> &'static str { "MeshletMesh3d" }
            fn type_ident() -> ::core::option::Option<&'static str> {
                ::core::option::Option::Some("MeshletMesh3d")
            }
            fn crate_name() -> ::core::option::Option<&'static str> {
                ::core::option::Option::Some("bevy_pbr::meshlet".split(':').next().unwrap())
            }
            fn module_path() -> ::core::option::Option<&'static str> {
                ::core::option::Option::Some("bevy_pbr::meshlet")
            }
        }
        impl bevy_reflect::Reflect for MeshletMesh3d where  {
            #[inline]
            fn into_any(self:
                    bevy_reflect::__macro_exports::alloc_utils::Box<Self>)
                ->
                    bevy_reflect::__macro_exports::alloc_utils::Box<dyn ::core::any::Any> {
                self
            }
            #[inline]
            fn as_any(&self) -> &dyn ::core::any::Any { self }
            #[inline]
            fn as_any_mut(&mut self) -> &mut dyn ::core::any::Any { self }
            #[inline]
            fn into_reflect(self:
                    bevy_reflect::__macro_exports::alloc_utils::Box<Self>)
                ->
                    bevy_reflect::__macro_exports::alloc_utils::Box<dyn bevy_reflect::Reflect> {
                self
            }
            #[inline]
            fn as_reflect(&self) -> &dyn bevy_reflect::Reflect { self }
            #[inline]
            fn as_reflect_mut(&mut self) -> &mut dyn bevy_reflect::Reflect {
                self
            }
            #[inline]
            fn set(&mut self,
                value:
                    bevy_reflect::__macro_exports::alloc_utils::Box<dyn bevy_reflect::Reflect>)
                ->
                    ::core::result::Result<(),
                    bevy_reflect::__macro_exports::alloc_utils::Box<dyn bevy_reflect::Reflect>> {
                *self = <dyn bevy_reflect::Reflect>::take(value)?;
                ::core::result::Result::Ok(())
            }
        }
        #[allow(non_upper_case_globals)]
        const _: () =
            {
                static __INVENTORY: ::inventory::Node =
                    ::inventory::Node {
                        value: &{
                                bevy_reflect::__macro_exports::auto_register::AutomaticReflectRegistrations(<MeshletMesh3d
                                        as
                                        bevy_reflect::__macro_exports::auto_register::RegisterForReflection>::__register)
                            },
                        next: ::inventory::__private::UnsafeCell::new(::inventory::__private::Option::None),
                    };
                #[link_section = ".text.startup"]
                unsafe extern "C" fn __ctor() {
                    unsafe {
                        ::inventory::ErasedNode::submit(__INVENTORY.value,
                            &__INVENTORY)
                    }
                }
                #[used]
                #[link_section = ".init_array"]
                static __CTOR: unsafe extern "C" fn() = __ctor;
            };
        impl bevy_reflect::tuple_struct::TupleStruct for MeshletMesh3d where
            {
            fn field(&self, index: usize)
                -> ::core::option::Option<&dyn bevy_reflect::PartialReflect> {
                match index {
                    0usize => ::core::option::Option::Some(&self.0),
                    _ => ::core::option::Option::None,
                }
            }
            fn field_mut(&mut self, index: usize)
                ->
                    ::core::option::Option<&mut dyn bevy_reflect::PartialReflect> {
                match index {
                    0usize => ::core::option::Option::Some(&mut self.0),
                    _ => ::core::option::Option::None,
                }
            }
            #[inline]
            fn field_len(&self) -> usize { 1usize }
            #[inline]
            fn iter_fields(&self)
                -> bevy_reflect::tuple_struct::TupleStructFieldIter {
                bevy_reflect::tuple_struct::TupleStructFieldIter::new(self)
            }
            fn to_dynamic_tuple_struct(&self)
                ->
                    ::core::result::Result<bevy_reflect::tuple_struct::DynamicTupleStruct,
                    bevy_reflect::ReflectCloneError> {
                let mut dynamic:
                        bevy_reflect::tuple_struct::DynamicTupleStruct =
                    ::core::default::Default::default();
                dynamic.set_represented_type(bevy_reflect::PartialReflect::get_represented_type_info(self));
                dynamic.insert_boxed(bevy_reflect::PartialReflect::to_dynamic(&self.0)?);
                ::core::result::Result::Ok(dynamic)
            }
        }
        impl bevy_reflect::PartialReflect for MeshletMesh3d where  {
            #[inline]
            fn get_represented_type_info(&self)
                -> ::core::option::Option<&'static bevy_reflect::TypeInfo> {
                ::core::option::Option::Some(<Self as
                            bevy_reflect::Typed>::type_info())
            }
            #[inline]
            fn try_apply(&mut self, value: &dyn bevy_reflect::PartialReflect)
                -> ::core::result::Result<(), bevy_reflect::ApplyError> {
                if let bevy_reflect::ReflectRef::TupleStruct(struct_value) =
                        bevy_reflect::PartialReflect::reflect_ref(value) {
                    for (i, value) in
                        ::core::iter::Iterator::enumerate(bevy_reflect::tuple_struct::TupleStruct::iter_fields(struct_value))
                        {
                        if let ::core::option::Option::Some(v) =
                                bevy_reflect::tuple_struct::TupleStruct::field_mut(self, i)
                            {
                            bevy_reflect::PartialReflect::try_apply(v, value)?;
                        }
                    }
                } else {
                    return ::core::result::Result::Err(bevy_reflect::ApplyError::MismatchedKinds {
                                from_kind: bevy_reflect::PartialReflect::reflect_kind(value),
                                to_kind: bevy_reflect::ReflectKind::TupleStruct,
                            });
                }
                ::core::result::Result::Ok(())
            }
            #[inline]
            fn reflect_kind(&self) -> bevy_reflect::ReflectKind {
                bevy_reflect::ReflectKind::TupleStruct
            }
            #[inline]
            fn reflect_ref(&self) -> bevy_reflect::ReflectRef {
                bevy_reflect::ReflectRef::TupleStruct(self)
            }
            #[inline]
            fn reflect_mut(&mut self) -> bevy_reflect::ReflectMut {
                bevy_reflect::ReflectMut::TupleStruct(self)
            }
            #[inline]
            fn reflect_owned(self:
                    bevy_reflect::__macro_exports::alloc_utils::Box<Self>)
                -> bevy_reflect::ReflectOwned {
                bevy_reflect::ReflectOwned::TupleStruct(self)
            }
            #[inline]
            fn try_into_reflect(self:
                    bevy_reflect::__macro_exports::alloc_utils::Box<Self>)
                ->
                    ::core::result::Result<bevy_reflect::__macro_exports::alloc_utils::Box<dyn bevy_reflect::Reflect>,
                    bevy_reflect::__macro_exports::alloc_utils::Box<dyn bevy_reflect::PartialReflect>> {
                ::core::result::Result::Ok(self)
            }
            #[inline]
            fn try_as_reflect(&self)
                -> ::core::option::Option<&dyn bevy_reflect::Reflect> {
                ::core::option::Option::Some(self)
            }
            #[inline]
            fn try_as_reflect_mut(&mut self)
                -> ::core::option::Option<&mut dyn bevy_reflect::Reflect> {
                ::core::option::Option::Some(self)
            }
            #[inline]
            fn into_partial_reflect(self:
                    bevy_reflect::__macro_exports::alloc_utils::Box<Self>)
                ->
                    bevy_reflect::__macro_exports::alloc_utils::Box<dyn bevy_reflect::PartialReflect> {
                self
            }
            #[inline]
            fn as_partial_reflect(&self)
                -> &dyn bevy_reflect::PartialReflect {
                self
            }
            #[inline]
            fn as_partial_reflect_mut(&mut self)
                -> &mut dyn bevy_reflect::PartialReflect {
                self
            }
            fn reflect_partial_eq(&self,
                value: &dyn bevy_reflect::PartialReflect)
                -> ::core::option::Option<bool> {
                let value =
                    <dyn bevy_reflect::PartialReflect>::try_downcast_ref::<Self>(value);
                if let ::core::option::Option::Some(value) = value {
                    ::core::option::Option::Some(::core::cmp::PartialEq::eq(self,
                            value))
                } else { ::core::option::Option::Some(false) }
            }
            fn reflect_partial_cmp(&self,
                value: &dyn bevy_reflect::PartialReflect)
                -> ::core::option::Option<::core::cmp::Ordering> {
                (bevy_reflect::tuple_struct::tuple_struct_partial_cmp)(self,
                    value)
            }
            #[inline]
            fn reflect_clone(&self)
                ->
                    ::core::result::Result<bevy_reflect::__macro_exports::alloc_utils::Box<dyn bevy_reflect::Reflect>,
                    bevy_reflect::ReflectCloneError> {
                ::core::result::Result::Ok(bevy_reflect::__macro_exports::alloc_utils::Box::new(::core::clone::Clone::clone(self)))
            }
        }
        impl bevy_reflect::FromReflect for MeshletMesh3d where  {
            fn from_reflect(reflect: &dyn bevy_reflect::PartialReflect)
                -> ::core::option::Option<Self> {
                if let bevy_reflect::ReflectRef::TupleStruct(__ref_struct) =
                        bevy_reflect::PartialReflect::reflect_ref(reflect) {
                    let mut __this =
                        <Self as ::core::default::Default>::default();
                    if let ::core::option::Option::Some(__field) =
                            (||
                                        <Handle<MeshletMesh> as
                                                bevy_reflect::FromReflect>::from_reflect(bevy_reflect::tuple_struct::TupleStruct::field(__ref_struct,
                                                    0)?))() {
                        __this.0 = __field;
                    }
                    ::core::option::Option::Some(__this)
                } else { ::core::option::Option::None }
            }
        }
    };Reflect, #[automatically_derived]
impl ::core::marker::StructuralPartialEq for MeshletMesh3d { }
#[automatically_derived]
impl ::core::cmp::PartialEq for MeshletMesh3d {
    #[inline]
    fn eq(&self, other: &Self) -> bool { self.0 == other.0 }
}PartialEq, #[automatically_derived]
impl ::core::cmp::Eq for MeshletMesh3d {
    #[inline]
    #[doc(hidden)]
    #[coverage(off)]
    fn assert_fields_are_eq(&self) {
        let _: ::core::cmp::AssertParamIsEq<Handle<MeshletMesh>>;
    }
}Eq, #[allow(deprecated)]
#[allow(unreachable_code)]
#[automatically_derived]
impl derive_more::core::convert::From<(Handle<MeshletMesh>)> for MeshletMesh3d
    {
    #[inline]
    fn from(value: (Handle<MeshletMesh>)) -> Self { MeshletMesh3d(value) }
}From,
231)]
232#[reflect(Component, Default, Clone, PartialEq)]
233#[require(Transform, PreviousGlobalTransform, Visibility, VisibilityClass)]
234#[component(on_add = visibility::add_visibility_class::<MeshletMesh3d>)]
235pub struct MeshletMesh3d(pub Handle<MeshletMesh>);
236
237impl From<MeshletMesh3d> for AssetId<MeshletMesh> {
238    fn from(mesh: MeshletMesh3d) -> Self {
239        mesh.id()
240    }
241}
242
243impl From<&MeshletMesh3d> for AssetId<MeshletMesh> {
244    fn from(mesh: &MeshletMesh3d) -> Self {
245        mesh.id()
246    }
247}
248
249fn configure_meshlet_views(
250    mut views_3d: Query<(
251        Entity,
252        &Msaa,
253        Has<NormalPrepass>,
254        Has<MotionVectorPrepass>,
255        Has<DeferredPrepass>,
256    )>,
257    mut commands: Commands,
258) {
259    for (entity, msaa, normal_prepass, motion_vector_prepass, deferred_prepass) in &mut views_3d {
260        if *msaa != Msaa::Off {
261            {
    use ::tracing::__macro_support::Callsite as _;
    static __CALLSITE: ::tracing::callsite::DefaultCallsite =
        {
            static META: ::tracing::Metadata<'static> =
                {
                    ::tracing_core::metadata::Metadata::new("event src/meshlet/mod.rs:261",
                        "bevy_pbr::meshlet", ::tracing::Level::ERROR,
                        ::tracing_core::__macro_support::Option::Some("src/meshlet/mod.rs"),
                        ::tracing_core::__macro_support::Option::Some(261u32),
                        ::tracing_core::__macro_support::Option::Some("bevy_pbr::meshlet"),
                        ::tracing_core::field::FieldSet::new(&["message"],
                            ::tracing_core::callsite::Identifier(&__CALLSITE)),
                        ::tracing::metadata::Kind::EVENT)
                };
            ::tracing::callsite::DefaultCallsite::new(&META)
        };
    let enabled =
        ::tracing::Level::ERROR <= ::tracing::level_filters::STATIC_MAX_LEVEL
                &&
                ::tracing::Level::ERROR <=
                    ::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!("MeshletPlugin can\'t be used with MSAA. Add Msaa::Off to your camera to use this plugin.")
                                            as &dyn ::tracing::field::Value))])
            });
    } else { ; }
};error!("MeshletPlugin can't be used with MSAA. Add Msaa::Off to your camera to use this plugin.");
262            std::process::exit(1);
263        }
264
265        if !(normal_prepass || motion_vector_prepass || deferred_prepass) {
266            commands
267                .entity(entity)
268                .insert(MeshletViewMaterialsMainOpaquePass::default());
269        } else {
270            // TODO: Should we add both Prepass and DeferredGBufferPrepass materials here, and in other systems/nodes?
271            commands.entity(entity).insert((
272                MeshletViewMaterialsMainOpaquePass::default(),
273                MeshletViewMaterialsPrepass::default(),
274                MeshletViewMaterialsDeferredGBufferPrepass::default(),
275            ));
276        }
277    }
278}