1mod 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
78pub struct MeshletPlugin {
107 pub cluster_buffer_slots: u32,
116}
117
118impl MeshletPlugin {
119 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 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#[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 commands.entity(entity).insert((
272 MeshletViewMaterialsMainOpaquePass::default(),
273 MeshletViewMaterialsPrepass::default(),
274 MeshletViewMaterialsDeferredGBufferPrepass::default(),
275 ));
276 }
277 }
278}