1//! GPU buffers that support sparse updates if only a small number of elements
2//! have changed.
34use alloc::sync::{Arc, Weak};
5use core::{
6iter, slice,
7 sync::atomic::{AtomicU64, Ordering},
8};
910use bevy_app::{App, Plugin};
11use bevy_asset::{embedded_asset, load_embedded_asset, Handle};
12use bevy_derive::{Deref, DerefMut};
13use bevy_ecs::{
14resource::Resource,
15schedule::IntoScheduleConfigsas _,
16 system::{Res, ResMut},
17 world::{FromWorld, World},
18};
19use bevy_log::{error, info};
20use bevy_material::{
21 bind_group_layout_entries::{
22 binding_types::{storage_buffer, storage_buffer_read_only, uniform_buffer},
23BindGroupLayoutEntries,
24 },
25 descriptor::{BindGroupLayoutDescriptor, CachedComputePipelineId, ComputePipelineDescriptor},
26};
27use bevy_shader::Shader;
28use bytemuck::{Pod, Zeroable};
29use encase::ShaderType;
30use weak_table::WeakKeyHashMap;
31use wgpu::{BufferDescriptor, BufferUsages, ComputePassDescriptor, ShaderStages};
3233use crate::{
34 diagnostic::{DiagnosticsRecorder, RecordDiagnosticsas _},
35 render_resource::{
36AtomicPod, BindGroup, BindGroupEntries, Buffer, PipelineCache, RawBufferVec,
37SpecializedComputePipeline, SpecializedComputePipelines, UniformBuffer,
38 },
39 renderer::{RenderDevice, RenderGraph, RenderGraphSystems, RenderQueue},
40ExtractSchedule, RenderApp,
41};
4243/// A plugin that allows sparse updates of GPU buffers if only a small number of
44/// elements have changed.
45pub struct SparseBufferPlugin;
4647impl Pluginfor SparseBufferPlugin {
48fn build(&self, app: &mut App) {
49{
{
let mut embedded =
app.world_mut().resource_mut::<::bevy_asset::io::embedded::EmbeddedAssetRegistry>();
let path =
{
let crate_name =
"bevy_render::render_resource::sparse_buffer_vec".split(':').next().unwrap();
::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
"src".as_ref(),
"src/render_resource/sparse_buffer_vec.rs".as_ref(),
"sparse_buffer_update.wgsl".as_ref())
};
let watched_path =
::bevy_asset::io::embedded::watched_path("src/render_resource/sparse_buffer_vec.rs",
"sparse_buffer_update.wgsl");
embedded.insert_asset(watched_path, &path,
b"// A compute shader that performs scattering sparse updates to\n// `AtomicSparseBufferVec` types.\n//\n// This shader isn\'t used for every update. Only if the number of updates is\n// small is this shader used. Otherwise, the standard `write_buffer` `wgpu`\n// command is used to update the buffer in bulk.\n//\n// We issue one thread per *word*, not per element. That allows us to achieve\n// maximum parallelism, without any loops.\n\n// Metadata that describes the update.\nstruct SparseBufferUpdateMetadata {\n // The size of a single element in words.\n element_size: u32,\n // The total number of pages to be updated.\n updated_page_count: u32,\n // The base-2 logarithm of the page size.\n page_size_log2: u32,\n};\n\n// The buffer we\'re copying to.\n@group(0) @binding(0) var<storage, read_write> dest_buffer: array<u32>;\n// The buffer we\'re copying from.\n@group(0) @binding(1) var<storage> src_buffer: array<u32>;\n// For each page in `src_buffer`, the page in `dest_buffer` that we should copy\n// it to.\n@group(0) @binding(2) var<storage> indices: array<u32>;\n// Metadata that describes the operation.\n@group(0) @binding(3) var<uniform> metadata: SparseBufferUpdateMetadata;\n\n@workgroup_size(256, 1, 1)\n@compute\nfn main(@builtin(global_invocation_id) global_id: vec3<u32>) {\n // Calculate which word we are. Remember that this shader executes with one\n // thread per word.\n let invocation_index = global_id.x;\n let total_word_count = (metadata.updated_page_count << metadata.page_size_log2) *\n metadata.element_size;\n if (invocation_index >= total_word_count) {\n return;\n }\n\n // Calculate which element we are.\n let element_index = invocation_index / metadata.element_size;\n // Calculate which word *within* that element we\'re looking at.\n let word_index = invocation_index % metadata.element_size;\n // Calculate which page we\'re copying.\n let update_index = element_index >> metadata.page_size_log2;\n // Determine which element we\'re copying within that page.\n let element_index_in_page = element_index & ((1u << metadata.page_size_log2) - 1u);\n\n // Look up our destination page.\n let page_index = indices[update_index];\n // Calculate where we should write our word.\n let dest_index = ((page_index << metadata.page_size_log2) + element_index_in_page) *\n metadata.element_size + word_index;\n if (dest_index >= arrayLength(&dest_buffer)) {\n return;\n }\n\n // Copy the word over.\n let src_index = element_index * metadata.element_size + word_index;\n dest_buffer[dest_index] = src_buffer[src_index];\n}\n");
}
};embedded_asset!(app, "sparse_buffer_update.wgsl");
50 }
5152fn finish(&self, app: &mut App) {
53let Some(render_app) = app.get_sub_app_mut(RenderApp) else {
54return;
55 };
5657render_app58 .init_resource::<SparseBufferUpdateJobs>()
59 .init_resource::<SparseBufferUpdatePipelines>()
60 .init_resource::<SpecializedComputePipelines<SparseBufferUpdatePipelines>>()
61 .init_resource::<SparseBufferUpdateBindGroups>()
62 .add_systems(ExtractSchedule, clear_sparse_buffer_jobs)
63 .add_systems(
64RenderGraph,
65// We perform sparse buffer updates very early so that sparse
66 // buffers can be used in any render pass.
67update_sparse_buffers.in_set(RenderGraphSystems::Begin),
68 );
69 }
70}
7172/// A globally-unique ID that identifies this sparse buffer.
73#[derive(#[automatically_derived]
impl ::core::clone::Clone for SparseBufferId {
#[inline]
fn clone(&self) -> SparseBufferId {
let _: ::core::clone::AssertParamIsClone<u64>;
*self
}
}Clone, #[automatically_derived]
impl ::core::marker::Copy for SparseBufferId { }Copy, #[automatically_derived]
impl ::core::cmp::PartialEq for SparseBufferId {
#[inline]
fn eq(&self, other: &SparseBufferId) -> bool { self.0 == other.0 }
}PartialEq, #[automatically_derived]
impl ::core::cmp::Eq for SparseBufferId {
#[inline]
#[doc(hidden)]
#[coverage(off)]
fn assert_fields_are_eq(&self) {
let _: ::core::cmp::AssertParamIsEq<u64>;
}
}Eq, #[automatically_derived]
impl ::core::cmp::PartialOrd for SparseBufferId {
#[inline]
fn partial_cmp(&self, other: &SparseBufferId)
-> ::core::option::Option<::core::cmp::Ordering> {
::core::option::Option::Some(::core::cmp::Ord::cmp(self, other))
}
}PartialOrd, #[automatically_derived]
impl ::core::cmp::Ord for SparseBufferId {
#[inline]
fn cmp(&self, other: &SparseBufferId) -> ::core::cmp::Ordering {
::core::cmp::Ord::cmp(&self.0, &other.0)
}
}Ord, #[automatically_derived]
impl ::core::hash::Hash for SparseBufferId {
#[inline]
fn hash<__H: ::core::hash::Hasher>(&self, state: &mut __H) {
::core::hash::Hash::hash(&self.0, state)
}
}Hash, #[automatically_derived]
impl ::core::fmt::Debug for SparseBufferId {
#[inline]
fn fmt(&self, f: &mut ::core::fmt::Formatter) -> ::core::fmt::Result {
::core::fmt::Formatter::debug_tuple_field1_finish(f, "SparseBufferId",
&&self.0)
}
}Debug, impl ::core::ops::Deref for SparseBufferId {
type Target = u64;
fn deref(&self) -> &Self::Target { &self.0 }
}Deref, impl ::core::ops::DerefMut for SparseBufferId {
fn deref_mut(&mut self) -> &mut Self::Target { &mut self.0 }
}DerefMut)]
74pub struct SparseBufferId(pub u64);
7576/// An object that allows the sparse buffer ID to be query and holds the bind
77/// group for that sparse buffer alive.
78///
79/// Each sparse buffer holds a strong reference to this handle, and the
80/// [`SparseBufferUpdateBindGroups`] resource contains a weak map from this
81/// handle to the bind group. This setup ensures that, when the sparse buffer is
82/// freed, the bind groups for that sparse buffer are freed as well.
83pub type SparseBufferHandle = Arc<SparseBufferId>;
8485/// The next sparse buffer ID to be assigned.
86static NEXT_SPARSE_BUFFER_ID: AtomicU64 = AtomicU64::new(0);
8788/// The size of a single workgroup in the sparse buffer shader.
89const SPARSE_BUFFER_UPDATE_WORKGROUP_SIZE: u32 = 256;
9091/// The fraction of the buffer that may be changed before we fall back to full
92/// reupload.
93///
94/// This is set to 15% by default. This was obtained experimentally by testing
95/// very large scenes and roughly matches the values used by other engines that
96/// perform sparse buffer updates.
97const SPARSE_UPLOAD_THRESHOLD: f64 = 0.15;
9899/// The WebGPU limit on the number of workgroups that can be dispatched.
100const MAX_WORKGROUPS: u32 = 65535;
101102/// We round all allocations up to the nearest power of this.
103const REALLOCATION_FACTOR: f64 = 1.5;
104/// We round all allocations up to the nearest multiple of this.
105const REALLOCATION_SIZE_MULTIPLE: usize = 256;
106107/// The number of dirty-page bits packed into each [`AtomicU64`] word.
108const PAGES_PER_DIRTY_WORD: u32 = 64;
109110/// Pipelines for the sparse buffer update shader.
111///
112/// This shader is shared among all sparse buffer vectors.
113#[derive(impl bevy_ecs::resource::Resource for SparseBufferUpdatePipelines where
Self: ::core::marker::Send + ::core::marker::Sync + 'static {}Resource)]
114pub struct SparseBufferUpdatePipelines {
115/// The bind group layout.
116 ///
117 /// We only have one bind group layout shared among all sparse buffer
118 /// vectors.
119bind_group_layout: Option<BindGroupLayoutDescriptor>,
120/// The shader that performs the scatter operation.
121shader: Option<Handle<Shader>>,
122}
123124/// A resource, part of the render world, that stores the bind groups for each
125/// sparse buffer.
126#[derive(impl bevy_ecs::resource::Resource for SparseBufferUpdateBindGroups where
Self: ::core::marker::Send + ::core::marker::Sync + 'static {}Resource)]
127pub struct SparseBufferUpdateBindGroups {
128/// The bind groups for each sparse buffer.
129 ///
130 /// These are stored in a weak map so that when the sparse buffer goes away,
131 /// the bind group for that buffer goes away as well.
132bind_groups: WeakKeyHashMap<Weak<SparseBufferId>, SparseBufferUpdateBindGroup>,
133/// The ID of the update shader pipeline shared among all sparse buffers.
134pipeline_id: CachedComputePipelineId,
135}
136137/// A single bind group for the sparse buffer update shader.
138pub struct SparseBufferUpdateBindGroup {
139/// The actual bind group.
140bind_group: BindGroup,
141}
142143/// A resource, part of the render world, that stores all pending sparse updates
144/// to buffers.
145#[derive(impl bevy_ecs::resource::Resource for SparseBufferUpdateJobs where
Self: ::core::marker::Send + ::core::marker::Sync + 'static {}Resource, #[automatically_derived]
impl ::core::default::Default for SparseBufferUpdateJobs {
#[inline]
fn default() -> SparseBufferUpdateJobs {
SparseBufferUpdateJobs(::core::default::Default::default())
}
}Default, impl ::core::ops::Deref for SparseBufferUpdateJobs {
type Target = Vec<SparseBufferUpdateJob>;
fn deref(&self) -> &Self::Target { &self.0 }
}Deref, impl ::core::ops::DerefMut for SparseBufferUpdateJobs {
fn deref_mut(&mut self) -> &mut Self::Target { &mut self.0 }
}DerefMut)]
146pub struct SparseBufferUpdateJobs(pub Vec<SparseBufferUpdateJob>);
147148/// Describes a sparse update operation for a buffer.
149pub struct SparseBufferUpdateJob {
150/// A handle to the buffer to be updated.
151sparse_buffer_handle: SparseBufferHandle,
152/// The number of pages to update.
153updated_page_count: u32,
154/// The base-2 logarithm of the size of a page for the buffer.
155 ///
156 /// The actual page size can be computed as `1 << page_size_log2`.
157page_size_log2: u32,
158/// The size of each element in 32-bit words.
159element_word_size: u32,
160/// A debugging label for the buffer.
161label: Arc<str>,
162}
163164impl SparseBufferUpdateJob {
165/// The number of elements per page.
166fn page_size(&self) -> u32 {
1671 << self.page_size_log2
168 }
169170/// Calculates the number of words that need to be updated.
171fn words_to_update(&self) -> u32 {
172self.updated_page_count * self.page_size() * self.element_word_size
173 }
174175/// Calculates the number of workgroups that need to be dispatched.
176fn workgroup_count(&self) -> u32 {
177self.words_to_update()
178 .div_ceil(SPARSE_BUFFER_UPDATE_WORKGROUP_SIZE)
179 }
180}
181182/// A GPU type that describes a sparse update that is to be performed.
183#[derive(#[automatically_derived]
impl ::core::clone::Clone for GpuSparseBufferUpdateMetadata {
#[inline]
fn clone(&self) -> GpuSparseBufferUpdateMetadata {
let _: ::core::clone::AssertParamIsClone<u32>;
*self
}
}Clone, #[automatically_derived]
impl ::core::marker::Copy for GpuSparseBufferUpdateMetadata { }Copy, #[automatically_derived]
impl ::core::default::Default for GpuSparseBufferUpdateMetadata {
#[inline]
fn default() -> GpuSparseBufferUpdateMetadata {
GpuSparseBufferUpdateMetadata {
element_size: ::core::default::Default::default(),
updated_page_count: ::core::default::Default::default(),
page_size_log2: ::core::default::Default::default(),
}
}
}Default, impl ::encase::private::ShaderSize for GpuSparseBufferUpdateMetadata where
u32: ::encase::private::ShaderSize, u32: ::encase::private::ShaderSize,
u32: ::encase::private::ShaderSize {}ShaderType, unsafe impl ::bytemuck::Pod for GpuSparseBufferUpdateMetadata {}Pod, unsafe impl ::bytemuck::Zeroable for GpuSparseBufferUpdateMetadata {}Zeroable)]
184#[repr(C)]
185struct GpuSparseBufferUpdateMetadata {
186/// The size of a single element in 32-bit words.
187element_size: u32,
188/// The number of pages that need to be updated.
189updated_page_count: u32,
190/// The base-2 logarithm of the page size.
191 ///
192 /// That is, the page size is `1 << page_size_log2`.
193page_size_log2: u32,
194}
195196/// A system, part of the render graph, that performs sparse buffer updates to
197/// buffers for which only a small number of elements have changed.
198///
199/// This runs as early in the pipeline as possible so that sparse buffers can be
200/// used for any subsequent pass.
201fn update_sparse_buffers(
202 sparse_buffer_update_jobs: Res<SparseBufferUpdateJobs>,
203 sparse_buffer_update_bind_groups: Res<SparseBufferUpdateBindGroups>,
204 pipeline_cache: Res<PipelineCache>,
205mut diagnostics: Option<ResMut<DiagnosticsRecorder>>,
206 render_device: Res<RenderDevice>,
207 render_queue: Res<RenderQueue>,
208) {
209// Bail if we have nothing to do.
210if sparse_buffer_update_jobs.is_empty() {
211return;
212 }
213214// We need to create a command encoder since this pass isn't associated with
215 // a view.
216let mut command_encoder =
217render_device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
218 label: Some("sparse buffer update"),
219 });
220221let time_span = diagnostics222 .as_mut()
223 .map(|diagnostics| diagnostics.time_span(&mut command_encoder, "sparse buffer update"));
224225command_encoder.push_debug_group("sparse buffer update");
226227let Some(compute_pipeline) =
228pipeline_cache.get_compute_pipeline(sparse_buffer_update_bind_groups.pipeline_id)
229else {
230return;
231 };
232233// Process each sparse buffer update job.
234for sparse_buffer_update_job in sparse_buffer_update_jobs.iter() {
235let Some(sparse_buffer_update_bind_group) = sparse_buffer_update_bind_groups
236 .bind_groups
237 .get(&sparse_buffer_update_job.sparse_buffer_handle)
238else {
239continue;
240 };
241242let mut sparse_buffer_update_pass =
243 command_encoder.begin_compute_pass(&ComputePassDescriptor {
244 label: Some(&*::alloc::__export::must_use({
::alloc::fmt::format(format_args!("sparse buffer update ({0})",
&sparse_buffer_update_job.label))
})format!(
245"sparse buffer update ({})",
246&sparse_buffer_update_job.label
247 )),
248 timestamp_writes: None,
249 });
250 sparse_buffer_update_pass.set_pipeline(compute_pipeline);
251 sparse_buffer_update_pass.set_bind_group(
2520,
253&sparse_buffer_update_bind_group.bind_group,
254&[],
255 );
256 sparse_buffer_update_pass.dispatch_workgroups(
257 sparse_buffer_update_job.workgroup_count(),
2581,
2591,
260 );
261 }
262263command_encoder.pop_debug_group();
264if let Some(time_span) = time_span {
265time_span.end(&mut command_encoder);
266 }
267268render_queue.submit([command_encoder.finish()]);
269}
270271/// A system that clears out the sparse buffer update jobs in preparation for a
272/// new frame.
273fn clear_sparse_buffer_jobs(mut sparse_buffer_update_jobs: ResMut<SparseBufferUpdateJobs>) {
274sparse_buffer_update_jobs.clear();
275}
276277impl FromWorldfor SparseBufferUpdatePipelines {
278fn from_world(world: &mut World) -> Self {
279let render_device = world.resource::<RenderDevice>();
280let limit = render_device.limits().max_storage_buffers_per_shader_stage;
281282if limit < 3 {
283{
use ::tracing::__macro_support::Callsite as _;
static __CALLSITE: ::tracing::callsite::DefaultCallsite =
{
static META: ::tracing::Metadata<'static> =
{
::tracing_core::metadata::Metadata::new("event src/render_resource/sparse_buffer_vec.rs:283",
"bevy_render::render_resource::sparse_buffer_vec",
::tracing::Level::INFO,
::tracing_core::__macro_support::Option::Some("src/render_resource/sparse_buffer_vec.rs"),
::tracing_core::__macro_support::Option::Some(283u32),
::tracing_core::__macro_support::Option::Some("bevy_render::render_resource::sparse_buffer_vec"),
::tracing_core::field::FieldSet::new(&["message"],
::tracing_core::callsite::Identifier(&__CALLSITE)),
::tracing::metadata::Kind::EVENT)
};
::tracing::callsite::DefaultCallsite::new(&META)
};
let enabled =
::tracing::Level::INFO <= ::tracing::level_filters::STATIC_MAX_LEVEL
&&
::tracing::Level::INFO <=
::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!("Sparse buffer updates disabled. RenderDevice lacks support: max_storage_buffers_per_shader_stage ({0}) < 3.",
limit) as &dyn ::tracing::field::Value))])
});
} else { ; }
};info!(
284"Sparse buffer updates disabled. RenderDevice lacks support: max_storage_buffers_per_shader_stage ({}) < 3.",
285 limit
286 );
287288return SparseBufferUpdatePipelines {
289 bind_group_layout: None,
290 shader: None,
291 };
292 }
293294let bind_group_layout = BindGroupLayoutDescriptor::new(
295"sparse buffer update bind group layout",
296&BindGroupLayoutEntries::sequential(
297ShaderStages::COMPUTE,
298 (
299// @group(0) @binding(0) var<storage, read_write> dest_buffer: array<u32>;
300storage_buffer::<u32>(false),
301// @group(0) @binding(1) var<storage> src_buffer: array<u32>;
302storage_buffer_read_only::<u32>(false),
303// @group(0) @binding(2) var<storage> indices: array<u32>;
304storage_buffer_read_only::<u32>(false),
305// @group(0) @binding(3) var<uniform> metadata:
306 // SparseBufferUpdateMetadata;
307uniform_buffer::<GpuSparseBufferUpdateMetadata>(false),
308 ),
309 ),
310 );
311312SparseBufferUpdatePipelines {
313 bind_group_layout: Some(bind_group_layout),
314 shader: Some({
let (path, asset_server) =
{
let path =
{
{
let crate_name =
"bevy_render::render_resource::sparse_buffer_vec".split(':').next().unwrap();
::bevy_asset::io::embedded::_embedded_asset_path(crate_name,
"src".as_ref(),
"src/render_resource/sparse_buffer_vec.rs".as_ref(),
"sparse_buffer_update.wgsl".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(world);
(path, asset_server)
};
asset_server.load(path)
}load_embedded_asset!(world, "sparse_buffer_update.wgsl")),
315 }
316 }
317}
318319impl SpecializedComputePipelinefor SparseBufferUpdatePipelines {
320type Key = ();
321322fn specialize(&self, _: Self::Key) -> ComputePipelineDescriptor {
323ComputePipelineDescriptor {
324 label: Some("sparse buffer update pipeline".into()),
325 layout: self.bind_group_layout.clone().into_iter().collect(),
326 shader: self.shader.clone().unwrap_or_default(),
327 shader_defs: ::alloc::vec::Vec::new()vec![],
328 ..ComputePipelineDescriptor::default()
329 }
330 }
331}
332333/// The buffers that we use to sparsely scatter new data to the GPU.
334///
335/// There's one such set of buffers per sparse buffer vector.
336struct SparseBufferStagingBuffers {
337/// All pages that have changed and need to be updated.
338source_data: RawBufferVec<u32>,
339340/// The index at which we write each page in [`Self::source_data`].
341 ///
342 /// The length of this buffer is equal to [`Self::source_data`] divided by
343 /// 2^[`Self::page_size_log2`].
344indices: RawBufferVec<u32>,
345346/// The size of each element in 32-bit words.
347element_word_size: u32,
348349/// The base-2 logarithm of the page size in elements.
350 ///
351 /// That is, the page size in elements is `1 << page_size_log2`.
352page_size_log2: u32,
353}
354355impl SparseBufferStagingBuffers {
356/// The number of elements per page.
357fn page_size(&self) -> usize {
3581 << self.page_size_log2
359 }
360361/// Creates a new set of staging buffers for a sparse buffer vector.
362fn new(label: &str, element_word_size: u32, page_size_log2: u32) -> SparseBufferStagingBuffers {
363let mut source_data_buffer =
364RawBufferVec::new(BufferUsages::COPY_DST | BufferUsages::STORAGE);
365source_data_buffer.set_label(Some(&*::alloc::__export::must_use({
::alloc::fmt::format(format_args!("{0} staging buffer", label))
})format!("{} staging buffer", label)));
366367let mut indices_buffer = RawBufferVec::new(BufferUsages::COPY_DST | BufferUsages::STORAGE);
368indices_buffer.set_label(Some(&*::alloc::__export::must_use({
::alloc::fmt::format(format_args!("{0} index buffer", label))
})format!("{} index buffer", label)));
369370SparseBufferStagingBuffers {
371 source_data: source_data_buffer,
372 indices: indices_buffer,
373element_word_size,
374page_size_log2,
375 }
376 }
377378/// Returns the number of updated pages.
379fn updated_page_count(&self) -> u32 {
380// Note that we don't have to round up here because data is always
381 // uploaded in increments of a whole page.
382let element_count = self.source_data.len() / self.element_word_size as usize;
383 (element_count / self.page_size()) as u32384 }
385386/// Writes the buffers that contain all the data necessary to perform a
387 /// sparse upload to the GPU.
388 ///
389 /// This includes the buffer associated with the supplied
390 /// `metadata_uniform`.
391fn write_buffers(
392&mut self,
393 metadata_uniform: &mut UniformBuffer<GpuSparseBufferUpdateMetadata>,
394 render_device: &RenderDevice,
395 render_queue: &RenderQueue,
396 ) {
397metadata_uniform.get_mut().updated_page_count = self.updated_page_count();
398metadata_uniform.write_buffer(render_device, render_queue);
399400self.source_data.write_buffer(render_device, render_queue);
401self.indices.write_buffer(render_device, render_queue);
402 }
403404/// Returns true if a sparse buffer update should *not* be performed because
405 /// too many words changed.
406fn should_perform_full_reupload(&self, changed_page_count: u32, buffer_length: usize) -> bool {
407// Calculate the number of changed words. If it's greater than the
408 // maximum number of workgroups as defined by `wgpu`, we must perform a
409 // full reupload.
410let total_changed_word_count =
411changed_page_count * self.page_size() as u32 * self.element_word_size;
412if total_changed_word_count > MAX_WORKGROUPS * SPARSE_BUFFER_UPDATE_WORKGROUP_SIZE {
413return true;
414 }
415416// Don't perform a sparse upload if too many words changed, as it'll end
417 // up being slower than just uploading the whole buffer afresh.
418let sparse_upload_fraction =
419changed_page_countas f64 / buffer_length.div_ceil(self.page_size()) as f64;
420sparse_upload_fraction > SPARSE_UPLOAD_THRESHOLD421 }
422}
423424/// A GPU buffer that can grow, can be updated atomically from multiple threads
425/// on the CPU, and is sparsely updated on the GPU if only a small number of
426/// elements have changed.
427///
428/// This type is similar to
429/// [`crate::render_resource::buffer_vec::AtomicRawBufferVec`], but instead of
430/// reuploading the entire buffer to the GPU when it's changed, it tracks
431/// changes on a per-page level and uploads only the pages that changed if the
432/// number of such pages is small. It uses a compute shader to scatter the
433/// changed pages.
434///
435/// As the stored data is [`AtomicPod`], multiple threads may update the buffer
436/// simultaneously. Note that, like
437/// [`crate::render_resource::buffer_vec::AtomicRawBufferVec`], only existing
438/// elements may be updated from multiple threads; new data still requires
439/// exclusive access.
440///
441/// `T` must have a size that's a multiple of 4.
442pub struct AtomicSparseBufferVec<T>
443where
444T: AtomicPod,
445{
446/// An ID that uniquely identifies this [`AtomicSparseBufferVec`].
447handle: SparseBufferHandle,
448/// The underlying values.
449 ///
450 /// These are stored as their blob representation to allow for thread-safe
451 /// update.
452values: Vec<T::Blob>,
453/// The GPU buffer, if allocated.
454data_buffer: Option<Buffer>,
455/// The GPU buffers that data is copied into in preparation to be scattered
456 /// to the [`Self::data_buffer`].
457staging_buffers: SparseBufferStagingBuffers,
458/// A GPU buffer that stores information such as the element size and stride
459 /// that's needed to perform sparse updates.
460metadata_uniform: UniformBuffer<GpuSparseBufferUpdateMetadata>,
461/// The capacity of the GPU buffer in elements.
462capacity: usize,
463/// The allowed `wgpu` buffer usages for the GPU buffer.
464buffer_usages: BufferUsages,
465/// An optional debug label to identify this buffer.
466label: Arc<str>,
467/// A bit set of dirty pages.
468 ///
469 /// The size of this vector in bits is the number of elements divided by the
470 /// page size, rounded up. A 1 in a bit indicates that the page has changed
471 /// since the last upload, while a 0 indicates that the page hasn't changed.
472dirty_pages: Vec<AtomicU64>,
473/// True if the entire buffer needs to be reuploaded because it resized.
474needs_full_reupload: bool,
475/// True if a sparse update is to be performed.
476sparse_update_scheduled: bool,
477}
478479impl<T> AtomicSparseBufferVec<T>
480where
481T: AtomicPod,
482{
483/// The number of elements per page.
484fn page_size(&self) -> u32 {
4851 << self.staging_buffers.page_size_log2
486 }
487488/// Creates a new [`AtomicSparseBufferVec`] with the given set of buffer
489 /// usages, page size, and label.
490 ///
491 /// `buffer_usages` specifies the set of allowed `wgpu` buffer usages for
492 /// the buffer that [`AtomicSparseBufferVec`] manages.
493 /// `BufferUsages::COPY_DST` is automatically added to this set.
494 ///
495 /// The `page_size_log2` parameter is the base-2 logarithm of the page size.
496 /// That is, the page size is `1 << page_size_log2`.
497pub fn new(buffer_usages: BufferUsages, page_size_log2: u32, label: Arc<str>) -> Self {
498// Make sure the value is word-aligned.
499if true {
{
match (&(size_of::<T>() % 4), &0) {
(left_val, right_val) => {
if !(*left_val == *right_val) {
let kind = ::core::panicking::AssertKind::Eq;
::core::panicking::assert_failed(kind, &*left_val,
&*right_val, ::core::option::Option::None);
}
}
}
};
};debug_assert_eq!(size_of::<T>() % 4, 0);
500let element_word_size = size_of::<T>() / 4;
501502// Create a unique ID.
503let id = Arc::new(SparseBufferId(
504NEXT_SPARSE_BUFFER_ID.fetch_add(1, Ordering::Relaxed),
505 ));
506507Self {
508 handle: id,
509 values: ::alloc::vec::Vec::new()vec![],
510 data_buffer: None,
511 staging_buffers: SparseBufferStagingBuffers::new(
512&label,
513element_word_sizeas u32,
514page_size_log2,
515 ),
516 metadata_uniform: UniformBuffer::from(GpuSparseBufferUpdateMetadata::new::<T>(
517page_size_log2,
518 )),
519 capacity: 0,
520 buffer_usages: buffer_usages | BufferUsages::COPY_DST,
521label,
522 dirty_pages: ::alloc::vec::Vec::new()vec![],
523 needs_full_reupload: false,
524 sparse_update_scheduled: false,
525 }
526 }
527528/// Returns the number of elements in the CPU side copy of the buffer.
529pub fn len(&self) -> u32 {
530self.values.len() as u32531 }
532533/// Returns true if there are no elements in the CPU side copy of the buffer.
534pub fn is_empty(&self) -> bool {
535self.values.is_empty()
536 }
537538/// Returns a handle to the buffer, if the data has been uploaded.
539pub fn buffer(&self) -> Option<&Buffer> {
540self.data_buffer.as_ref()
541 }
542543/// Removes all elements from the buffer.
544pub fn clear(&mut self) {
545self.truncate(0);
546 }
547548/// Copies a value out of the buffer.
549pub fn get(&self, index: u32) -> T {
550 T::read_from_blob(&self.values[indexas usize])
551 }
552553/// Sets the value at the given index.
554 ///
555 /// If the index isn't in range of the buffer, this method panics.
556 ///
557 /// Internally, the value is converted to its blob representation.
558 ///
559 /// Note that this method is thread-safe and doesn't require `&mut self`.
560 /// It's your responsibility, however, to ensure synchronization; though
561 /// this method is memory-safe, it's possible for other threads to observe
562 /// partially-overwritten values if [`Self::get`] or similar methods are
563 /// called while the write operation is occurring.
564pub fn set(&self, index: u32, value: T) {
565value.write_to_blob(&self.values[indexas usize]);
566self.note_changed_index(index);
567 }
568569/// Adds a new value and returns its index.
570pub fn push(&mut self, value: T) -> u32 {
571let index = self.values.len() as u32;
572self.values.push(T::Blob::default());
573value.write_to_blob(&self.values[indexas usize]);
574575let page_word = (self.index_to_page(index) / PAGES_PER_DIRTY_WORD) as usize;
576while self.dirty_pages.len() < page_word + 1 {
577self.dirty_pages.push(AtomicU64::default());
578 }
579self.note_changed_index(index);
580581index582 }
583584/// Marks the page corresponding to the given element index as dirty so that
585 /// we know that we need to upload it.
586fn note_changed_index(&self, index: u32) {
587let page = self.index_to_page(index);
588let (page_word, page_in_word) = (page / PAGES_PER_DIRTY_WORD, page % PAGES_PER_DIRTY_WORD);
589self.dirty_pages[page_wordas usize].fetch_or(1 << page_in_word, Ordering::Relaxed);
590 }
591592/// Returns the page corresponding to the given element index.
593fn index_to_page(&self, index: u32) -> u32 {
594index / self.page_size()
595 }
596597/// Ensures that the backing buffer for this buffer vector is present and
598 /// appropriately sized on the GPU.
599pub fn reserve(&mut self, new_capacity: usize, render_device: &RenderDevice) {
600reserve(
601new_capacity,
602&mut self.capacity,
603&self.label,
604&mut self.data_buffer,
605self.buffer_usages,
606&mut self.needs_full_reupload,
607size_of::<T::Blob>(),
608render_device,
609 );
610 }
611612/// Grows the buffer by adding default values so that it's at least the
613 /// given size.
614 ///
615 /// If the buffer is already large enough, this method does nothing.
616pub fn grow(&mut self, new_len: u32) {
617let old_len = self.values.len() as u32;
618if old_len >= new_len {
619return;
620 }
621622self.values.reserve(new_lenas usize - old_lenas usize);
623self.values.resize_with(new_lenas usize, T::Blob::default);
624625// This is a bit tricky. We want to set the dirty bits corresponding to
626 // all pages that we added, if any. First, we compute the index of the
627 // last page word before the append operation.
628let old_final_page = self.index_to_page(old_len);
629let old_final_page_word_index = old_final_page / PAGES_PER_DIRTY_WORD;
630let old_final_page_in_word = old_final_page % PAGES_PER_DIRTY_WORD;
631632// Next, we set the bits corresponding to every page that we added to
633 // that final page word. Note that this might set bits corresponding to
634 // pages past the end of our buffer; that's OK as we ignore them.
635if old_final_page_in_word != 0
636&& let Some(ref mut old_final_atomic_page_word) =
637self.dirty_pages.get_mut(old_final_page_word_indexas usize)
638 {
639*old_final_atomic_page_word.get_mut() |= !((1u64 << old_final_page_in_word) - 1);
640 }
641642// Finally, we add any new page words, with all bits set.
643let new_page_count = self.index_to_page(new_len);
644self.dirty_pages.resize_with(
645 (new_page_countas usize).div_ceil(PAGES_PER_DIRTY_WORDas usize),
646 || AtomicU64::new(u64::MAX),
647 );
648 }
649650/// Truncates the buffer to the given length.
651 ///
652 /// If the buffer is already that length or shorter, this method does
653 /// nothing.
654pub fn truncate(&mut self, len: u32) {
655self.values.truncate(lenas usize);
656657let page = self.index_to_page(len);
658self.dirty_pages
659 .truncate(page.div_ceil(PAGES_PER_DIRTY_WORD) as usize);
660 }
661662/// Writes the data to the GPU, either via a sparse upload or a bulk data
663 /// upload.
664pub fn write_buffers(&mut self, render_device: &RenderDevice, render_queue: &RenderQueue) {
665if self.values.is_empty() {
666return;
667 }
668669// Round up the size to a good value to balance reallocation frequency
670 // against memory waste.
671let good_size = calculate_allocation_size(self.values.len());
672self.reserve(good_size, render_device);
673674if self.should_perform_full_reupload(render_device) {
675self.write_entire_buffer(render_queue);
676 } else {
677self.prepare_sparse_upload(render_device, render_queue);
678 }
679 }
680681/// Returns true if the sparse buffer should perform a full reupload, either
682 /// because it was resized or because too much data changed for a sparse
683 /// update to be worthwhile.
684fn should_perform_full_reupload(&self, render_device: &RenderDevice) -> bool {
685if self.needs_full_reupload {
686return true;
687 }
688689if render_device.limits().max_storage_buffers_per_shader_stage < 3 {
690return true;
691 }
692693// Calculate the number of changed pages via population count.
694let changed_page_count: u32 = self695 .dirty_pages
696 .iter()
697 .map(|atomic_page_word| atomic_page_word.load(Ordering::Relaxed).count_ones())
698 .sum();
699700self.staging_buffers
701 .should_perform_full_reupload(changed_page_count, self.values.len())
702 }
703704/// Writes the entire buffer in bulk.
705 ///
706 /// This is the method used when a sparse update is not used, either because
707 /// the buffer resized or because too much data changed for a sparse update
708 /// to be worthwhile.
709fn write_entire_buffer(&mut self, render_queue: &RenderQueue) {
710let Some(ref mut data_buffer) = self.data_buffer else {
711{
use ::tracing::__macro_support::Callsite as _;
static __CALLSITE: ::tracing::callsite::DefaultCallsite =
{
static META: ::tracing::Metadata<'static> =
{
::tracing_core::metadata::Metadata::new("event src/render_resource/sparse_buffer_vec.rs:711",
"bevy_render::render_resource::sparse_buffer_vec",
::tracing::Level::ERROR,
::tracing_core::__macro_support::Option::Some("src/render_resource/sparse_buffer_vec.rs"),
::tracing_core::__macro_support::Option::Some(711u32),
::tracing_core::__macro_support::Option::Some("bevy_render::render_resource::sparse_buffer_vec"),
::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!("Dirty sparse buffer should have created a data buffer by now")
as &dyn ::tracing::field::Value))])
});
} else { ; }
};error!("Dirty sparse buffer should have created a data buffer by now");
712return;
713 };
714715// SAFETY: We're just writing atomic data to the GPU. The worst that
716 // can happen is that we race with somebody, which is unfortunate
717 // but not memory-unsafe.
718unsafe {
719render_queue.write_buffer(
720data_buffer,
7210,
722 slice::from_raw_parts(
723self.values.as_ptr().cast::<u8>(),
724self.values.len() * size_of::<T::Blob>(),
725 ),
726 );
727 }
728729// Mark all pages as clean.
730for atomic_page_word in self.dirty_pages.iter() {
731 atomic_page_word.store(0, Ordering::Relaxed);
732 }
733self.sparse_update_scheduled = false;
734 }
735736/// Schedules a sparse upload of only the pages that changed.
737fn prepare_sparse_upload(&mut self, render_device: &RenderDevice, render_queue: &RenderQueue) {
738// Iterate over all dirty pages.
739for (page_word_index, atomic_page_word) in self.dirty_pages.iter().enumerate() {
740let page_word = atomic_page_word.load(Ordering::Relaxed);
741for page_index_in_word in BitIter::new(page_word) {
742let page = page_word_index as u32 * PAGES_PER_DIRTY_WORD + page_index_in_word;
743744// Write the index of the page so the shader will know where to
745 // scatter the data to.
746self.staging_buffers.indices.push(page);
747748// Copy the page to the GPU staging buffer.
749let page_size = self.staging_buffers.page_size();
750let page_start = page as usize * page_size;
751let page_end = page_start + page_size;
752for value_index in page_start..page_end {
753match self.values.get(value_index) {
754Some(blob) => {
755let value = T::read_from_blob(blob);
756self.staging_buffers
757 .source_data
758 .extend(bytemuck::cast_slice(&[value]).iter().copied());
759 }
760None => {
761self.staging_buffers.source_data.extend(iter::repeat_n(
7620,
763self.staging_buffers.element_word_size as usize,
764 ));
765 }
766 }
767 }
768769// Make sure we're aligned up to a full page.
770if true {
{
match (&(self.staging_buffers.source_data.len() %
(self.staging_buffers.element_word_size as usize *
self.staging_buffers.page_size())), &0) {
(left_val, right_val) => {
if !(*left_val == *right_val) {
let kind = ::core::panicking::AssertKind::Eq;
::core::panicking::assert_failed(kind, &*left_val,
&*right_val, ::core::option::Option::None);
}
}
}
};
};debug_assert_eq!(
771self.staging_buffers.source_data.len()
772 % (self.staging_buffers.element_word_size as usize
773 * self.staging_buffers.page_size()),
7740
775);
776 }
777778// Mark the page as clean.
779atomic_page_word.store(0, Ordering::Relaxed);
780 }
781782// Schedule a sparse update if there was something to do.
783self.sparse_update_scheduled = !self.staging_buffers.source_data.is_empty();
784if self.sparse_update_scheduled {
785self.staging_buffers.write_buffers(
786&mut self.metadata_uniform,
787render_device,
788render_queue,
789 );
790 }
791 }
792793/// If a sparse update has been scheduled, prepares all GPU resources
794 /// necessary to perform a sparse buffer update, other than updating the
795 /// metadata uniform.
796pub fn prepare_to_populate_buffers(
797&mut self,
798 render_device: &RenderDevice,
799 pipeline_cache: &PipelineCache,
800 sparse_buffer_update_jobs: &mut SparseBufferUpdateJobs,
801 sparse_buffer_update_bind_groups: &mut SparseBufferUpdateBindGroups,
802 sparse_buffer_update_pipelines: &SparseBufferUpdatePipelines,
803 ) {
804if self.sparse_update_scheduled {
805match (&self.data_buffer, self.metadata_uniform.buffer()) {
806 (Some(data_buffer), Some(metadata_buffer)) => {
807prepare_to_populate_buffers(
808self.handle.clone(),
809&self.label,
810data_buffer,
811&mut self.staging_buffers,
812metadata_buffer,
813render_device,
814pipeline_cache,
815sparse_buffer_update_jobs,
816sparse_buffer_update_bind_groups,
817sparse_buffer_update_pipelines,
818 );
819 }
820_ => {
821{
use ::tracing::__macro_support::Callsite as _;
static __CALLSITE: ::tracing::callsite::DefaultCallsite =
{
static META: ::tracing::Metadata<'static> =
{
::tracing_core::metadata::Metadata::new("event src/render_resource/sparse_buffer_vec.rs:821",
"bevy_render::render_resource::sparse_buffer_vec",
::tracing::Level::ERROR,
::tracing_core::__macro_support::Option::Some("src/render_resource/sparse_buffer_vec.rs"),
::tracing_core::__macro_support::Option::Some(821u32),
::tracing_core::__macro_support::Option::Some("bevy_render::render_resource::sparse_buffer_vec"),
::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!("Buffers should have been created by now")
as &dyn ::tracing::field::Value))])
});
} else { ; }
};error!("Buffers should have been created by now");
822 }
823 }
824 }
825826// Clear out the staging buffers, now that we know the data is already
827 // on the GPU.
828self.staging_buffers.source_data.clear();
829self.staging_buffers.indices.clear();
830831// Reset the `needs_full_reupload` and `needs_sparse_update` flags.
832self.needs_full_reupload = false;
833self.sparse_update_scheduled = false;
834 }
835}
836837impl FromWorldfor SparseBufferUpdateBindGroups {
838fn from_world(world: &mut World) -> Self {
839world.resource_scope::<SpecializedComputePipelines<SparseBufferUpdatePipelines>, _>(
840 |world, mut specialized_sparse_buffer_update_pipelines| {
841let pipeline_cache = world.resource::<PipelineCache>();
842let sparse_buffer_update_pipelines =
843world.resource::<SparseBufferUpdatePipelines>();
844let pipeline_id = specialized_sparse_buffer_update_pipelines.specialize(
845pipeline_cache,
846sparse_buffer_update_pipelines,
847 (),
848 );
849850SparseBufferUpdateBindGroups {
851 bind_groups: WeakKeyHashMap::default(),
852pipeline_id,
853 }
854 },
855 )
856 }
857}
858859/// Prepares all GPU resources necessary to perform a sparse buffer update,
860/// other than updating the metadata uniform.
861///
862/// This function creates the [`SparseBufferUpdateJob`] and ensures the bind
863/// group and pipeline are up to date.
864fn prepare_to_populate_buffers(
865 sparse_buffer_handle: SparseBufferHandle,
866 label: &Arc<str>,
867 data_buffer: &Buffer,
868 staging_buffers: &mut SparseBufferStagingBuffers,
869 metadata_buffer: &Buffer,
870 render_device: &RenderDevice,
871 pipeline_cache: &PipelineCache,
872 sparse_buffer_update_jobs: &mut SparseBufferUpdateJobs,
873 sparse_buffer_update_bind_groups: &mut SparseBufferUpdateBindGroups,
874 sparse_buffer_update_pipelines: &SparseBufferUpdatePipelines,
875) {
876let (Some(source_data_staging_buffer), Some(indices_staging_buffer)) = (
877staging_buffers.source_data.buffer(),
878staging_buffers.indices.buffer(),
879 ) else {
880{
use ::tracing::__macro_support::Callsite as _;
static __CALLSITE: ::tracing::callsite::DefaultCallsite =
{
static META: ::tracing::Metadata<'static> =
{
::tracing_core::metadata::Metadata::new("event src/render_resource/sparse_buffer_vec.rs:880",
"bevy_render::render_resource::sparse_buffer_vec",
::tracing::Level::ERROR,
::tracing_core::__macro_support::Option::Some("src/render_resource/sparse_buffer_vec.rs"),
::tracing_core::__macro_support::Option::Some(880u32),
::tracing_core::__macro_support::Option::Some("bevy_render::render_resource::sparse_buffer_vec"),
::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!("Staging buffers should have been created by now")
as &dyn ::tracing::field::Value))])
});
} else { ; }
};error!("Staging buffers should have been created by now");
881return;
882 };
883884let Some(bind_group_layout) = &sparse_buffer_update_pipelines.bind_group_layout else {
885return;
886 };
887888// Record the update job.
889sparse_buffer_update_jobs.push(SparseBufferUpdateJob {
890 sparse_buffer_handle: sparse_buffer_handle.clone(),
891 page_size_log2: staging_buffers.page_size_log2,
892 updated_page_count: staging_buffers.updated_page_count(),
893 element_word_size: staging_buffers.element_word_size,
894 label: (*label).clone(),
895 });
896897// Create the bind group.
898let bind_group = render_device.create_bind_group(
899Some(&*::alloc::__export::must_use({
::alloc::fmt::format(format_args!("{0} bind group", label))
})format!("{} bind group", label)),
900&pipeline_cache.get_bind_group_layout(bind_group_layout),
901&BindGroupEntries::sequential((
902// @group(0) @binding(0) var<storage, read_write> dest_buffer: array<u32>;
903data_buffer.as_entire_binding(),
904// @group(0) @binding(1) var<storage> src_buffer: array<u32>;
905source_data_staging_buffer.as_entire_binding(),
906// @group(0) @binding(2) var<storage> indices: array<u32>;
907indices_staging_buffer.as_entire_binding(),
908// @group(0) @binding(3) var<uniform> metadata:
909 // SparseBufferUpdateMetadata;
910metadata_buffer.as_entire_binding(),
911 )),
912 );
913sparse_buffer_update_bind_groups.bind_groups.insert(
914sparse_buffer_handle,
915SparseBufferUpdateBindGroup { bind_group },
916 );
917}
918919/// Ensures that the backing buffer for an [`AtomicSparseBufferVec`] is present
920/// on the GPU.
921///
922/// The `capacity`, `data_buffer`, and `needs_full_reupload` fields are updated
923/// to reflect the new buffer.
924fn reserve(
925 new_capacity: usize,
926 capacity: &mut usize,
927 label: &str,
928 data_buffer: &mut Option<Buffer>,
929 buffer_usages: BufferUsages,
930 needs_full_reupload: &mut bool,
931 element_size: usize,
932 render_device: &RenderDevice,
933) {
934// If the buffer is already big enough, do nothing.
935if new_capacity == 0 || new_capacity <= *capacity {
936return;
937 }
938939*capacity = new_capacity;
940*data_buffer = Some(render_device.create_buffer(&BufferDescriptor {
941 label: Some(label),
942 size: element_sizeas u64 * new_capacityas u64,
943 usage: buffer_usages,
944 mapped_at_creation: false,
945 }));
946947// Since we resized the buffer, we need to reupload it.
948*needs_full_reupload = true;
949}
950951impl GpuSparseBufferUpdateMetadata {
952/// Returns a new [`GpuSparseBufferUpdateMetadata`] for the given type and
953 /// page size.
954fn new<T>(page_size_log2: u32) -> GpuSparseBufferUpdateMetadata {
955{
match (&(size_of::<T>() % 4), &0) {
(left_val, right_val) => {
if !(*left_val == *right_val) {
let kind = ::core::panicking::AssertKind::Eq;
::core::panicking::assert_failed(kind, &*left_val,
&*right_val, ::core::option::Option::None);
}
}
}
};assert_eq!(size_of::<T>() % 4, 0);
956GpuSparseBufferUpdateMetadata {
957 element_size: (size_of::<T>() / 4) as u32,
958 updated_page_count: 0,
959page_size_log2,
960 }
961 }
962}
963964/// Iterates over the bits in a single `u64`, from the least significant bit to
965/// the most significant bit.
966struct BitIter(u64);
967968impl BitIter {
969fn new(bits: u64) -> BitIter {
970BitIter(bits)
971 }
972}
973974impl Iteratorfor BitIter {
975type Item = u32;
976977fn next(&mut self) -> Option<Self::Item> {
978let trailing_zeros = self.0.trailing_zeros();
979if trailing_zeros == 64 {
980return None;
981 }
982self.0 &= !(1 << trailing_zeros);
983Some(trailing_zeros)
984 }
985}
986987/// Calculates the size that a buffer should be in order to balance reallocation
988/// frequency against memory waste.
989fn calculate_allocation_size(length: usize) -> usize {
990let exponent = (lengthas f64).log(REALLOCATION_FACTOR).ceil();
991let size = REALLOCATION_FACTOR.powf(exponent) as usize;
992size.next_multiple_of(REALLOCATION_SIZE_MULTIPLE)
993}