From c9202ee54a1a7b4aa152e9e68cc051d75fc46127 Mon Sep 17 00:00:00 2001 From: Elie Michel Date: Wed, 2 Oct 2024 10:21:30 +0200 Subject: [PATCH 01/19] Remove redundant let binding (#6356) --- wgpu-core/src/device/resource.rs | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs index 03b183e085..9f0cbf33f1 100644 --- a/wgpu-core/src/device/resource.rs +++ b/wgpu-core/src/device/resource.rs @@ -437,13 +437,11 @@ impl Device { .last_successful_submission_index .load(Ordering::Acquire); - if let wgt::Maintain::WaitForSubmissionIndex(submission_index) = maintain { - if submission_index > last_successful_submission_index { - return Err(WaitIdleError::WrongSubmissionIndex( - submission_index, - last_successful_submission_index, - )); - } + if submission_index > last_successful_submission_index { + return Err(WaitIdleError::WrongSubmissionIndex( + submission_index, + last_successful_submission_index, + )); } submission_index From 025787bbadb28464a630373604e796d93fbb490a Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Thu, 18 Jul 2024 20:02:38 +0200 Subject: [PATCH 02/19] remove duplicate validation --- wgpu-core/src/binding_model.rs | 4 ---- 1 file changed, 4 deletions(-) diff --git a/wgpu-core/src/binding_model.rs b/wgpu-core/src/binding_model.rs index 2357a8c776..6f43ac5234 100644 --- a/wgpu-core/src/binding_model.rs +++ b/wgpu-core/src/binding_model.rs @@ -373,10 +373,6 @@ impl BindingTypeMaxCountValidator { limits.max_sampled_textures_per_shader_stage, BindingTypeMaxCountErrorKind::SampledTextures, )?; - self.storage_buffers.validate( - limits.max_storage_buffers_per_shader_stage, - BindingTypeMaxCountErrorKind::StorageBuffers, - )?; self.samplers.validate( limits.max_samplers_per_shader_stage, BindingTypeMaxCountErrorKind::Samplers, From 76af20348d0fefaffc734339dc3f295c51ebebb4 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Fri, 19 Jul 2024 11:26:40 +0200 Subject: [PATCH 03/19] remove old comment --- wgpu-core/src/command/compute.rs | 2 -- 1 file changed, 2 deletions(-) diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs index 5de03917d2..7f396d11b6 100644 --- a/wgpu-core/src/command/compute.rs +++ b/wgpu-core/src/command/compute.rs @@ -495,8 +495,6 @@ impl Global { state.raw_encoder.begin_compute_pass(&hal_desc); } - // TODO: We should be draining the commands here, avoiding extra copies in the process. - // (A command encoder can't be executed twice!) for command in base.commands { match command { ArcComputeCommand::SetBindGroup { From 0903ba6d47e952ad16c76f2fde0403b959a99a6d Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 19 Sep 2024 15:02:56 -0400 Subject: [PATCH 04/19] refactor(core): delete unused `CreateDeviceError` --- wgpu-core/src/device/resource.rs | 9 --------- 1 file changed, 9 deletions(-) diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs index 9f0cbf33f1..18fdf5708b 100644 --- a/wgpu-core/src/device/resource.rs +++ b/wgpu-core/src/device/resource.rs @@ -38,7 +38,6 @@ use arrayvec::ArrayVec; use once_cell::sync::OnceCell; use smallvec::SmallVec; -use thiserror::Error; use wgt::{ math::align_to, DeviceLostReason, TextureFormat, TextureSampleType, TextureViewDimension, }; @@ -187,14 +186,6 @@ impl Drop for Device { } } -#[derive(Clone, Debug, Error)] -pub enum CreateDeviceError { - #[error("Not enough memory left to create device")] - OutOfMemory, - #[error("Failed to create internal buffer for initializing textures")] - FailedToCreateZeroBuffer(#[from] DeviceError), -} - impl Device { pub(crate) fn raw(&self) -> &dyn hal::DynDevice { self.raw.as_ref() From 7ac533a3128497ea5c06d4062f83154cfa832b64 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Thu, 18 Jul 2024 21:00:14 +0200 Subject: [PATCH 05/19] add missing indirect buffer offset validation --- wgpu-core/src/command/compute.rs | 6 ++++++ wgpu-core/src/command/render.rs | 10 ++++++++++ 2 files changed, 16 insertions(+) diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs index 7f396d11b6..1540a5f504 100644 --- a/wgpu-core/src/command/compute.rs +++ b/wgpu-core/src/command/compute.rs @@ -136,6 +136,8 @@ pub enum ComputePassErrorInner { BindGroupIndexOutOfRange { index: u32, max: u32 }, #[error(transparent)] DestroyedResource(#[from] DestroyedResourceError), + #[error("Indirect buffer offset {0:?} is not a multiple of 4")] + UnalignedIndirectBufferOffset(BufferAddress), #[error("Indirect buffer uses bytes {offset}..{end_offset} which overruns indirect buffer of size {buffer_size}")] IndirectBufferOverrun { offset: u64, @@ -845,6 +847,10 @@ fn dispatch_indirect( .merge_single(&buffer, hal::BufferUses::INDIRECT)?; buffer.check_usage(wgt::BufferUsages::INDIRECT)?; + if offset % 4 != 0 { + return Err(ComputePassErrorInner::UnalignedIndirectBufferOffset(offset)); + } + let end_offset = offset + size_of::() as u64; if end_offset > buffer.size { return Err(ComputePassErrorInner::IndirectBufferOverrun { diff --git a/wgpu-core/src/command/render.rs b/wgpu-core/src/command/render.rs index b6680333c2..d22eb5f0d6 100644 --- a/wgpu-core/src/command/render.rs +++ b/wgpu-core/src/command/render.rs @@ -638,6 +638,8 @@ pub enum RenderPassErrorInner { MissingFeatures(#[from] MissingFeatures), #[error(transparent)] MissingDownlevelFlags(#[from] MissingDownlevelFlags), + #[error("Indirect buffer offset {0:?} is not a multiple of 4")] + UnalignedIndirectBufferOffset(BufferAddress), #[error("Indirect draw uses bytes {offset}..{end_offset} {} which overruns indirect buffer of size {buffer_size}", count.map_or_else(String::new, |v| format!("(using count {v})")))] IndirectBufferOverrun { @@ -2450,6 +2452,10 @@ fn multi_draw_indirect( let actual_count = count.map_or(1, |c| c.get()); + if offset % 4 != 0 { + return Err(RenderPassErrorInner::UnalignedIndirectBufferOffset(offset)); + } + let end_offset = offset + stride as u64 * actual_count as u64; if end_offset > indirect_buffer.size { return Err(RenderPassErrorInner::IndirectBufferOverrun { @@ -2534,6 +2540,10 @@ fn multi_draw_indirect_count( count_buffer.check_usage(BufferUsages::INDIRECT)?; let count_raw = count_buffer.try_raw(state.snatch_guard)?; + if offset % 4 != 0 { + return Err(RenderPassErrorInner::UnalignedIndirectBufferOffset(offset)); + } + let end_offset = offset + stride * max_count as u64; if end_offset > indirect_buffer.size { return Err(RenderPassErrorInner::IndirectBufferOverrun { From 38a13b94f0a779727a21b37cede6c5e1c7bcb161 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Fri, 19 Jul 2024 10:23:13 +0200 Subject: [PATCH 06/19] refactor: make `Snatchable::snatch` take `_guard` by `&mut _` --- wgpu-core/src/device/resource.rs | 5 +++-- wgpu-core/src/present.rs | 14 ++++++++++---- wgpu-core/src/resource.rs | 6 ++---- wgpu-core/src/snatch.rs | 2 +- 4 files changed, 16 insertions(+), 11 deletions(-) diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs index 18fdf5708b..5a726ef70f 100644 --- a/wgpu-core/src/device/resource.rs +++ b/wgpu-core/src/device/resource.rs @@ -363,7 +363,7 @@ impl Device { let Some(view) = view.upgrade() else { continue; }; - let Some(raw_view) = view.raw.snatch(self.snatchable_lock.write()) else { + let Some(raw_view) = view.raw.snatch(&mut self.snatchable_lock.write()) else { continue; }; @@ -377,7 +377,8 @@ impl Device { let Some(bind_group) = bind_group.upgrade() else { continue; }; - let Some(raw_bind_group) = bind_group.raw.snatch(self.snatchable_lock.write()) + let Some(raw_bind_group) = + bind_group.raw.snatch(&mut self.snatchable_lock.write()) else { continue; }; diff --git a/wgpu-core/src/present.rs b/wgpu-core/src/present.rs index c9d0124bf4..b12e35328a 100644 --- a/wgpu-core/src/present.rs +++ b/wgpu-core/src/present.rs @@ -288,8 +288,11 @@ impl Global { .textures .remove(texture.tracker_index()); let suf = surface.raw(device.backend()).unwrap(); - let exclusive_snatch_guard = device.snatchable_lock.write(); - match texture.inner.snatch(exclusive_snatch_guard).unwrap() { + match texture + .inner + .snatch(&mut device.snatchable_lock.write()) + .unwrap() + { resource::TextureInner::Surface { raw, parent_id } => { if surface_id != parent_id { log::error!("Presented frame is from a different surface"); @@ -359,8 +362,11 @@ impl Global { .textures .remove(texture.tracker_index()); let suf = surface.raw(device.backend()); - let exclusive_snatch_guard = device.snatchable_lock.write(); - match texture.inner.snatch(exclusive_snatch_guard).unwrap() { + match texture + .inner + .snatch(&mut device.snatchable_lock.write()) + .unwrap() + { resource::TextureInner::Surface { raw, parent_id } => { if surface_id == parent_id { unsafe { suf.unwrap().discard_texture(raw) }; diff --git a/wgpu-core/src/resource.rs b/wgpu-core/src/resource.rs index 5df285da54..32fde3dd57 100644 --- a/wgpu-core/src/resource.rs +++ b/wgpu-core/src/resource.rs @@ -737,8 +737,7 @@ impl Buffer { let device = &self.device; let temp = { - let snatch_guard = device.snatchable_lock.write(); - let raw = match self.raw.snatch(snatch_guard) { + let raw = match self.raw.snatch(&mut device.snatchable_lock.write()) { Some(raw) => raw, None => { return Err(DestroyError::AlreadyDestroyed); @@ -1185,8 +1184,7 @@ impl Texture { let device = &self.device; let temp = { - let snatch_guard = device.snatchable_lock.write(); - let raw = match self.inner.snatch(snatch_guard) { + let raw = match self.inner.snatch(&mut device.snatchable_lock.write()) { Some(TextureInner::Native { raw }) => raw, Some(TextureInner::Surface { .. }) => { return Ok(()); diff --git a/wgpu-core/src/snatch.rs b/wgpu-core/src/snatch.rs index 9866b77723..a817e2068c 100644 --- a/wgpu-core/src/snatch.rs +++ b/wgpu-core/src/snatch.rs @@ -38,7 +38,7 @@ impl Snatchable { } /// Take the value. Requires a the snatchable lock's write guard. - pub fn snatch(&self, _guard: ExclusiveSnatchGuard) -> Option { + pub fn snatch(&self, _guard: &mut ExclusiveSnatchGuard) -> Option { unsafe { (*self.value.get()).take() } } From 71b4f361e244d8b4be79c67207e19a2c267c73f2 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Fri, 19 Jul 2024 10:23:13 +0200 Subject: [PATCH 07/19] chore(core): more detail on `TODO` for push constants --- wgpu-core/src/command/compute.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs index 1540a5f504..133a5af35a 100644 --- a/wgpu-core/src/command/compute.rs +++ b/wgpu-core/src/command/compute.rs @@ -780,7 +780,7 @@ fn set_push_constant( .binder .pipeline_layout .as_ref() - //TODO: don't error here, lazily update the push constants + // TODO: don't error here, lazily update the push constants using `state.push_constants` .ok_or(ComputePassErrorInner::Dispatch( DispatchError::MissingPipeline, ))?; From 2d82054ae4af209c3cbe297039469941b2f3b4e6 Mon Sep 17 00:00:00 2001 From: sagudev <16504129+sagudev@users.noreply.github.com> Date: Sun, 1 Sep 2024 16:12:06 +0200 Subject: [PATCH 08/19] [wgsl-in, spv-out] Allow dynamic indexing of arrays by value. Bring https://github.com/gfx-rs/naga/pull/723 back from the dead. Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> Co-authored-by: Dzmitry Malyshau Co-authored-by: Jim Blandy --- naga/src/back/spv/block.rs | 28 +- naga/src/back/spv/mod.rs | 1 + naga/src/back/spv/writer.rs | 53 ++++ naga/src/lib.rs | 21 +- naga/src/proc/mod.rs | 8 +- naga/src/proc/typifier.rs | 7 + naga/src/valid/expression.rs | 7 +- naga/tests/in/access.wgsl | 11 + naga/tests/out/analysis/access.info.ron | 186 ++++++++++++ .../access.assign_through_ptr.Compute.glsl | 12 +- naga/tests/out/glsl/access.foo.Vertex.glsl | 52 ++++ .../out/glsl/access.foo_frag.Fragment.glsl | 12 +- .../out/glsl/access.foo_vert.Vertex.glsl | 16 +- naga/tests/out/hlsl/access.hlsl | 24 +- naga/tests/out/hlsl/access.ron | 4 + naga/tests/out/ir/access.compact.ron | 117 ++++++++ naga/tests/out/ir/access.ron | 117 ++++++++ naga/tests/out/msl/access.msl | 33 ++- naga/tests/out/spv/access.spvasm | 273 ++++++++++-------- naga/tests/out/wgsl/access.wgsl | 23 +- naga/tests/wgsl_errors.rs | 5 - 21 files changed, 837 insertions(+), 173 deletions(-) create mode 100644 naga/tests/out/glsl/access.foo.Vertex.glsl diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index f0c3bfa848..97db1cb549 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -343,6 +343,32 @@ impl<'w> BlockContext<'w> { load_id } + crate::TypeInner::Array { + base: ty_element, .. + } => { + let index_id = self.cached[index]; + let base_id = self.cached[base]; + let base_ty = match self.fun_info[base].ty { + TypeResolution::Handle(handle) => handle, + TypeResolution::Value(_) => { + return Err(Error::Validation( + "Array types should always be in the arena", + )) + } + }; + let (id, variable) = self.writer.promote_access_expression_to_variable( + &self.ir_module.types, + result_type_id, + base_id, + base_ty, + index_id, + ty_element, + block, + )?; + self.function.internal_variables.push(variable); + id + } + // wgpu#4337: Support `crate::TypeInner::Matrix` ref other => { log::error!( "Unable to access base {:?} of type {:?}", @@ -350,7 +376,7 @@ impl<'w> BlockContext<'w> { other ); return Err(Error::Validation( - "only vectors may be dynamically indexed by value", + "only vectors and arrays may be dynamically indexed by value", )); } } diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 32bd1fcecf..77f915abb1 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -144,6 +144,7 @@ struct Function { signature: Option, parameters: Vec, variables: crate::FastHashMap, LocalVariable>, + internal_variables: Vec, blocks: Vec, entry_point_context: Option, } diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 678dcb4246..7f41caf295 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -32,6 +32,9 @@ impl Function { for local_var in self.variables.values() { local_var.instruction.to_words(sink); } + for internal_var in self.internal_variables.iter() { + internal_var.instruction.to_words(sink); + } } for instruction in block.body.iter() { instruction.to_words(sink); @@ -135,6 +138,56 @@ impl Writer { self.capabilities_used.insert(spirv::Capability::Shader); } + #[allow(clippy::too_many_arguments)] + pub(super) fn promote_access_expression_to_variable( + &mut self, + ir_types: &UniqueArena, + result_type_id: Word, + container_id: Word, + container_ty: Handle, + index_id: Word, + element_ty: Handle, + block: &mut Block, + ) -> Result<(Word, LocalVariable), Error> { + let pointer_type_id = + self.get_pointer_id(ir_types, container_ty, spirv::StorageClass::Function)?; + + let variable = { + let id = self.id_gen.next(); + LocalVariable { + id, + instruction: Instruction::variable( + pointer_type_id, + id, + spirv::StorageClass::Function, + None, + ), + } + }; + block + .body + .push(Instruction::store(variable.id, container_id, None)); + + let element_pointer_id = self.id_gen.next(); + let element_pointer_type_id = + self.get_pointer_id(ir_types, element_ty, spirv::StorageClass::Function)?; + block.body.push(Instruction::access_chain( + element_pointer_type_id, + element_pointer_id, + variable.id, + &[index_id], + )); + let id = self.id_gen.next(); + block.body.push(Instruction::load( + result_type_id, + id, + element_pointer_id, + None, + )); + + Ok((id, variable)) + } + /// Indicate that the code requires any one of the listed capabilities. /// /// If nothing in `capabilities` appears in the available capabilities diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 85fd7a4508..038e215a6a 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -1402,21 +1402,20 @@ pub enum Expression { /// ## Dynamic indexing restrictions /// /// To accommodate restrictions in some of the shader languages that Naga - /// targets, it is not permitted to subscript a matrix or array with a - /// dynamically computed index unless that matrix or array appears behind a - /// pointer. In other words, if the inner type of `base` is [`Array`] or - /// [`Matrix`], then `index` must be a constant. But if the type of `base` - /// is a [`Pointer`] to an array or matrix or a [`ValuePointer`] with a - /// `size`, then the index may be any expression of integer type. + /// targets, it is not permitted to subscript a matrix with a dynamically + /// computed index unless that matrix appears behind a pointer. In other + /// words, if the inner type of `base` is [`Matrix`], then `index` must be a + /// constant. But if the type of `base` is a [`Pointer`] to an matrix, then + /// the index may be any expression of integer type. /// /// You can use the [`Expression::is_dynamic_index`] method to determine - /// whether a given index expression requires matrix or array base operands - /// to be behind a pointer. + /// whether a given index expression requires matrix base operands to be + /// behind a pointer. /// /// (It would be simpler to always require the use of `AccessIndex` when - /// subscripting arrays and matrices that are not behind pointers, but to - /// accommodate existing front ends, Naga also permits `Access`, with a - /// restricted `index`.) + /// subscripting matrices that are not behind pointers, but to accommodate + /// existing front ends, Naga also permits `Access`, with a restricted + /// `index`.) /// /// [`Vector`]: TypeInner::Vector /// [`Matrix`]: TypeInner::Matrix diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index a5b3ea4e38..abbe0c7e46 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -521,12 +521,12 @@ impl crate::Expression { } } - /// Return true if this expression is a dynamic array index, for [`Access`]. + /// Return true if this expression is a dynamic array/vector/matrix index, + /// for [`Access`]. /// /// This method returns true if this expression is a dynamically computed - /// index, and as such can only be used to index matrices and arrays when - /// they appear behind a pointer. See the documentation for [`Access`] for - /// details. + /// index, and as such can only be used to index matrices when they appear + /// behind a pointer. See the documentation for [`Access`] for details. /// /// Note, this does not check the _type_ of the given expression. It's up to /// the caller to establish that the `Access` expression is well-typed diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs index d8af0cd236..f29ff40f02 100644 --- a/naga/src/proc/typifier.rs +++ b/naga/src/proc/typifier.rs @@ -92,6 +92,13 @@ pub enum TypeResolution { /// available in the associated arena. However, the `TypeInner` itself may /// contain `Handle` values referring to types from the arena. /// + /// The inner type must only be one of the following variants: + /// - TypeInner::Pointer + /// - TypeInner::ValuePointer + /// - TypeInner::Matrix (generated by matrix multiplication) + /// - TypeInner::Vector + /// - TypeInner::Scalar + /// /// [`TypeInner`]: crate::TypeInner Value(crate::TypeInner), } diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index 0b0d115c57..2b479d3a73 100644 --- a/naga/src/valid/expression.rs +++ b/naga/src/valid/expression.rs @@ -240,9 +240,10 @@ impl super::Validator { let base_type = &resolver[base]; // See the documentation for `Expression::Access`. let dynamic_indexing_restricted = match *base_type { - Ti::Vector { .. } => false, - Ti::Matrix { .. } | Ti::Array { .. } => true, - Ti::Pointer { .. } + Ti::Matrix { .. } => true, + Ti::Vector { .. } + | Ti::Array { .. } + | Ti::Pointer { .. } | Ti::ValuePointer { size: Some(_), .. } | Ti::BindingArray { .. } => false, ref other => { diff --git a/naga/tests/in/access.wgsl b/naga/tests/in/access.wgsl index 956a694aaa..3336522fd9 100644 --- a/naga/tests/in/access.wgsl +++ b/naga/tests/in/access.wgsl @@ -167,3 +167,14 @@ fn assign_through_ptr() { var arr = array, 2>(vec4(6.0), vec4(7.0)); assign_array_through_ptr_fn(&arr); } + +@vertex +fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { + let arr = array(1, 2, 3, 4, 5); + let value = arr[vi]; + return vec4(vec4(value)); +} + +fn array_by_value(a: array, i: i32) -> i32 { + return a[i]; +} diff --git a/naga/tests/out/analysis/access.info.ron b/naga/tests/out/analysis/access.info.ron index 308bb1a8b6..830831cb1f 100644 --- a/naga/tests/out/analysis/access.info.ron +++ b/naga/tests/out/analysis/access.info.ron @@ -2735,6 +2735,54 @@ sampling: [], dual_source_blending: false, ), + ( + flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE"), + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + (""), + (""), + (""), + (""), + (""), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(25), + ), + ( + uniformity: ( + non_uniform_result: Some(1), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(2), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(2), + ), + ], + sampling: [], + dual_source_blending: false, + ), ], entry_points: [ ( @@ -3981,6 +4029,144 @@ sampling: [], dual_source_blending: false, ), + ( + flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE"), + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + (""), + (""), + (""), + (""), + (""), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(0), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(25), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(2), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + scalar: ( + kind: Sint, + width: 4, + ), + )), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + scalar: ( + kind: Float, + width: 4, + ), + )), + ), + ], + sampling: [], + dual_source_blending: false, + ), ], const_expression_types: [ Value(Scalar(( diff --git a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl index 2e51bbde63..4a4791c994 100644 --- a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl +++ b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl @@ -20,8 +20,8 @@ struct MatCx2InArray { mat4x2 am[2]; }; -float read_from_private(inout float foo_1) { - float _e1 = foo_1; +float read_from_private(inout float foo_2) { + float _e1 = foo_2; return _e1; } @@ -34,11 +34,15 @@ void assign_through_ptr_fn(inout uint p) { return; } -void assign_array_through_ptr_fn(inout vec4 foo_2[2]) { - foo_2 = vec4[2](vec4(1.0), vec4(2.0)); +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); return; } +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + void main() { uint val = 33u; vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0)); diff --git a/naga/tests/out/glsl/access.foo.Vertex.glsl b/naga/tests/out/glsl/access.foo.Vertex.glsl new file mode 100644 index 0000000000..e1f3138403 --- /dev/null +++ b/naga/tests/out/glsl/access.foo.Vertex.glsl @@ -0,0 +1,52 @@ +#version 310 es + +precision highp float; +precision highp int; + +struct GlobalConst { + uint a; + uvec3 b; + int c; +}; +struct AlignedWrapper { + int value; +}; +struct Baz { + mat3x2 m; +}; +struct MatCx2InArray { + mat4x2 am[2]; +}; + +float read_from_private(inout float foo_2) { + float _e1 = foo_2; + return _e1; +} + +float test_arr_as_arg(float a[5][10]) { + return a[4][9]; +} + +void assign_through_ptr_fn(inout uint p) { + p = 42u; + return; +} + +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); + return; +} + +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + +void main() { + uint vi_1 = uint(gl_VertexID); + int arr_1[5] = int[5](1, 2, 3, 4, 5); + int value = arr_1[vi_1]; + gl_Position = vec4(ivec4(value)); + gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w); + return; +} + diff --git a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl index aacdda0130..eca6bc54c5 100644 --- a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -30,8 +30,8 @@ layout(std430) buffer type_13_block_1Fragment { ivec2 _group_0_binding_2_fs; }; layout(location = 0) out vec4 _fs2p_location0; -float read_from_private(inout float foo_1) { - float _e1 = foo_1; +float read_from_private(inout float foo_2) { + float _e1 = foo_2; return _e1; } @@ -44,11 +44,15 @@ void assign_through_ptr_fn(inout uint p) { return; } -void assign_array_through_ptr_fn(inout vec4 foo_2[2]) { - foo_2 = vec4[2](vec4(1.0), vec4(2.0)); +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); return; } +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + void main() { _group_0_binding_0_fs._matrix[1][2] = 1.0; _group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0)); diff --git a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl index d4a9b92945..a926eadf78 100644 --- a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -103,8 +103,8 @@ void test_matrix_within_array_within_struct_accesses() { return; } -float read_from_private(inout float foo_1) { - float _e1 = foo_1; +float read_from_private(inout float foo_2) { + float _e1 = foo_2; return _e1; } @@ -117,11 +117,15 @@ void assign_through_ptr_fn(inout uint p) { return; } -void assign_array_through_ptr_fn(inout vec4 foo_2[2]) { - foo_2 = vec4[2](vec4(1.0), vec4(2.0)); +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); return; } +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + void main() { uint vi = uint(gl_VertexID); float foo = 0.0; @@ -133,10 +137,10 @@ void main() { mat4x3 _matrix = _group_0_binding_0_vs._matrix; uvec2 arr_1[2] = _group_0_binding_0_vs.arr; float b = _group_0_binding_0_vs._matrix[3u][0]; - int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; + int a_2 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; ivec2 c = _group_0_binding_2_vs; float _e33 = read_from_private(foo); - c2_ = int[5](a_1, int(b), 3, 4, 5); + c2_ = int[5](a_2, int(b), 3, 4, 5); c2_[(vi + 1u)] = 42; int value = c2_[vi]; float _e47 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0))); diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl index 142083be68..543b841967 100644 --- a/naga/tests/out/hlsl/access.hlsl +++ b/naga/tests/out/hlsl/access.hlsl @@ -201,9 +201,9 @@ void test_matrix_within_array_within_struct_accesses() return; } -float read_from_private(inout float foo_1) +float read_from_private(inout float foo_2) { - float _e1 = foo_1; + float _e1 = foo_2; return _e1; } @@ -224,12 +224,17 @@ ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) { return ret; } -void assign_array_through_ptr_fn(inout float4 foo_2[2]) +void assign_array_through_ptr_fn(inout float4 foo_3[2]) { - foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx); + foo_3 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx); return; } +int array_by_value(int a_1[5], int i) +{ + return a_1[i]; +} + typedef int ret_Constructarray5_int_[5]; ret_Constructarray5_int_ Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4) { int ret[5] = { arg0, arg1, arg2, arg3, arg4 }; @@ -266,10 +271,10 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48))); uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))); float b = asfloat(bar.Load(0+3u*16+0)); - int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160)); + int a_2 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160)); int2 c = asint(qux.Load2(0)); const float _e33 = read_from_private(foo); - c2_ = Constructarray5_int_(a_1, int(b), 3, 4, 5); + c2_ = Constructarray5_int_(a_2, int(b), 3, 4, 5); c2_[(vi + 1u)] = 42; int value = c2_[vi]; const float _e47 = test_arr_as_arg(ZeroValuearray5_array10_float__()); @@ -310,3 +315,10 @@ void assign_through_ptr() assign_array_through_ptr_fn(arr); return; } + +float4 foo_1(uint vi_1 : SV_VertexID) : SV_Position +{ + int arr_2[5] = Constructarray5_int_(1, 2, 3, 4, 5); + int value_1 = arr_2[vi_1]; + return float4((value_1).xxxx); +} diff --git a/naga/tests/out/hlsl/access.ron b/naga/tests/out/hlsl/access.ron index 73c9e44448..8960a612ed 100644 --- a/naga/tests/out/hlsl/access.ron +++ b/naga/tests/out/hlsl/access.ron @@ -4,6 +4,10 @@ entry_point:"foo_vert", target_profile:"vs_5_1", ), + ( + entry_point:"foo_1", + target_profile:"vs_5_1", + ), ], fragment:[ ( diff --git a/naga/tests/out/ir/access.compact.ron b/naga/tests/out/ir/access.compact.ron index 1b95742ff2..2d066b8ffa 100644 --- a/naga/tests/out/ir/access.compact.ron +++ b/naga/tests/out/ir/access.compact.ron @@ -1655,6 +1655,47 @@ ), ], ), + ( + name: Some("array_by_value"), + arguments: [ + ( + name: Some("a"), + ty: 25, + binding: None, + ), + ( + name: Some("i"), + ty: 2, + binding: None, + ), + ], + result: Some(( + ty: 2, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + FunctionArgument(1), + Access( + base: 0, + index: 1, + ), + ], + named_expressions: { + 0: "a", + 1: "i", + }, + body: [ + Emit(( + start: 2, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), ], entry_points: [ ( @@ -2230,5 +2271,81 @@ ], ), ), + ( + name: "foo", + stage: Vertex, + early_depth_test: None, + workgroup_size: (0, 0, 0), + function: ( + name: Some("foo"), + arguments: [ + ( + name: Some("vi"), + ty: 0, + binding: Some(BuiltIn(VertexIndex)), + ), + ], + result: Some(( + ty: 24, + binding: Some(BuiltIn(Position( + invariant: false, + ))), + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + Literal(I32(1)), + Literal(I32(2)), + Literal(I32(3)), + Literal(I32(4)), + Literal(I32(5)), + Compose( + ty: 25, + components: [ + 1, + 2, + 3, + 4, + 5, + ], + ), + Access( + base: 6, + index: 0, + ), + Splat( + size: Quad, + value: 7, + ), + As( + expr: 8, + kind: Float, + convert: Some(4), + ), + ], + named_expressions: { + 0: "vi", + 6: "arr", + 7: "value", + }, + body: [ + Emit(( + start: 6, + end: 7, + )), + Emit(( + start: 7, + end: 8, + )), + Emit(( + start: 8, + end: 10, + )), + Return( + value: Some(9), + ), + ], + ), + ), ], ) \ No newline at end of file diff --git a/naga/tests/out/ir/access.ron b/naga/tests/out/ir/access.ron index 1b95742ff2..2d066b8ffa 100644 --- a/naga/tests/out/ir/access.ron +++ b/naga/tests/out/ir/access.ron @@ -1655,6 +1655,47 @@ ), ], ), + ( + name: Some("array_by_value"), + arguments: [ + ( + name: Some("a"), + ty: 25, + binding: None, + ), + ( + name: Some("i"), + ty: 2, + binding: None, + ), + ], + result: Some(( + ty: 2, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + FunctionArgument(1), + Access( + base: 0, + index: 1, + ), + ], + named_expressions: { + 0: "a", + 1: "i", + }, + body: [ + Emit(( + start: 2, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), ], entry_points: [ ( @@ -2230,5 +2271,81 @@ ], ), ), + ( + name: "foo", + stage: Vertex, + early_depth_test: None, + workgroup_size: (0, 0, 0), + function: ( + name: Some("foo"), + arguments: [ + ( + name: Some("vi"), + ty: 0, + binding: Some(BuiltIn(VertexIndex)), + ), + ], + result: Some(( + ty: 24, + binding: Some(BuiltIn(Position( + invariant: false, + ))), + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + Literal(I32(1)), + Literal(I32(2)), + Literal(I32(3)), + Literal(I32(4)), + Literal(I32(5)), + Compose( + ty: 25, + components: [ + 1, + 2, + 3, + 4, + 5, + ], + ), + Access( + base: 6, + index: 0, + ), + Splat( + size: Quad, + value: 7, + ), + As( + expr: 8, + kind: Float, + convert: Some(4), + ), + ], + named_expressions: { + 0: "vi", + 6: "arr", + 7: "value", + }, + body: [ + Emit(( + start: 6, + end: 7, + )), + Emit(( + start: 7, + end: 8, + )), + Emit(( + start: 8, + end: 10, + )), + Return( + value: Some(9), + ), + ], + ), + ), ], ) \ No newline at end of file diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl index 65dba4910e..924b604e4c 100644 --- a/naga/tests/out/msl/access.msl +++ b/naga/tests/out/msl/access.msl @@ -133,9 +133,9 @@ void test_matrix_within_array_within_struct_accesses( } float read_from_private( - thread float& foo_1 + thread float& foo_2 ) { - float _e1 = foo_1; + float _e1 = foo_2; return _e1; } @@ -153,12 +153,19 @@ void assign_through_ptr_fn( } void assign_array_through_ptr_fn( - thread type_22& foo_2 + thread type_22& foo_3 ) { - foo_2 = type_22 {metal::float4(1.0), metal::float4(2.0)}; + foo_3 = type_22 {metal::float4(1.0), metal::float4(2.0)}; return; } +int array_by_value( + type_20 a_1, + int i +) { + return a_1.inner[i]; +} + struct foo_vertInput { }; struct foo_vertOutput { @@ -181,10 +188,10 @@ vertex foo_vertOutput foo_vert( metal::float4x3 _matrix = bar._matrix; type_10 arr_1 = bar.arr; float b = bar._matrix[3u].x; - int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; + int a_2 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; metal::int2 c = qux; float _e33 = read_from_private(foo); - c2_ = type_20 {a_1, static_cast(b), 3, 4, 5}; + c2_ = type_20 {a_2, static_cast(b), 3, 4, 5}; c2_.inner[vi + 1u] = 42; int value = c2_.inner[vi]; float _e47 = test_arr_as_arg(type_18 {}); @@ -217,3 +224,17 @@ kernel void assign_through_ptr( assign_array_through_ptr_fn(arr); return; } + + +struct foo_1Input { +}; +struct foo_1Output { + metal::float4 member_3 [[position]]; +}; +vertex foo_1Output foo_1( + uint vi_1 [[vertex_id]] +) { + type_20 arr_2 = type_20 {1, 2, 3, 4, 5}; + int value_1 = arr_2.inner[vi_1]; + return foo_1Output { static_cast(metal::int4(value_1)) }; +} diff --git a/naga/tests/out/spv/access.spvasm b/naga/tests/out/spv/access.spvasm index ab0112870f..6cb87a3b10 100644 --- a/naga/tests/out/spv/access.spvasm +++ b/naga/tests/out/spv/access.spvasm @@ -1,16 +1,17 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 301 +; Bound: 323 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %219 "foo_vert" %214 %217 -OpEntryPoint Fragment %273 "foo_frag" %272 -OpEntryPoint GLCompute %291 "assign_through_ptr" -OpExecutionMode %273 OriginUpperLeft -OpExecutionMode %291 LocalSize 1 1 1 +OpEntryPoint Vertex %229 "foo_vert" %224 %227 +OpEntryPoint Fragment %282 "foo_frag" %281 +OpEntryPoint GLCompute %300 "assign_through_ptr" +OpEntryPoint Vertex %314 "foo" %311 %313 +OpExecutionMode %282 OriginUpperLeft +OpExecutionMode %300 LocalSize 1 1 1 OpMemberName %6 0 "a" OpMemberName %6 1 "b" OpMemberName %6 2 "c" @@ -47,14 +48,19 @@ OpName %200 "p" OpName %201 "assign_through_ptr_fn" OpName %206 "foo" OpName %207 "assign_array_through_ptr_fn" -OpName %214 "vi" -OpName %219 "foo_vert" -OpName %231 "foo" -OpName %232 "c2" -OpName %273 "foo_frag" -OpName %291 "assign_through_ptr" -OpName %296 "val" -OpName %297 "arr" +OpName %214 "a" +OpName %215 "i" +OpName %216 "array_by_value" +OpName %224 "vi" +OpName %229 "foo_vert" +OpName %241 "foo" +OpName %242 "c2" +OpName %282 "foo_frag" +OpName %300 "assign_through_ptr" +OpName %305 "val" +OpName %306 "arr" +OpName %311 "vi" +OpName %314 "foo" OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %6 1 Offset 16 OpMemberDecorate %6 2 Offset 28 @@ -99,9 +105,11 @@ OpDecorate %50 DescriptorSet 0 OpDecorate %50 Binding 3 OpDecorate %51 Block OpMemberDecorate %51 0 Offset 0 -OpDecorate %214 BuiltIn VertexIndex -OpDecorate %217 BuiltIn Position -OpDecorate %272 Location 0 +OpDecorate %224 BuiltIn VertexIndex +OpDecorate %227 BuiltIn Position +OpDecorate %281 Location 0 +OpDecorate %311 BuiltIn VertexIndex +OpDecorate %313 BuiltIn Position %2 = OpTypeVoid %3 = OpTypeInt 32 0 %4 = OpTypeVector %3 3 @@ -209,44 +217,49 @@ OpDecorate %272 Location 0 %209 = OpConstantComposite %31 %59 %59 %59 %59 %210 = OpConstantComposite %31 %61 %61 %61 %61 %211 = OpConstantComposite %34 %209 %210 -%215 = OpTypePointer Input %3 -%214 = OpVariable %215 Input -%218 = OpTypePointer Output %31 -%217 = OpVariable %218 Output -%221 = OpTypePointer StorageBuffer %23 -%224 = OpConstant %8 0.0 -%225 = OpConstant %3 3 -%226 = OpConstant %5 3 -%227 = OpConstant %5 4 -%228 = OpConstant %5 5 -%229 = OpConstant %5 42 -%230 = OpConstantNull %29 -%233 = OpTypePointer Function %32 -%234 = OpConstantNull %32 -%239 = OpTypePointer StorageBuffer %9 -%242 = OpTypePointer StorageBuffer %18 -%243 = OpConstant %3 4 -%246 = OpTypePointer StorageBuffer %10 -%247 = OpTypePointer StorageBuffer %8 -%250 = OpTypePointer StorageBuffer %19 -%253 = OpTypePointer StorageBuffer %7 -%254 = OpTypePointer StorageBuffer %5 -%266 = OpTypeVector %5 4 -%272 = OpVariable %218 Output -%275 = OpConstantComposite %10 %224 %224 %224 -%276 = OpConstantComposite %10 %59 %59 %59 -%277 = OpConstantComposite %10 %61 %61 %61 -%278 = OpConstantComposite %10 %63 %63 %63 -%279 = OpConstantComposite %9 %275 %276 %277 %278 -%280 = OpConstantComposite %17 %36 %36 -%281 = OpConstantComposite %17 %99 %99 -%282 = OpConstantComposite %18 %280 %281 -%283 = OpConstantNull %23 -%284 = OpConstantComposite %31 %224 %224 %224 %224 -%292 = OpConstant %3 33 -%293 = OpConstantComposite %31 %67 %67 %67 %67 -%294 = OpConstantComposite %31 %137 %137 %137 %137 -%295 = OpConstantComposite %34 %293 %294 +%217 = OpTypeFunction %5 %32 %5 +%219 = OpTypePointer Function %32 +%225 = OpTypePointer Input %3 +%224 = OpVariable %225 Input +%228 = OpTypePointer Output %31 +%227 = OpVariable %228 Output +%231 = OpTypePointer StorageBuffer %23 +%234 = OpConstant %8 0.0 +%235 = OpConstant %3 3 +%236 = OpConstant %5 3 +%237 = OpConstant %5 4 +%238 = OpConstant %5 5 +%239 = OpConstant %5 42 +%240 = OpConstantNull %29 +%243 = OpConstantNull %32 +%248 = OpTypePointer StorageBuffer %9 +%251 = OpTypePointer StorageBuffer %18 +%252 = OpConstant %3 4 +%255 = OpTypePointer StorageBuffer %10 +%256 = OpTypePointer StorageBuffer %8 +%259 = OpTypePointer StorageBuffer %19 +%262 = OpTypePointer StorageBuffer %7 +%263 = OpTypePointer StorageBuffer %5 +%275 = OpTypeVector %5 4 +%281 = OpVariable %228 Output +%284 = OpConstantComposite %10 %234 %234 %234 +%285 = OpConstantComposite %10 %59 %59 %59 +%286 = OpConstantComposite %10 %61 %61 %61 +%287 = OpConstantComposite %10 %63 %63 %63 +%288 = OpConstantComposite %9 %284 %285 %286 %287 +%289 = OpConstantComposite %17 %36 %36 +%290 = OpConstantComposite %17 %99 %99 +%291 = OpConstantComposite %18 %289 %290 +%292 = OpConstantNull %23 +%293 = OpConstantComposite %31 %234 %234 %234 %234 +%301 = OpConstant %3 33 +%302 = OpConstantComposite %31 %67 %67 %67 %67 +%303 = OpConstantComposite %31 %137 %137 %137 %137 +%304 = OpConstantComposite %34 %302 %303 +%311 = OpVariable %225 Input +%313 = OpVariable %228 Output +%315 = OpConstant %5 2 +%316 = OpConstantComposite %32 %58 %315 %236 %237 %238 %54 = OpFunction %2 None %55 %53 = OpLabel %82 = OpVariable %83 Function %58 @@ -390,72 +403,98 @@ OpBranch %212 OpStore %206 %211 OpReturn OpFunctionEnd -%219 = OpFunction %2 None %55 +%216 = OpFunction %5 None %217 +%214 = OpFunctionParameter %32 +%215 = OpFunctionParameter %5 %213 = OpLabel -%231 = OpVariable %27 Function %224 -%232 = OpVariable %233 Function %234 -%216 = OpLoad %3 %214 -%220 = OpAccessChain %56 %44 %36 -%222 = OpAccessChain %221 %47 %36 -%223 = OpAccessChain %131 %50 %36 -OpBranch %235 -%235 = OpLabel -%236 = OpLoad %8 %231 -OpStore %231 %59 -%237 = OpFunctionCall %2 %54 -%238 = OpFunctionCall %2 %130 -%240 = OpAccessChain %239 %42 %36 -%241 = OpLoad %9 %240 -%244 = OpAccessChain %242 %42 %243 -%245 = OpLoad %18 %244 -%248 = OpAccessChain %247 %42 %36 %225 %36 -%249 = OpLoad %8 %248 -%251 = OpArrayLength %3 %42 5 -%252 = OpISub %3 %251 %14 -%255 = OpAccessChain %254 %42 %30 %252 %36 -%256 = OpLoad %5 %255 -%257 = OpLoad %23 %222 -%258 = OpFunctionCall %8 %188 %231 -%259 = OpConvertFToS %5 %249 -%260 = OpCompositeConstruct %32 %256 %259 %226 %227 %228 -OpStore %232 %260 -%261 = OpIAdd %3 %216 %99 -%262 = OpAccessChain %83 %232 %261 -OpStore %262 %229 -%263 = OpAccessChain %83 %232 %216 -%264 = OpLoad %5 %263 -%265 = OpFunctionCall %8 %194 %230 -%267 = OpCompositeConstruct %266 %264 %264 %264 %264 -%268 = OpConvertSToF %31 %267 -%269 = OpMatrixTimesVector %10 %241 %268 -%270 = OpCompositeConstruct %31 %269 %61 -OpStore %217 %270 +%220 = OpVariable %219 Function +OpBranch %218 +%218 = OpLabel +OpStore %220 %214 +%221 = OpAccessChain %83 %220 %215 +%222 = OpLoad %5 %221 +OpReturnValue %222 +OpFunctionEnd +%229 = OpFunction %2 None %55 +%223 = OpLabel +%241 = OpVariable %27 Function %234 +%242 = OpVariable %219 Function %243 +%226 = OpLoad %3 %224 +%230 = OpAccessChain %56 %44 %36 +%232 = OpAccessChain %231 %47 %36 +%233 = OpAccessChain %131 %50 %36 +OpBranch %244 +%244 = OpLabel +%245 = OpLoad %8 %241 +OpStore %241 %59 +%246 = OpFunctionCall %2 %54 +%247 = OpFunctionCall %2 %130 +%249 = OpAccessChain %248 %42 %36 +%250 = OpLoad %9 %249 +%253 = OpAccessChain %251 %42 %252 +%254 = OpLoad %18 %253 +%257 = OpAccessChain %256 %42 %36 %235 %36 +%258 = OpLoad %8 %257 +%260 = OpArrayLength %3 %42 5 +%261 = OpISub %3 %260 %14 +%264 = OpAccessChain %263 %42 %30 %261 %36 +%265 = OpLoad %5 %264 +%266 = OpLoad %23 %232 +%267 = OpFunctionCall %8 %188 %241 +%268 = OpConvertFToS %5 %258 +%269 = OpCompositeConstruct %32 %265 %268 %236 %237 %238 +OpStore %242 %269 +%270 = OpIAdd %3 %226 %99 +%271 = OpAccessChain %83 %242 %270 +OpStore %271 %239 +%272 = OpAccessChain %83 %242 %226 +%273 = OpLoad %5 %272 +%274 = OpFunctionCall %8 %194 %240 +%276 = OpCompositeConstruct %275 %273 %273 %273 %273 +%277 = OpConvertSToF %31 %276 +%278 = OpMatrixTimesVector %10 %250 %277 +%279 = OpCompositeConstruct %31 %278 %61 +OpStore %227 %279 +OpReturn +OpFunctionEnd +%282 = OpFunction %2 None %55 +%280 = OpLabel +%283 = OpAccessChain %231 %47 %36 +OpBranch %294 +%294 = OpLabel +%295 = OpAccessChain %256 %42 %36 %99 %14 +OpStore %295 %59 +%296 = OpAccessChain %248 %42 %36 +OpStore %296 %288 +%297 = OpAccessChain %251 %42 %252 +OpStore %297 %291 +%298 = OpAccessChain %263 %42 %30 %99 %36 +OpStore %298 %58 +OpStore %283 %292 +OpStore %281 %293 OpReturn OpFunctionEnd -%273 = OpFunction %2 None %55 -%271 = OpLabel -%274 = OpAccessChain %221 %47 %36 -OpBranch %285 -%285 = OpLabel -%286 = OpAccessChain %247 %42 %36 %99 %14 -OpStore %286 %59 -%287 = OpAccessChain %239 %42 %36 -OpStore %287 %279 -%288 = OpAccessChain %242 %42 %243 -OpStore %288 %282 -%289 = OpAccessChain %254 %42 %30 %99 %36 -OpStore %289 %58 -OpStore %274 %283 -OpStore %272 %284 +%300 = OpFunction %2 None %55 +%299 = OpLabel +%305 = OpVariable %33 Function %301 +%306 = OpVariable %35 Function %304 +OpBranch %307 +%307 = OpLabel +%308 = OpFunctionCall %2 %201 %305 +%309 = OpFunctionCall %2 %207 %306 OpReturn OpFunctionEnd -%291 = OpFunction %2 None %55 -%290 = OpLabel -%296 = OpVariable %33 Function %292 -%297 = OpVariable %35 Function %295 -OpBranch %298 -%298 = OpLabel -%299 = OpFunctionCall %2 %201 %296 -%300 = OpFunctionCall %2 %207 %297 +%314 = OpFunction %2 None %55 +%310 = OpLabel +%318 = OpVariable %219 Function +%312 = OpLoad %3 %311 +OpBranch %317 +%317 = OpLabel +OpStore %318 %316 +%319 = OpAccessChain %83 %318 %312 +%320 = OpLoad %5 %319 +%321 = OpCompositeConstruct %275 %320 %320 %320 %320 +%322 = OpConvertSToF %31 %321 +OpStore %313 %322 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/access.wgsl b/naga/tests/out/wgsl/access.wgsl index 1409e80b11..0c29e5d605 100644 --- a/naga/tests/out/wgsl/access.wgsl +++ b/naga/tests/out/wgsl/access.wgsl @@ -107,8 +107,8 @@ fn test_matrix_within_array_within_struct_accesses() { return; } -fn read_from_private(foo_1: ptr) -> f32 { - let _e1 = (*foo_1); +fn read_from_private(foo_2: ptr) -> f32 { + let _e1 = (*foo_2); return _e1; } @@ -121,11 +121,15 @@ fn assign_through_ptr_fn(p: ptr) { return; } -fn assign_array_through_ptr_fn(foo_2: ptr, 2>>) { - (*foo_2) = array, 2>(vec4(1f), vec4(2f)); +fn assign_array_through_ptr_fn(foo_3: ptr, 2>>) { + (*foo_3) = array, 2>(vec4(1f), vec4(2f)); return; } +fn array_by_value(a_1: array, i: i32) -> i32 { + return a_1[i]; +} + @vertex fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32 = 0f; @@ -138,11 +142,11 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { let _matrix = bar._matrix; let arr_1 = bar.arr; let b = bar._matrix[3u][0]; - let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value; + let a_2 = bar.data[(arrayLength((&bar.data)) - 2u)].value; let c = qux; let data_pointer = (&bar.data[0].value); let _e33 = read_from_private((&foo)); - c2_ = array(a_1, i32(b), 3i, 4i, 5i); + c2_ = array(a_2, i32(b), 3i, 4i, 5i); c2_[(vi + 1u)] = 42i; let value = c2_[vi]; let _e47 = test_arr_as_arg(array, 5>()); @@ -168,3 +172,10 @@ fn assign_through_ptr() { assign_array_through_ptr_fn((&arr)); return; } + +@vertex +fn foo_1(@builtin(vertex_index) vi_1: u32) -> @builtin(position) vec4 { + const arr_2 = array(1i, 2i, 3i, 4i, 5i); + let value_1 = arr_2[vi_1]; + return vec4(vec4(value_1)); +} diff --git a/naga/tests/wgsl_errors.rs b/naga/tests/wgsl_errors.rs index 2d91ba01cf..e5fb77644d 100644 --- a/naga/tests/wgsl_errors.rs +++ b/naga/tests/wgsl_errors.rs @@ -1359,11 +1359,6 @@ fn missing_bindings2() { #[test] fn invalid_access() { check_validation! { - " - fn array_by_value(a: array, i: i32) -> i32 { - return a[i]; - } - ", " fn matrix_by_value(m: mat4x4, i: i32) -> vec4 { return m[i]; From 215f0fc8874bb8238a357ff9280fade77f62e886 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Mon, 30 Sep 2024 13:54:02 -0700 Subject: [PATCH 09/19] [naga spv-out] Internal doc fixes for `back::spv::index`. In `naga::back::spv::index`, clarify documentation for: - `BoundsCheckResult` - `write_index_comparison` - `write_bounds_check` --- naga/src/back/spv/index.rs | 39 +++++++++++++++++++++++++++++++++----- 1 file changed, 34 insertions(+), 5 deletions(-) diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs index 0295d895b2..618770a5d9 100644 --- a/naga/src/back/spv/index.rs +++ b/naga/src/back/spv/index.rs @@ -11,16 +11,31 @@ use crate::{arena::Handle, proc::BoundsCheckPolicy}; /// The results of performing a bounds check. /// -/// On success, `write_bounds_check` returns a value of this type. +/// On success, `write_bounds_check` returns a value of this type. The +/// caller can assume that the right policy has been applied, and +/// simply do what the variant says. pub(super) enum BoundsCheckResult { /// The index is statically known and in bounds, with the given value. KnownInBounds(u32), /// The given instruction computes the index to be used. + /// + /// When [`BoundsCheckPolicy::Restrict`] is in force, this is a + /// clamped version of the index the user supplied. + /// + /// When [`BoundsCheckPolicy::Unchecked`] is in force, this is + /// simply the index the user supplied. This variant indicates + /// that we couldn't prove statically that the index was in + /// bounds; otherwise we would have returned [`KnownInBounds`]. + /// + /// [`KnownInBounds`]: BoundsCheckResult::KnownInBounds Computed(Word), /// The given instruction computes a boolean condition which is true /// if the index is in bounds. + /// + /// This is returned when [`BoundsCheckPolicy::ReadZeroSkipWrite`] + /// is in force. Conditional(Word), } @@ -357,6 +372,8 @@ impl<'w> BlockContext<'w> { /// Write an index bounds comparison to `block`, if needed. /// + /// This is used to implement [`BoundsCheckPolicy::ReadZeroSkipWrite`]. + /// /// If we're able to determine statically that `index` is in bounds for /// `sequence`, return `KnownInBounds(value)`, where `value` is the actual /// value of the index. (In principle, one could know that the index is in @@ -477,11 +494,23 @@ impl<'w> BlockContext<'w> { /// Emit code for bounds checks for an array, vector, or matrix access. /// - /// This implements either `index_bounds_check_policy` or - /// `buffer_bounds_check_policy`, depending on the address space of the - /// pointer being accessed. + /// This tries to handle all the critical steps for bounds checks: + /// + /// - First, select the appropriate bounds check policy for `base`, + /// depending on its address space. + /// + /// - Next, analyze `index` to see if its value is known at + /// compile time, in which case we can decide statically whether + /// the index is in bounds. + /// + /// - If the index's value is not known at compile time, emit code to: + /// + /// - restrict its value (for [`BoundsCheckPolicy::Restrict`]), or + /// + /// - check whether it's in bounds (for + /// [`BoundsCheckPolicy::ReadZeroSkipWrite`]). /// - /// Return a `BoundsCheckResult` indicating how the index should be + /// Return a [`BoundsCheckResult`] indicating how the index should be /// consumed. See that type's documentation for details. pub(super) fn write_bounds_check( &mut self, From 3d584f99edfc57055000fdc88c3131f5bacf3ee7 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Wed, 2 Oct 2024 22:16:53 -0400 Subject: [PATCH 10/19] refactor(spv-out): linkify docs. ref. to `write_bounds_check` --- naga/src/back/spv/index.rs | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs index 618770a5d9..832f806fd3 100644 --- a/naga/src/back/spv/index.rs +++ b/naga/src/back/spv/index.rs @@ -11,9 +11,9 @@ use crate::{arena::Handle, proc::BoundsCheckPolicy}; /// The results of performing a bounds check. /// -/// On success, `write_bounds_check` returns a value of this type. The -/// caller can assume that the right policy has been applied, and -/// simply do what the variant says. +/// On success, [`write_bounds_check`](BlockContext::write_bounds_check) +/// returns a value of this type. The caller can assume that the right +/// policy has been applied, and simply do what the variant says. pub(super) enum BoundsCheckResult { /// The index is statically known and in bounds, with the given value. KnownInBounds(u32), From 78928654a29acd87b262e211b7bea124821f238e Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Tue, 1 Oct 2024 11:37:53 -0400 Subject: [PATCH 11/19] style(wgsl-in): unblock `rustfmt` in `Error::as_parse_error` --- naga/src/front/wgsl/error.rs | 186 ++++++++++++++++++++--------------- 1 file changed, 107 insertions(+), 79 deletions(-) diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs index a7986ec898..7c65d93de3 100644 --- a/naga/src/front/wgsl/error.rs +++ b/naga/src/front/wgsl/error.rs @@ -298,32 +298,42 @@ impl<'a> Error<'a> { match *self { Error::Unexpected(unexpected_span, expected) => { let expected_str = match expected { - ExpectedToken::Token(token) => { - match token { - Token::Separator(c) => format!("'{c}'"), - Token::Paren(c) => format!("'{c}'"), - Token::Attribute => "@".to_string(), - Token::Number(_) => "number".to_string(), - Token::Word(s) => s.to_string(), - Token::Operation(c) => format!("operation ('{c}')"), - Token::LogicalOperation(c) => format!("logical operation ('{c}')"), - Token::ShiftOperation(c) => format!("bitshift ('{c}{c}')"), - Token::AssignmentOperation(c) if c=='<' || c=='>' => format!("bitshift ('{c}{c}=')"), - Token::AssignmentOperation(c) => format!("operation ('{c}=')"), - Token::IncrementOperation => "increment operation".to_string(), - Token::DecrementOperation => "decrement operation".to_string(), - Token::Arrow => "->".to_string(), - Token::Unknown(c) => format!("unknown ('{c}')"), - Token::Trivia => "trivia".to_string(), - Token::End => "end".to_string(), + ExpectedToken::Token(token) => match token { + Token::Separator(c) => format!("'{c}'"), + Token::Paren(c) => format!("'{c}'"), + Token::Attribute => "@".to_string(), + Token::Number(_) => "number".to_string(), + Token::Word(s) => s.to_string(), + Token::Operation(c) => format!("operation ('{c}')"), + Token::LogicalOperation(c) => format!("logical operation ('{c}')"), + Token::ShiftOperation(c) => format!("bitshift ('{c}{c}')"), + Token::AssignmentOperation(c) if c == '<' || c == '>' => { + format!("bitshift ('{c}{c}=')") } - } + Token::AssignmentOperation(c) => format!("operation ('{c}=')"), + Token::IncrementOperation => "increment operation".to_string(), + Token::DecrementOperation => "decrement operation".to_string(), + Token::Arrow => "->".to_string(), + Token::Unknown(c) => format!("unknown ('{c}')"), + Token::Trivia => "trivia".to_string(), + Token::End => "end".to_string(), + }, ExpectedToken::Identifier => "identifier".to_string(), ExpectedToken::PrimaryExpression => "expression".to_string(), ExpectedToken::Assignment => "assignment or increment/decrement".to_string(), - ExpectedToken::SwitchItem => "switch item ('case' or 'default') or a closing curly bracket to signify the end of the switch statement ('}')".to_string(), - ExpectedToken::WorkgroupSizeSeparator => "workgroup size separator (',') or a closing parenthesis".to_string(), - ExpectedToken::GlobalItem => "global item ('struct', 'const', 'var', 'alias', ';', 'fn') or the end of the file".to_string(), + ExpectedToken::SwitchItem => concat!( + "switch item ('case' or 'default') or a closing curly bracket ", + "to signify the end of the switch statement ('}')" + ) + .to_string(), + ExpectedToken::WorkgroupSizeSeparator => { + "workgroup size separator (',') or a closing parenthesis".to_string() + } + ExpectedToken::GlobalItem => concat!( + "global item ('struct', 'const', 'var', 'alias', ';', 'fn') ", + "or the end of the file" + ) + .to_string(), ExpectedToken::Type => "type".to_string(), ExpectedToken::Variable => "variable access".to_string(), ExpectedToken::Function => "function name".to_string(), @@ -384,9 +394,11 @@ impl<'a> Error<'a> { notes: vec![], }, Error::BadIncrDecrReferenceType(span) => ParseError { - message: - "increment/decrement operation requires reference type to be one of i32 or u32" - .to_string(), + message: concat!( + "increment/decrement operation requires ", + "reference type to be one of i32 or u32" + ) + .to_string(), labels: vec![(span, "must be a reference type of i32 or u32".into())], notes: vec![], }, @@ -527,25 +539,24 @@ impl<'a> Error<'a> { labels: vec![(span, "type can't be inferred".into())], notes: vec![], }, - Error::InitializationTypeMismatch { name, ref expected, ref got } => { - ParseError { - message: format!( - "the type of `{}` is expected to be `{}`, but got `{}`", - &source[name], expected, got, - ), - labels: vec![( - name, - format!("definition of `{}`", &source[name]).into(), - )], - notes: vec![], - } - } + Error::InitializationTypeMismatch { + name, + ref expected, + ref got, + } => ParseError { + message: format!( + "the type of `{}` is expected to be `{}`, but got `{}`", + &source[name], expected, got, + ), + labels: vec![(name, format!("definition of `{}`", &source[name]).into())], + notes: vec![], + }, Error::DeclMissingTypeAndInit(name_span) => ParseError { - message: format!("declaration of `{}` needs a type specifier or initializer", &source[name_span]), - labels: vec![( - name_span, - "needs a type specifier or initializer".into(), - )], + message: format!( + "declaration of `{}` needs a type specifier or initializer", + &source[name_span] + ), + labels: vec![(name_span, "needs a type specifier or initializer".into())], notes: vec![], }, Error::MissingAttribute(name, name_span) => ParseError { @@ -725,7 +736,11 @@ impl<'a> Error<'a> { notes: vec![message.into()], }, Error::ExpectedConstExprConcreteIntegerScalar(span) => ParseError { - message: "must be a const-expression that resolves to a concrete integer scalar (u32 or i32)".to_string(), + message: concat!( + "must be a const-expression that ", + "resolves to a concrete integer scalar (u32 or i32)" + ) + .to_string(), labels: vec![(span, "must resolve to u32 or i32".into())], notes: vec![], }, @@ -754,9 +769,17 @@ impl<'a> Error<'a> { }, Error::AutoConversion(ref error) => { // destructuring ensures all fields are handled - let AutoConversionError { dest_span, ref dest_type, source_span, ref source_type } = **error; + let AutoConversionError { + dest_span, + ref dest_type, + source_span, + ref source_type, + } = **error; ParseError { - message: format!("automatic conversions cannot convert `{source_type}` to `{dest_type}`"), + message: format!( + "automatic conversions cannot convert `{}` to `{}`", + source_type, dest_type + ), labels: vec![ ( dest_span, @@ -765,72 +788,77 @@ impl<'a> Error<'a> { ( source_span, format!("this expression has type {source_type}").into(), - ) + ), ], notes: vec![], } - }, + } Error::AutoConversionLeafScalar(ref error) => { - let AutoConversionLeafScalarError { dest_span, ref dest_scalar, source_span, ref source_type } = **error; + let AutoConversionLeafScalarError { + dest_span, + ref dest_scalar, + source_span, + ref source_type, + } = **error; ParseError { - message: format!("automatic conversions cannot convert elements of `{source_type}` to `{dest_scalar}`"), + message: format!( + "automatic conversions cannot convert elements of `{}` to `{}`", + source_type, dest_scalar + ), labels: vec![ ( dest_span, - format!("a value with elements of type {dest_scalar} is required here").into(), + format!( + "a value with elements of type {} is required here", + dest_scalar + ) + .into(), ), ( source_span, format!("this expression has type {source_type}").into(), - ) + ), ], notes: vec![], } - }, + } Error::ConcretizationFailed(ref error) => { - let ConcretizationFailedError { expr_span, ref expr_type, ref scalar, ref inner } = **error; + let ConcretizationFailedError { + expr_span, + ref expr_type, + ref scalar, + ref inner, + } = **error; ParseError { message: format!("failed to convert expression to a concrete type: {inner}"), - labels: vec![ - ( - expr_span, - format!("this expression has type {expr_type}").into(), - ) - ], - notes: vec![ - format!("the expression should have been converted to have {} scalar type", scalar), - ] + labels: vec![( + expr_span, + format!("this expression has type {expr_type}").into(), + )], + notes: vec![format!( + "the expression should have been converted to have {} scalar type", + scalar + )], } - }, + } Error::ExceededLimitForNestedBraces { span, limit } => ParseError { message: "brace nesting limit reached".into(), labels: vec![(span, "limit reached at this brace".into())], - notes: vec![ - format!("nesting limit is currently set to {limit}"), - ], + notes: vec![format!("nesting limit is currently set to {limit}")], }, Error::PipelineConstantIDValue(span) => ParseError { message: "pipeline constant ID must be between 0 and 65535 inclusive".to_string(), - labels: vec![( - span, - "must be between 0 and 65535 inclusive".into(), - )], + labels: vec![(span, "must be between 0 and 65535 inclusive".into())], notes: vec![], }, Error::NotBool(span) => ParseError { message: "must be a const-expression that resolves to a bool".to_string(), - labels: vec![( - span, - "must resolve to bool".into(), - )], + labels: vec![(span, "must resolve to bool".into())], notes: vec![], }, Error::ConstAssertFailed(span) => ParseError { message: "const_assert failure".to_string(), - labels: vec![( - span, - "evaluates to false".into(), - )], + labels: vec![(span, "evaluates to false".into())], notes: vec![], }, } From 25b2b6afa16090ab6bf80b3b2e317107b3ff187a Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Mon, 30 Sep 2024 09:29:13 -0400 Subject: [PATCH 12/19] style: newline b/w `Span::{UNDEFINED, new}` --- naga/src/span.rs | 1 + 1 file changed, 1 insertion(+) diff --git a/naga/src/span.rs b/naga/src/span.rs index 82cfbe5a4b..f8a0f67fbe 100644 --- a/naga/src/span.rs +++ b/naga/src/span.rs @@ -11,6 +11,7 @@ pub struct Span { impl Span { pub const UNDEFINED: Self = Self { start: 0, end: 0 }; + /// Creates a new `Span` from a range of byte indices /// /// Note: end is exclusive, it doesn't belong to the `Span` From 9f275f7655f1c1bcf1116fdc486daf230a230c6c Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Tue, 1 Oct 2024 15:02:51 -0400 Subject: [PATCH 13/19] style: remove comments from `MathFunction` type res. that stop `rustfmt` --- naga/src/proc/typifier.rs | 281 +++++++++++++++++++------------------- 1 file changed, 144 insertions(+), 137 deletions(-) diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs index f29ff40f02..04c776365f 100644 --- a/naga/src/proc/typifier.rs +++ b/naga/src/proc/typifier.rs @@ -638,41 +638,37 @@ impl<'a> ResolveContext<'a> { use crate::MathFunction as Mf; let res_arg = past(arg)?; match fun { - // comparison - Mf::Abs | - Mf::Min | - Mf::Max | - Mf::Clamp | - Mf::Saturate | - // trigonometry - Mf::Cos | - Mf::Cosh | - Mf::Sin | - Mf::Sinh | - Mf::Tan | - Mf::Tanh | - Mf::Acos | - Mf::Asin | - Mf::Atan | - Mf::Atan2 | - Mf::Asinh | - Mf::Acosh | - Mf::Atanh | - Mf::Radians | - Mf::Degrees | - // decomposition - Mf::Ceil | - Mf::Floor | - Mf::Round | - Mf::Fract | - Mf::Trunc | - Mf::Ldexp | - // exponent - Mf::Exp | - Mf::Exp2 | - Mf::Log | - Mf::Log2 | - Mf::Pow => res_arg.clone(), + Mf::Abs + | Mf::Min + | Mf::Max + | Mf::Clamp + | Mf::Saturate + | Mf::Cos + | Mf::Cosh + | Mf::Sin + | Mf::Sinh + | Mf::Tan + | Mf::Tanh + | Mf::Acos + | Mf::Asin + | Mf::Atan + | Mf::Atan2 + | Mf::Asinh + | Mf::Acosh + | Mf::Atanh + | Mf::Radians + | Mf::Degrees + | Mf::Ceil + | Mf::Floor + | Mf::Round + | Mf::Fract + | Mf::Trunc + | Mf::Ldexp + | Mf::Exp + | Mf::Exp2 + | Mf::Log + | Mf::Log2 + | Mf::Pow => res_arg.clone(), Mf::Modf | Mf::Frexp => { let (size, width) = match res_arg.inner_with(types) { &Ti::Scalar(crate::Scalar { @@ -680,77 +676,81 @@ impl<'a> ResolveContext<'a> { width, }) => (None, width), &Ti::Vector { - scalar: crate::Scalar { - kind: crate::ScalarKind::Float, - width, - }, + scalar: + crate::Scalar { + kind: crate::ScalarKind::Float, + width, + }, size, } => (Some(size), width), - ref other => - return Err(ResolveError::IncompatibleOperands(format!("{fun:?}({other:?}, _)"))) + ref other => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({other:?}, _)" + ))) + } }; let result = self - .special_types - .predeclared_types - .get(&if fun == Mf::Modf { - crate::PredeclaredType::ModfResult { size, width } - } else { - crate::PredeclaredType::FrexpResult { size, width } - }) - .ok_or(ResolveError::MissingSpecialType)?; + .special_types + .predeclared_types + .get(&if fun == Mf::Modf { + crate::PredeclaredType::ModfResult { size, width } + } else { + crate::PredeclaredType::FrexpResult { size, width } + }) + .ok_or(ResolveError::MissingSpecialType)?; TypeResolution::Handle(*result) - }, - // geometry + } Mf::Dot => match *res_arg.inner_with(types) { - Ti::Vector { - size: _, - scalar, - } => TypeResolution::Value(Ti::Scalar(scalar)), - ref other => - return Err(ResolveError::IncompatibleOperands( - format!("{fun:?}({other:?}, _)") - )), + Ti::Vector { size: _, scalar } => TypeResolution::Value(Ti::Scalar(scalar)), + ref other => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({other:?}, _)" + ))) + } }, Mf::Outer => { - let arg1 = arg1.ok_or_else(|| ResolveError::IncompatibleOperands( - format!("{fun:?}(_, None)") - ))?; + let arg1 = arg1.ok_or_else(|| { + ResolveError::IncompatibleOperands(format!("{fun:?}(_, None)")) + })?; match (res_arg.inner_with(types), past(arg1)?.inner_with(types)) { ( - &Ti::Vector { size: columns, scalar }, - &Ti::Vector{ size: rows, .. } + &Ti::Vector { + size: columns, + scalar, + }, + &Ti::Vector { size: rows, .. }, ) => TypeResolution::Value(Ti::Matrix { columns, rows, scalar, }), - (left, right) => - return Err(ResolveError::IncompatibleOperands( - format!("{fun:?}({left:?}, {right:?})") - )), + (left, right) => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({left:?}, {right:?})" + ))) + } } - }, + } Mf::Cross => res_arg.clone(), - Mf::Distance | - Mf::Length => match *res_arg.inner_with(types) { - Ti::Scalar(scalar) | - Ti::Vector {scalar,size:_} => TypeResolution::Value(Ti::Scalar(scalar)), - ref other => return Err(ResolveError::IncompatibleOperands( - format!("{fun:?}({other:?})") - )), + Mf::Distance | Mf::Length => match *res_arg.inner_with(types) { + Ti::Scalar(scalar) | Ti::Vector { scalar, size: _ } => { + TypeResolution::Value(Ti::Scalar(scalar)) + } + ref other => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({other:?})" + ))) + } }, - Mf::Normalize | - Mf::FaceForward | - Mf::Reflect | - Mf::Refract => res_arg.clone(), + Mf::Normalize | Mf::FaceForward | Mf::Reflect | Mf::Refract => res_arg.clone(), // computational - Mf::Sign | - Mf::Fma | - Mf::Mix | - Mf::Step | - Mf::SmoothStep | - Mf::Sqrt | - Mf::InverseSqrt => res_arg.clone(), + Mf::Sign + | Mf::Fma + | Mf::Mix + | Mf::Step + | Mf::SmoothStep + | Mf::Sqrt + | Mf::InverseSqrt => res_arg.clone(), Mf::Transpose => match *res_arg.inner_with(types) { Ti::Matrix { columns, @@ -761,9 +761,11 @@ impl<'a> ResolveContext<'a> { rows: columns, scalar, }), - ref other => return Err(ResolveError::IncompatibleOperands( - format!("{fun:?}({other:?})") - )), + ref other => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({other:?})" + ))) + } }, Mf::Inverse => match *res_arg.inner_with(types) { Ti::Matrix { @@ -775,70 +777,75 @@ impl<'a> ResolveContext<'a> { rows, scalar, }), - ref other => return Err(ResolveError::IncompatibleOperands( - format!("{fun:?}({other:?})") - )), + ref other => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({other:?})" + ))) + } }, Mf::Determinant => match *res_arg.inner_with(types) { - Ti::Matrix { - scalar, - .. - } => TypeResolution::Value(Ti::Scalar(scalar)), - ref other => return Err(ResolveError::IncompatibleOperands( - format!("{fun:?}({other:?})") - )), + Ti::Matrix { scalar, .. } => TypeResolution::Value(Ti::Scalar(scalar)), + ref other => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({other:?})" + ))) + } }, // bits - Mf::CountTrailingZeros | - Mf::CountLeadingZeros | - Mf::CountOneBits | - Mf::ReverseBits | - Mf::ExtractBits | - Mf::InsertBits | - Mf::FirstTrailingBit | - Mf::FirstLeadingBit => match *res_arg.inner_with(types) { - Ti::Scalar(scalar @ crate::Scalar { - kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint, - .. - }) => TypeResolution::Value(Ti::Scalar(scalar)), - Ti::Vector { - size, - scalar: scalar @ crate::Scalar { + Mf::CountTrailingZeros + | Mf::CountLeadingZeros + | Mf::CountOneBits + | Mf::ReverseBits + | Mf::ExtractBits + | Mf::InsertBits + | Mf::FirstTrailingBit + | Mf::FirstLeadingBit => match *res_arg.inner_with(types) { + Ti::Scalar( + scalar @ crate::Scalar { kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint, .. - } + }, + ) => TypeResolution::Value(Ti::Scalar(scalar)), + Ti::Vector { + size, + scalar: + scalar @ crate::Scalar { + kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint, + .. + }, } => TypeResolution::Value(Ti::Vector { size, scalar }), - ref other => return Err(ResolveError::IncompatibleOperands( - format!("{fun:?}({other:?})") - )), + ref other => { + return Err(ResolveError::IncompatibleOperands(format!( + "{fun:?}({other:?})" + ))) + } }, // data packing - Mf::Pack4x8snorm | - Mf::Pack4x8unorm | - Mf::Pack2x16snorm | - Mf::Pack2x16unorm | - Mf::Pack2x16float | - Mf::Pack4xI8 | - Mf::Pack4xU8 => TypeResolution::Value(Ti::Scalar(crate::Scalar::U32)), + Mf::Pack4x8snorm + | Mf::Pack4x8unorm + | Mf::Pack2x16snorm + | Mf::Pack2x16unorm + | Mf::Pack2x16float + | Mf::Pack4xI8 + | Mf::Pack4xU8 => TypeResolution::Value(Ti::Scalar(crate::Scalar::U32)), // data unpacking - Mf::Unpack4x8snorm | - Mf::Unpack4x8unorm => TypeResolution::Value(Ti::Vector { + Mf::Unpack4x8snorm | Mf::Unpack4x8unorm => TypeResolution::Value(Ti::Vector { size: crate::VectorSize::Quad, - scalar: crate::Scalar::F32 - }), - Mf::Unpack2x16snorm | - Mf::Unpack2x16unorm | - Mf::Unpack2x16float => TypeResolution::Value(Ti::Vector { - size: crate::VectorSize::Bi, - scalar: crate::Scalar::F32 + scalar: crate::Scalar::F32, }), + Mf::Unpack2x16snorm | Mf::Unpack2x16unorm | Mf::Unpack2x16float => { + TypeResolution::Value(Ti::Vector { + size: crate::VectorSize::Bi, + scalar: crate::Scalar::F32, + }) + } Mf::Unpack4xI8 => TypeResolution::Value(Ti::Vector { size: crate::VectorSize::Quad, - scalar: crate::Scalar::I32 + scalar: crate::Scalar::I32, }), Mf::Unpack4xU8 => TypeResolution::Value(Ti::Vector { size: crate::VectorSize::Quad, - scalar: crate::Scalar::U32 + scalar: crate::Scalar::U32, }), } } From 6e528aaf2416eb591e727596c3752f56b5d77df1 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Tue, 1 Oct 2024 15:02:51 -0400 Subject: [PATCH 14/19] style: split (most) string literals over 100-column line width limit --- naga/src/back/pipeline_constants.rs | 5 +++- naga/src/front/glsl/context.rs | 9 +++++--- naga/src/front/glsl/functions.rs | 3 ++- naga/src/front/glsl/parser/expressions.rs | 8 ++++++- naga/src/front/glsl/variables.rs | 19 ++++++++------- naga/src/front/spv/error.rs | 8 ++++++- wgpu-hal/examples/ray-traced-triangle/main.rs | 11 +++++++-- wgpu-hal/src/auxil/renderdoc.rs | 6 +++-- wgpu-hal/src/gles/mod.rs | 23 +++++++++++++++---- wgpu-hal/src/gles/web.rs | 7 +++--- wgpu-hal/src/vulkan/instance.rs | 12 ++++++++-- wgpu-hal/src/vulkan/mod.rs | 8 ++++++- 12 files changed, 89 insertions(+), 30 deletions(-) diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index 5f82862f72..0005cbcb0e 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -14,7 +14,10 @@ use thiserror::Error; pub enum PipelineConstantError { #[error("Missing value for pipeline-overridable constant with identifier string: '{0}'")] MissingValue(String), - #[error("Source f64 value needs to be finite (NaNs and Inifinites are not allowed) for number destinations")] + #[error( + "Source f64 value needs to be finite ({}) for number destinations", + "NaNs and Inifinites are not allowed" + )] SrcNeedsToBeFinite, #[error("Source f64 value doesn't fit in destination")] DstRangeTooSmall, diff --git a/naga/src/front/glsl/context.rs b/naga/src/front/glsl/context.rs index ee1fcc04ba..b4cb1c874e 100644 --- a/naga/src/front/glsl/context.rs +++ b/naga/src/front/glsl/context.rs @@ -630,7 +630,8 @@ impl<'a> Context<'a> { frontend.errors.push(Error { kind: ErrorKind::SemanticError( format!( - "Cannot apply operation to {left_inner:?} and {right_inner:?}" + "Cannot apply operation to {:?} and {:?}", + left_inner, right_inner ) .into(), ), @@ -828,7 +829,8 @@ impl<'a> Context<'a> { frontend.errors.push(Error { kind: ErrorKind::SemanticError( format!( - "Cannot apply operation to {left_inner:?} and {right_inner:?}" + "Cannot apply operation to {:?} and {:?}", + left_inner, right_inner ) .into(), ), @@ -908,7 +910,8 @@ impl<'a> Context<'a> { frontend.errors.push(Error { kind: ErrorKind::SemanticError( format!( - "Cannot apply operation to {left_inner:?} and {right_inner:?}" + "Cannot apply operation to {:?} and {:?}", + left_inner, right_inner ) .into(), ), diff --git a/naga/src/front/glsl/functions.rs b/naga/src/front/glsl/functions.rs index a1a6038263..2a63c7a028 100644 --- a/naga/src/front/glsl/functions.rs +++ b/naga/src/front/glsl/functions.rs @@ -634,7 +634,8 @@ impl Frontend { self.errors.push(Error { kind: ErrorKind::SemanticError( format!( - "'{name}': image needs {overload_access:?} access but only {call_access:?} was provided" + "'{}': image needs {:?} access but only {:?} was provided", + name, overload_access, call_access ) .into(), ), diff --git a/naga/src/front/glsl/parser/expressions.rs b/naga/src/front/glsl/parser/expressions.rs index 594ad6a6cd..c218e7b115 100644 --- a/naga/src/front/glsl/parser/expressions.rs +++ b/naga/src/front/glsl/parser/expressions.rs @@ -38,7 +38,13 @@ impl<'source> ParsingContext<'source> { TokenValue::FloatConstant(float) => { if float.width != 32 { frontend.errors.push(Error { - kind: ErrorKind::SemanticError("Unsupported floating-point value (expected single-precision floating-point number)".into()), + kind: ErrorKind::SemanticError( + concat!( + "Unsupported floating-point value ", + "(expected single-precision floating-point number)" + ) + .into(), + ), meta: token.meta, }); } diff --git a/naga/src/front/glsl/variables.rs b/naga/src/front/glsl/variables.rs index 6b74b254bd..16c5bb65d6 100644 --- a/naga/src/front/glsl/variables.rs +++ b/naga/src/front/glsl/variables.rs @@ -294,14 +294,17 @@ impl Frontend { .any(|i| components[i..].contains(&components[i - 1])); if not_unique { self.errors.push(Error { - kind: - ErrorKind::SemanticError( - format!( - "swizzle cannot have duplicate components in left-hand-side expression for \"{name:?}\"" - ) - .into(), - ), - meta , + kind: ErrorKind::SemanticError( + format!( + concat!( + "swizzle cannot have duplicate components in ", + "left-hand-side expression for \"{:?}\"" + ), + name + ) + .into(), + ), + meta, }) } } diff --git a/naga/src/front/spv/error.rs b/naga/src/front/spv/error.rs index 219048e102..898113d446 100644 --- a/naga/src/front/spv/error.rs +++ b/naga/src/front/spv/error.rs @@ -47,7 +47,13 @@ pub enum Error { UnsupportedBinaryOperator(spirv::Word), #[error("Naga supports OpTypeRuntimeArray in the StorageBuffer storage class only")] UnsupportedRuntimeArrayStorageClass, - #[error("unsupported matrix stride {stride} for a {columns}x{rows} matrix with scalar width={width}")] + #[error( + "unsupported matrix stride {} for a {}x{} matrix with scalar width={}", + stride, + columns, + rows, + width + )] UnsupportedMatrixStride { stride: u32, columns: u8, diff --git a/wgpu-hal/examples/ray-traced-triangle/main.rs b/wgpu-hal/examples/ray-traced-triangle/main.rs index dd91843734..4eedfe7817 100644 --- a/wgpu-hal/examples/ray-traced-triangle/main.rs +++ b/wgpu-hal/examples/ray-traced-triangle/main.rs @@ -124,7 +124,12 @@ impl AccelerationStructureInstance { &mut self, shader_binding_table_record_offset: u32, ) { - debug_assert!(shader_binding_table_record_offset <= Self::MAX_U24, "shader_binding_table_record_offset uses more than 24 bits! {shader_binding_table_record_offset} > {}", Self::MAX_U24); + debug_assert!( + shader_binding_table_record_offset <= Self::MAX_U24, + "shader_binding_table_record_offset uses more than 24 bits! {} > {}", + shader_binding_table_record_offset, + Self::MAX_U24 + ); self.shader_binding_table_record_offset_and_flags = (shader_binding_table_record_offset & Self::LOW_24_MASK) | (self.shader_binding_table_record_offset_and_flags & !Self::LOW_24_MASK) @@ -151,7 +156,9 @@ impl AccelerationStructureInstance { ); debug_assert!( shader_binding_table_record_offset <= Self::MAX_U24, - "shader_binding_table_record_offset uses more than 24 bits! {shader_binding_table_record_offset} > {}", Self::MAX_U24 + "shader_binding_table_record_offset uses more than 24 bits! {} > {}", + shader_binding_table_record_offset, + Self::MAX_U24 ); AccelerationStructureInstance { transform: Self::affine_to_rows(transform), diff --git a/wgpu-hal/src/auxil/renderdoc.rs b/wgpu-hal/src/auxil/renderdoc.rs index 3b08955fad..3879bb9545 100644 --- a/wgpu-hal/src/auxil/renderdoc.rs +++ b/wgpu-hal/src/auxil/renderdoc.rs @@ -74,7 +74,8 @@ impl RenderDoc { Err(e) => { return RenderDoc::NotAvailable { reason: format!( - "Unable to get RENDERDOC_GetAPI from renderdoc library '{renderdoc_filename}': {e:?}" + "Unable to get RENDERDOC_GetAPI from renderdoc library '{}': {e:?}", + renderdoc_filename ), } } @@ -89,7 +90,8 @@ impl RenderDoc { }, return_value => RenderDoc::NotAvailable { reason: format!( - "Unable to get API from renderdoc library '{renderdoc_filename}': {return_value}" + "Unable to get API from renderdoc library '{}': {}", + renderdoc_filename, return_value ), }, } diff --git a/wgpu-hal/src/gles/mod.rs b/wgpu-hal/src/gles/mod.rs index df59778065..8eb2800895 100644 --- a/wgpu-hal/src/gles/mod.rs +++ b/wgpu-hal/src/gles/mod.rs @@ -457,11 +457,24 @@ impl Texture { }; log::error!( - "wgpu-hal heuristics assumed that the view dimension will be equal to `{got}` rather than `{view_dimension:?}`.\n{}\n{}\n{}\n{}", - "`D2` textures with `depth_or_array_layers == 1` are assumed to have view dimension `D2`", - "`D2` textures with `depth_or_array_layers > 1` are assumed to have view dimension `D2Array`", - "`D2` textures with `depth_or_array_layers == 6` are assumed to have view dimension `Cube`", - "`D2` textures with `depth_or_array_layers > 6 && depth_or_array_layers % 6 == 0` are assumed to have view dimension `CubeArray`", + concat!( + "wgpu-hal heuristics assumed that ", + "the view dimension will be equal to `{}` rather than `{:?}`.\n", + "`D2` textures with ", + "`depth_or_array_layers == 1` ", + "are assumed to have view dimension `D2`\n", + "`D2` textures with ", + "`depth_or_array_layers > 1` ", + "are assumed to have view dimension `D2Array`\n", + "`D2` textures with ", + "`depth_or_array_layers == 6` ", + "are assumed to have view dimension `Cube`\n", + "`D2` textures with ", + "`depth_or_array_layers > 6 && depth_or_array_layers % 6 == 0` ", + "are assumed to have view dimension `CubeArray`\n", + ), + got, + view_dimension, ); } } diff --git a/wgpu-hal/src/gles/web.rs b/wgpu-hal/src/gles/web.rs index ae7f836223..06bd871247 100644 --- a/wgpu-hal/src/gles/web.rs +++ b/wgpu-hal/src/gles/web.rs @@ -64,9 +64,10 @@ impl Instance { // “not supported” could include “insufficient GPU resources” or “the GPU process // previously crashed”. So, we must return it as an `Err` since it could occur // for circumstances outside the application author's control. - return Err(crate::InstanceError::new(String::from( - "canvas.getContext() returned null; webgl2 not available or canvas already in use" - ))); + return Err(crate::InstanceError::new(String::from(concat!( + "canvas.getContext() returned null; ", + "webgl2 not available or canvas already in use" + )))); } Err(js_error) => { // diff --git a/wgpu-hal/src/vulkan/instance.rs b/wgpu-hal/src/vulkan/instance.rs index 5673859e45..6d56ecf964 100644 --- a/wgpu-hal/src/vulkan/instance.rs +++ b/wgpu-hal/src/vulkan/instance.rs @@ -745,7 +745,12 @@ impl crate::Instance for super::Instance { Ok(sdk_ver) => sdk_ver, Err(err) => { log::error!( - "Couldn't parse Android's ro.build.version.sdk system property ({val}): {err}" + concat!( + "Couldn't parse Android's ", + "ro.build.version.sdk system property ({}): {}", + ), + val, + err, ); 0 } @@ -931,7 +936,10 @@ impl crate::Instance for super::Instance { if version < (21, 2) { // See https://gitlab.freedesktop.org/mesa/mesa/-/issues/4688 log::warn!( - "Disabling presentation on '{}' (id {:?}) due to NV Optimus and Intel Mesa < v21.2", + concat!( + "Disabling presentation on '{}' (id {:?}) ", + "due to NV Optimus and Intel Mesa < v21.2" + ), exposed.info.name, exposed.adapter.raw ); diff --git a/wgpu-hal/src/vulkan/mod.rs b/wgpu-hal/src/vulkan/mod.rs index 843b836f46..3b0f026fde 100644 --- a/wgpu-hal/src/vulkan/mod.rs +++ b/wgpu-hal/src/vulkan/mod.rs @@ -418,7 +418,13 @@ impl Surface { swapchain.next_present_time = Some(present_timing); } else { // Ideally we'd use something like `device.required_features` here, but that's in `wgpu-core`, which we are a dependency of - panic!("Tried to set display timing properties without the corresponding feature ({features:?}) enabled."); + panic!( + concat!( + "Tried to set display timing properties ", + "without the corresponding feature ({:?}) enabled." + ), + features + ); } } } From d3e09dd63ad54e2982ef744a6d74df468e232ff0 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Fri, 27 Sep 2024 17:10:15 -0700 Subject: [PATCH 15/19] [naga spv-out] Consolidate explanation of global wrapping. Consolidate the explanation of why and how we wrap Naga global variables in structs to satisfy Vulkan's requirements, and include it in the documentation for `back::spv::GlobalVariable`. Clarify `GlobalVariable`'s members documentation. --- naga/src/back/spv/helpers.rs | 17 +------- naga/src/back/spv/index.rs | 5 ++- naga/src/back/spv/mod.rs | 77 ++++++++++++++++++++++++++---------- 3 files changed, 62 insertions(+), 37 deletions(-) diff --git a/naga/src/back/spv/helpers.rs b/naga/src/back/spv/helpers.rs index 15c241d44e..63144abc02 100644 --- a/naga/src/back/spv/helpers.rs +++ b/naga/src/back/spv/helpers.rs @@ -85,22 +85,9 @@ impl crate::AddressSpace { /// Return true if the global requires a type decorated with `Block`. /// -/// In the Vulkan spec 1.3.296, the section [Descriptor Set Interface][dsi] says: +/// See [`back::spv::GlobalVariable`] for details. /// -/// > Variables identified with the `Uniform` storage class are used to -/// > access transparent buffer backed resources. Such variables must -/// > be: -/// > -/// > - typed as `OpTypeStruct`, or an array of this type, -/// > -/// > - identified with a `Block` or `BufferBlock` decoration, and -/// > -/// > - laid out explicitly using the `Offset`, `ArrayStride`, and -/// > `MatrixStride` decorations as specified in §15.6.4, "Offset -/// > and Stride Assignment." -/// -/// [dsi]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#interfaces-resources-descset -// See `back::spv::GlobalVariable::access_id` for details. +/// [`back::spv::GlobalVariable`]: super::GlobalVariable pub fn global_needs_wrapper(ir_module: &crate::Module, var: &crate::GlobalVariable) -> bool { match var.space { crate::AddressSpace::Uniform diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs index 832f806fd3..1b7cce9978 100644 --- a/naga/src/back/spv/index.rs +++ b/naga/src/back/spv/index.rs @@ -67,13 +67,14 @@ impl<'w> BlockContext<'w> { /// - An optional [`Access`] or [`AccessIndex`], for case 3, applied to... /// - A [`GlobalVariable`]. /// - /// The SPIR-V generated takes into account wrapped globals; see - /// [`global_needs_wrapper`]. + /// The generated SPIR-V takes into account wrapped globals; see + /// [`back::spv::GlobalVariable`] for details. /// /// [`GlobalVariable`]: crate::Expression::GlobalVariable /// [`AccessIndex`]: crate::Expression::AccessIndex /// [`Access`]: crate::Expression::Access /// [`base`]: crate::Expression::Access::base + /// [`back::spv::GlobalVariable`]: super::GlobalVariable pub(super) fn write_runtime_array_length( &mut self, array: Handle, diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 77f915abb1..764d06c907 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -467,38 +467,75 @@ enum CachedConstant { ZeroValue(Word), } +/// The SPIR-V representation of a [`crate::GlobalVariable`]. +/// +/// In the Vulkan spec 1.3.296, the section [Descriptor Set Interface][dsi] says: +/// +/// > Variables identified with the `Uniform` storage class are used to access +/// > transparent buffer backed resources. Such variables *must* be: +/// > +/// > - typed as `OpTypeStruct`, or an array of this type, +/// > +/// > - identified with a `Block` or `BufferBlock` decoration, and +/// > +/// > - laid out explicitly using the `Offset`, `ArrayStride`, and `MatrixStride` +/// > decorations as specified in "Offset and Stride Assignment". +/// +/// This is followed by identical language for the `StorageBuffer`, +/// except that a `BufferBlock` decoration is not allowed. +/// +/// When we encounter a global variable in the [`Storage`] or [`Uniform`] +/// address spaces whose type is not already [`Struct`], this backend implicitly +/// wraps the global variable in a struct: we generate a SPIR-V global variable +/// holding an `OpTypeStruct` with a single member, whose type is what the Naga +/// global's type would suggest, decorated as required above. +/// +/// The [`helpers::global_needs_wrapper`] function determines whether a given +/// [`crate::GlobalVariable`] needs to be wrapped. +/// +/// [dsi]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#interfaces-resources-descset +/// [`Storage`]: crate::AddressSpace::Storage +/// [`Uniform`]: crate::AddressSpace::Uniform +/// [`Struct`]: crate::TypeInner::Struct #[derive(Clone)] struct GlobalVariable { - /// ID of the OpVariable that declares the global. + /// The SPIR-V id of the `OpVariable` that declares the global. + /// + /// If this global has been implicitly wrapped in an `OpTypeStruct`, this id + /// refers to the wrapper, not the original Naga value it contains. If you + /// need the Naga value, use [`access_id`] instead of this field. /// - /// If you need the variable's value, use [`access_id`] instead of this - /// field. If we wrapped the Naga IR `GlobalVariable`'s type in a struct to - /// comply with Vulkan's requirements, then this points to the `OpVariable` - /// with the synthesized struct type, whereas `access_id` points to the - /// field of said struct that holds the variable's actual value. + /// If this global is not implicitly wrapped, this is the same as + /// [`access_id`]. /// /// This is used to compute the `access_id` pointer in function prologues, - /// and used for `ArrayLength` expressions, which do need the struct. + /// and used for `ArrayLength` expressions, which need to pass the wrapper + /// struct. /// /// [`access_id`]: GlobalVariable::access_id var_id: Word, - /// For `AddressSpace::Handle` variables, this ID is recorded in the function - /// prelude block (and reset before every function) as `OpLoad` of the variable. - /// It is then used for all the global ops, such as `OpImageSample`. + /// The loaded value of a `AddressSpace::Handle` global variable. + /// + /// If the current function uses this global variable, this is the id of an + /// `OpLoad` instruction in the function's prologue that loads its value. + /// (This value is assigned as we write the prologue code of each function.) + /// It is then used for all operations on the global, such as `OpImageSample`. handle_id: Word, - /// Actual ID used to access this variable. - /// For wrapped buffer variables, this ID is `OpAccessChain` into the - /// wrapper. Otherwise, the same as `var_id`. + /// The SPIR-V id of a pointer to this variable's Naga IR value. + /// + /// If the current function uses this global variable, and it has been + /// implicitly wrapped in an `OpTypeStruct`, this is the id of an + /// `OpAccessChain` instruction in the function's prologue that refers to + /// the wrapped value inside the struct. (This value is assigned as we write + /// the prologue code of each function.) If you need the wrapper struct + /// itself, use [`var_id`] instead of this field. + /// + /// If this global is not implicitly wrapped, this is the same as + /// [`var_id`]. /// - /// Vulkan requires that globals in the `StorageBuffer` and `Uniform` storage - /// classes must be structs with the `Block` decoration, but WGSL and Naga IR - /// make no such requirement. So for such variables, we generate a wrapper struct - /// type with a single element of the type given by Naga, generate an - /// `OpAccessChain` for that member in the function prelude, and use that pointer - /// to refer to the global in the function body. This is the id of that access, - /// updated for each function in `write_function`. + /// [`var_id`]: GlobalVariable::var_id access_id: Word, } From c5a4b4ecc3729a641a96ada0ed8275d08c16de33 Mon Sep 17 00:00:00 2001 From: Imbris <2002109+Imberflur@users.noreply.github.com> Date: Thu, 3 Oct 2024 16:40:43 -0400 Subject: [PATCH 16/19] Fix unaligned slice::from_raw_parts in gles push contant handling. (#6341) --- CHANGELOG.md | 1 + Cargo.lock | 1 + Cargo.toml | 2 +- examples/Cargo.toml | 2 +- wgpu-hal/Cargo.toml | 2 ++ wgpu-hal/src/gles/queue.rs | 58 +++++++++++++++++--------------------- 6 files changed, 32 insertions(+), 34 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 9d564bb082..fc094bcfb0 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -126,6 +126,7 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216). #### GLES / OpenGL - Fix GL debug message callbacks not being properly cleaned up (causing UB). By @Imberflur in [#6114](https://github.com/gfx-rs/wgpu/pull/6114) +- Fix calling `slice::from_raw_parts` with unaligned pointers in push constant handling. By @Imberflur in [#6341](https://github.com/gfx-rs/wgpu/pull/6341) #### WebGPU diff --git a/Cargo.lock b/Cargo.lock index 3c8d1aacce..74910d9290 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -3681,6 +3681,7 @@ dependencies = [ "bit-set", "bitflags 2.6.0", "block", + "bytemuck", "cfg-if", "cfg_aliases", "core-graphics-types", diff --git a/Cargo.toml b/Cargo.toml index fbf06524c1..68c29b671b 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -76,7 +76,7 @@ arrayvec = "0.7" bincode = "1" bit-vec = "0.8" bitflags = "2.6" -bytemuck = { version = "1.18", features = ["derive"] } +bytemuck = { version = "1.18" } cfg_aliases = "0.1" cfg-if = "1" criterion = "0.5" diff --git a/examples/Cargo.toml b/examples/Cargo.toml index f5beed1d22..1f4d4951f5 100644 --- a/examples/Cargo.toml +++ b/examples/Cargo.toml @@ -29,7 +29,7 @@ webgl = ["wgpu/webgl"] webgpu = ["wgpu/webgpu"] [dependencies] -bytemuck.workspace = true +bytemuck = { workspace = true, features = ["derive"] } cfg-if.workspace = true encase = { workspace = true, features = ["glam"] } flume.workspace = true diff --git a/wgpu-hal/Cargo.toml b/wgpu-hal/Cargo.toml index d3a05f0790..03af10b96e 100644 --- a/wgpu-hal/Cargo.toml +++ b/wgpu-hal/Cargo.toml @@ -56,6 +56,7 @@ vulkan = [ ] gles = [ "naga/glsl-out", + "dep:bytemuck", "dep:glow", "dep:glutin_wgl_sys", "dep:khronos-egl", @@ -126,6 +127,7 @@ rustc-hash.workspace = true log.workspace = true # backend: Gles +bytemuck = { workspace = true, optional = true } glow = { workspace = true, optional = true } [dependencies.wgt] diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs index 39315f72b7..fc106eb23d 100644 --- a/wgpu-hal/src/gles/queue.rs +++ b/wgpu-hal/src/gles/queue.rs @@ -1603,19 +1603,13 @@ impl super::Queue { ref uniform, offset, } => { - // T must be POD - // - // This function is absolutely sketchy and we really should be using bytemuck. - unsafe fn get_data(data: &[u8], offset: u32) -> &[T; COUNT] { + fn get_data(data: &[u8], offset: u32) -> [T; COUNT] + where + [T; COUNT]: bytemuck::AnyBitPattern, + { let data_required = size_of::() * COUNT; - let raw = &data[(offset as usize)..][..data_required]; - - debug_assert_eq!(data_required, raw.len()); - - let slice: &[T] = unsafe { slice::from_raw_parts(raw.as_ptr().cast(), COUNT) }; - - slice.try_into().unwrap() + bytemuck::pod_read_unaligned(raw) } let location = Some(&uniform.location); @@ -1625,28 +1619,28 @@ impl super::Queue { // --- Float 1-4 Component --- // naga::TypeInner::Scalar(naga::Scalar::F32) => { - let data = unsafe { get_data::(data_bytes, offset)[0] }; + let data = get_data::(data_bytes, offset)[0]; unsafe { gl.uniform_1_f32(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Bi, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_2_f32_slice(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Tri, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_3_f32_slice(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Quad, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_4_f32_slice(location, data) }; } @@ -1654,28 +1648,28 @@ impl super::Queue { // --- Int 1-4 Component --- // naga::TypeInner::Scalar(naga::Scalar::I32) => { - let data = unsafe { get_data::(data_bytes, offset)[0] }; + let data = get_data::(data_bytes, offset)[0]; unsafe { gl.uniform_1_i32(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Bi, scalar: naga::Scalar::I32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_2_i32_slice(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Tri, scalar: naga::Scalar::I32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_3_i32_slice(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Quad, scalar: naga::Scalar::I32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_4_i32_slice(location, data) }; } @@ -1683,28 +1677,28 @@ impl super::Queue { // --- Uint 1-4 Component --- // naga::TypeInner::Scalar(naga::Scalar::U32) => { - let data = unsafe { get_data::(data_bytes, offset)[0] }; + let data = get_data::(data_bytes, offset)[0]; unsafe { gl.uniform_1_u32(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Bi, scalar: naga::Scalar::U32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_2_u32_slice(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Tri, scalar: naga::Scalar::U32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_3_u32_slice(location, data) }; } naga::TypeInner::Vector { size: naga::VectorSize::Quad, scalar: naga::Scalar::U32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_4_u32_slice(location, data) }; } @@ -1716,7 +1710,7 @@ impl super::Queue { rows: naga::VectorSize::Bi, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_matrix_2_f32_slice(location, false, data) }; } naga::TypeInner::Matrix { @@ -1725,7 +1719,7 @@ impl super::Queue { scalar: naga::Scalar::F32, } => { // repack 2 vec3s into 6 values. - let unpacked_data = unsafe { get_data::(data_bytes, offset) }; + let unpacked_data = &get_data::(data_bytes, offset); #[rustfmt::skip] let packed_data = [ unpacked_data[0], unpacked_data[1], unpacked_data[2], @@ -1738,7 +1732,7 @@ impl super::Queue { rows: naga::VectorSize::Quad, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_matrix_2x4_f32_slice(location, false, data) }; } @@ -1750,7 +1744,7 @@ impl super::Queue { rows: naga::VectorSize::Bi, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_matrix_3x2_f32_slice(location, false, data) }; } naga::TypeInner::Matrix { @@ -1759,7 +1753,7 @@ impl super::Queue { scalar: naga::Scalar::F32, } => { // repack 3 vec3s into 9 values. - let unpacked_data = unsafe { get_data::(data_bytes, offset) }; + let unpacked_data = &get_data::(data_bytes, offset); #[rustfmt::skip] let packed_data = [ unpacked_data[0], unpacked_data[1], unpacked_data[2], @@ -1773,7 +1767,7 @@ impl super::Queue { rows: naga::VectorSize::Quad, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_matrix_3x4_f32_slice(location, false, data) }; } @@ -1785,7 +1779,7 @@ impl super::Queue { rows: naga::VectorSize::Bi, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_matrix_4x2_f32_slice(location, false, data) }; } naga::TypeInner::Matrix { @@ -1794,7 +1788,7 @@ impl super::Queue { scalar: naga::Scalar::F32, } => { // repack 4 vec3s into 12 values. - let unpacked_data = unsafe { get_data::(data_bytes, offset) }; + let unpacked_data = &get_data::(data_bytes, offset); #[rustfmt::skip] let packed_data = [ unpacked_data[0], unpacked_data[1], unpacked_data[2], @@ -1809,7 +1803,7 @@ impl super::Queue { rows: naga::VectorSize::Quad, scalar: naga::Scalar::F32, } => { - let data = unsafe { get_data::(data_bytes, offset) }; + let data = &get_data::(data_bytes, offset); unsafe { gl.uniform_matrix_4_f32_slice(location, false, data) }; } _ => panic!("Unsupported uniform datatype: {:?}!", uniform.ty), From 04182c24ecd9db68d03ee591a9fecde38401fed1 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Wed, 2 Oct 2024 16:16:10 -0700 Subject: [PATCH 17/19] [naga spv-out] Make `write_block` and its types private. Make `naga::back::spv::BlockContext::write_block` private to `naga::back::spv::block`. Introduce a new `pub(super)` function, `write_function_body`, for `Writer::write_function` to call. Make `BlockExit` private to `naga::back::spv::block`. Move `LoopContext` from `naga::back::spv` into `naga::back::spv::block`, and make it private. --- naga/src/back/spv/block.rs | 29 +++++++++++++++++++++++++---- naga/src/back/spv/mod.rs | 6 ------ naga/src/back/spv/writer.rs | 10 ++-------- 3 files changed, 27 insertions(+), 18 deletions(-) diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 97db1cb549..4076fc5a81 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -4,8 +4,7 @@ Implementations for `BlockContext` methods. use super::{ helpers, index::BoundsCheckResult, make_local, selection::Selection, Block, BlockContext, - Dimension, Error, Instruction, LocalType, LookupType, LoopContext, ResultMember, Writer, - WriterFlags, + Dimension, Error, Instruction, LocalType, LookupType, ResultMember, Writer, WriterFlags, }; use crate::{arena::Handle, proc::TypeResolution, Statement}; use spirv::Word; @@ -39,7 +38,7 @@ enum ExpressionPointer { } /// The termination statement to be added to the end of the block -pub enum BlockExit { +enum BlockExit { /// Generates an OpReturn (void return) Return, /// Generates an OpBranch to the specified block @@ -60,6 +59,12 @@ pub enum BlockExit { }, } +#[derive(Clone, Copy, Default)] +struct LoopContext { + continuing_id: Option, + break_id: Option, +} + #[derive(Debug)] pub(crate) struct DebugInfoInner<'a> { pub source_code: &'a str, @@ -2060,7 +2065,7 @@ impl<'w> BlockContext<'w> { } } - pub(super) fn write_block( + fn write_block( &mut self, label_id: Word, naga_block: &crate::Block, @@ -2721,4 +2726,20 @@ impl<'w> BlockContext<'w> { self.function.consume(block, termination); Ok(()) } + + pub(super) fn write_function_body( + &mut self, + entry_id: Word, + debug_info: Option<&DebugInfoInner>, + ) -> Result<(), Error> { + self.write_block( + entry_id, + &self.ir_function.body, + super::block::BlockExit::Return, + LoopContext::default(), + debug_info, + )?; + + Ok(()) + } } diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 764d06c907..e6397017c5 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -665,12 +665,6 @@ impl BlockContext<'_> { } } -#[derive(Clone, Copy, Default)] -struct LoopContext { - continuing_id: Option, - break_id: Option, -} - pub struct Writer { physical_layout: PhysicalLayout, logical_layout: LogicalLayout, diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 7f41caf295..14f1fc0027 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -3,7 +3,7 @@ use super::{ helpers::{contains_builtin, global_needs_wrapper, map_storage_class}, make_local, Block, BlockContext, CachedConstant, CachedExpressions, DebugInfo, EntryPointContext, Error, Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction, - LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, LoopContext, Options, + LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, Options, PhysicalLayout, PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE, }; use crate::{ @@ -756,13 +756,7 @@ impl Writer { next_id }; - context.write_block( - main_id, - &ir_function.body, - super::block::BlockExit::Return, - LoopContext::default(), - debug_info.as_ref(), - )?; + context.write_function_body(main_id, debug_info.as_ref())?; // Consume the `BlockContext`, ending its borrows and letting the // `Writer` steal back its cached expression table and temp_list. From e432980a7319482c65bf0e54cb4a0ca148407b5c Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Wed, 2 Oct 2024 17:13:07 -0700 Subject: [PATCH 18/19] [naga spv-out] Don't emit unreachable blocks that jump into loops. When generating SPIR-V, avoid generating unreachable blocks following statements like `break`, `return`, and so on that cause non-local exits. These unreachable blocks can cause SPIR-V validation to fail. Fixes #6220. --- naga/src/back/spv/block.rs | 101 +++++++++++++++--- naga/tests/in/6220-break-from-loop.param.ron | 2 + naga/tests/in/6220-break-from-loop.wgsl | 43 ++++++++ .../tests/out/spv/6220-break-from-loop.spvasm | 44 ++++++++ naga/tests/snapshots.rs | 1 + 5 files changed, 176 insertions(+), 15 deletions(-) create mode 100644 naga/tests/in/6220-break-from-loop.param.ron create mode 100644 naga/tests/in/6220-break-from-loop.wgsl create mode 100644 naga/tests/out/spv/6220-break-from-loop.spvasm diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 4076fc5a81..c6adba85b1 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -59,6 +59,30 @@ enum BlockExit { }, } +/// What code generation did with a provided [`BlockExit`] value. +/// +/// A function that accepts a [`BlockExit`] argument should return a value of +/// this type, to indicate whether the code it generated ended up using the +/// provided exit, or ignored it and did a non-local exit of some other kind +/// (say, [`Break`] or [`Continue`]). Some callers must use this information to +/// decide whether to generate the target block at all. +/// +/// [`Break`]: Statement::Break +/// [`Continue`]: Statement::Continue +#[must_use] +enum BlockExitDisposition { + /// The generated code used the provided `BlockExit` value. If it included a + /// block label, the caller should be sure to actually emit the block it + /// refers to. + Used, + + /// The generated code did not use the provided `BlockExit` value. If it + /// included a block label, the caller should not bother to actually emit + /// the block it refers to, unless it knows the block is needed for + /// something else. + Discarded, +} + #[derive(Clone, Copy, Default)] struct LoopContext { continuing_id: Option, @@ -2065,6 +2089,22 @@ impl<'w> BlockContext<'w> { } } + /// Generate one or more SPIR-V blocks for `naga_block`. + /// + /// Use `label_id` as the label for the SPIR-V entry point block. + /// + /// If control reaches the end of the SPIR-V block, terminate it according + /// to `exit`. This function's return value indicates whether it acted on + /// this parameter or not; see [`BlockExitDisposition`]. + /// + /// If the block contains [`Break`] or [`Continue`] statements, + /// `loop_context` supplies the labels of the SPIR-V blocks to jump to. If + /// either of these labels are `None`, then it should have been a Naga + /// validation error for the corresponding statement to occur in this + /// context. + /// + /// [`Break`]: Statement::Break + /// [`Continue`]: Statement::Continue fn write_block( &mut self, label_id: Word, @@ -2072,7 +2112,7 @@ impl<'w> BlockContext<'w> { exit: BlockExit, loop_context: LoopContext, debug_info: Option<&DebugInfoInner>, - ) -> Result<(), Error> { + ) -> Result { let mut block = Block::new(label_id); for (statement, span) in naga_block.span_iter() { if let (Some(debug_info), false) = ( @@ -2108,7 +2148,7 @@ impl<'w> BlockContext<'w> { self.function.consume(block, Instruction::branch(scope_id)); let merge_id = self.gen_id(); - self.write_block( + let merge_used = self.write_block( scope_id, block_statements, BlockExit::Branch { target: merge_id }, @@ -2116,7 +2156,14 @@ impl<'w> BlockContext<'w> { debug_info, )?; - block = Block::new(merge_id); + match merge_used { + BlockExitDisposition::Used => { + block = Block::new(merge_id); + } + BlockExitDisposition::Discarded => { + return Ok(BlockExitDisposition::Discarded); + } + } } Statement::If { condition, @@ -2152,7 +2199,11 @@ impl<'w> BlockContext<'w> { ); if let Some(block_id) = accept_id { - self.write_block( + // We can ignore the `BlockExitDisposition` returned here because, + // even if `merge_id` is not actually reachable, it is always + // referred to by the `OpSelectionMerge` instruction we emitted + // earlier. + let _ = self.write_block( block_id, accept, BlockExit::Branch { target: merge_id }, @@ -2161,7 +2212,11 @@ impl<'w> BlockContext<'w> { )?; } if let Some(block_id) = reject_id { - self.write_block( + // We can ignore the `BlockExitDisposition` returned here because, + // even if `merge_id` is not actually reachable, it is always + // referred to by the `OpSelectionMerge` instruction we emitted + // earlier. + let _ = self.write_block( block_id, reject, BlockExit::Branch { target: merge_id }, @@ -2239,7 +2294,15 @@ impl<'w> BlockContext<'w> { } else { merge_id }; - self.write_block( + // We can ignore the `BlockExitDisposition` returned here because + // `case_finish_id` is always referred to by either: + // + // - the `OpSwitch`, if it's the next case's label for a + // fall-through, or + // + // - the `OpSelectionMerge`, if it's the switch's overall merge + // block because there's no fall-through. + let _ = self.write_block( *label_id, &case.body, BlockExit::Branch { @@ -2285,7 +2348,10 @@ impl<'w> BlockContext<'w> { )); self.function.consume(block, Instruction::branch(body_id)); - self.write_block( + // We can ignore the `BlockExitDisposition` returned here because, + // even if `continuing_id` is not actually reachable, it is always + // referred to by the `OpLoopMerge` instruction we emitted earlier. + let _ = self.write_block( body_id, body, BlockExit::Branch { @@ -2308,7 +2374,10 @@ impl<'w> BlockContext<'w> { }, }; - self.write_block( + // We can ignore the `BlockExitDisposition` returned here because, + // even if `merge_id` is not actually reachable, it is always referred + // to by the `OpLoopMerge` instruction we emitted earlier. + let _ = self.write_block( continuing_id, continuing, exit, @@ -2324,14 +2393,14 @@ impl<'w> BlockContext<'w> { Statement::Break => { self.function .consume(block, Instruction::branch(loop_context.break_id.unwrap())); - return Ok(()); + return Ok(BlockExitDisposition::Discarded); } Statement::Continue => { self.function.consume( block, Instruction::branch(loop_context.continuing_id.unwrap()), ); - return Ok(()); + return Ok(BlockExitDisposition::Discarded); } Statement::Return { value: Some(value) } => { let value_id = self.cached[value]; @@ -2350,15 +2419,15 @@ impl<'w> BlockContext<'w> { None => Instruction::return_value(value_id), }; self.function.consume(block, instruction); - return Ok(()); + return Ok(BlockExitDisposition::Discarded); } Statement::Return { value: None } => { self.function.consume(block, Instruction::return_void()); - return Ok(()); + return Ok(BlockExitDisposition::Discarded); } Statement::Kill => { self.function.consume(block, Instruction::kill()); - return Ok(()); + return Ok(BlockExitDisposition::Discarded); } Statement::Barrier(flags) => { self.writer.write_barrier(flags, &mut block); @@ -2724,7 +2793,7 @@ impl<'w> BlockContext<'w> { }; self.function.consume(block, termination); - Ok(()) + Ok(BlockExitDisposition::Used) } pub(super) fn write_function_body( @@ -2732,7 +2801,9 @@ impl<'w> BlockContext<'w> { entry_id: Word, debug_info: Option<&DebugInfoInner>, ) -> Result<(), Error> { - self.write_block( + // We can ignore the `BlockExitDisposition` returned here because + // `BlockExit::Return` doesn't refer to a block. + let _ = self.write_block( entry_id, &self.ir_function.body, super::block::BlockExit::Return, diff --git a/naga/tests/in/6220-break-from-loop.param.ron b/naga/tests/in/6220-break-from-loop.param.ron new file mode 100644 index 0000000000..72873dd667 --- /dev/null +++ b/naga/tests/in/6220-break-from-loop.param.ron @@ -0,0 +1,2 @@ +( +) diff --git a/naga/tests/in/6220-break-from-loop.wgsl b/naga/tests/in/6220-break-from-loop.wgsl new file mode 100644 index 0000000000..424886a757 --- /dev/null +++ b/naga/tests/in/6220-break-from-loop.wgsl @@ -0,0 +1,43 @@ +// #6220: Don't generate unreachable SPIR-V blocks that branch into +// structured control flow constructs. +// +// Suppose we have Naga code like this: +// +// Block { +// ... prelude +// Block { ... nested } +// ... postlude +// } +// +// The SPIR-V back end used to always generate three separate SPIR-V +// blocks for the sections labeled "prelude", "nested", and +// "postlude", each block ending with a branch to the next, even if +// they were empty. +// +// However, the function below generates code that includes the +// following structure: +// +// Loop { +// body: Block { +// ... prelude +// Block { Break } +// ... postlude +// } +// continuing: ... +// } +// +// In this case, even though the `Break` renders the "postlude" +// unreachable, we used to generate a SPIR-V block for it anyway, +// ending with a branch to the `Loop`'s "continuing" block. However, +// SPIR-V's structured control flow rules forbid branches to a loop +// construct's continue target from outside the loop, so the SPIR-V +// module containing the unreachable block didn't pass validation. +// +// One might assume that unreachable blocks shouldn't affect +// validation, but the spec doesn't clearly agree, and this doesn't +// seem to be the way validation has been implemented. +fn break_from_loop() { + for (var i = 0; i < 4; i += 1) { + break; + } +} diff --git a/naga/tests/out/spv/6220-break-from-loop.spvasm b/naga/tests/out/spv/6220-break-from-loop.spvasm new file mode 100644 index 0000000000..9dabbc2079 --- /dev/null +++ b/naga/tests/out/spv/6220-break-from-loop.spvasm @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 26 +OpCapability Shader +OpCapability Linkage +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +%2 = OpTypeVoid +%3 = OpTypeInt 32 1 +%6 = OpTypeFunction %2 +%7 = OpConstant %3 0 +%8 = OpConstant %3 4 +%9 = OpConstant %3 1 +%11 = OpTypePointer Function %3 +%18 = OpTypeBool +%5 = OpFunction %2 None %6 +%4 = OpLabel +%10 = OpVariable %11 Function %7 +OpBranch %12 +%12 = OpLabel +OpBranch %13 +%13 = OpLabel +OpLoopMerge %14 %16 None +OpBranch %15 +%15 = OpLabel +%17 = OpLoad %3 %10 +%19 = OpSLessThan %18 %17 %8 +OpSelectionMerge %20 None +OpBranchConditional %19 %20 %21 +%21 = OpLabel +OpBranch %14 +%20 = OpLabel +OpBranch %22 +%22 = OpLabel +OpBranch %14 +%16 = OpLabel +%24 = OpLoad %3 %10 +%25 = OpIAdd %3 %24 %9 +OpStore %10 %25 +OpBranch %13 +%14 = OpLabel +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 596e4cea14..adf67f8333 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -932,6 +932,7 @@ fn convert_wgsl() { "phony_assignment", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ), + ("6220-break-from-loop", Targets::SPIRV), ]; for &(name, targets) in inputs.iter() { From ee0d1703e5f4a267ce9b87d50b824190b45b5a79 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Fri, 4 Oct 2024 09:37:30 -0700 Subject: [PATCH 19/19] Roll back once_cell to 1.19.0. Roll back `wgpu`'s dependencies on `once_cell` from 1.20.1 to 1.19.0. Version 1.20.1 of `once_cell` added a more complex conditional dependency on `portable-atomic`, which causes `cargo metadata` to incorrectly list `portable-atomic` as a dependency even though the given `once_cell` features are not enabled. The Firefox source tree uses `cargo vet` to enforce supply-chain auditing. Since `cargo vet` depends on `cargo metadata` to tell it what crates are going to be included in the tree, the extraneous dependency above adds `portable-atomic` to the set of sources we must audit. Since `portable-atomic` is roughly 50kloc, we would like to avoid this. Nothing in `wgpu` actually needs `once_cell` 1.20; it was upgraded by Dependabot. So the simplest workaround for the moment is to roll back the version. --- Cargo.lock | 13 ++----------- Cargo.toml | 2 +- 2 files changed, 3 insertions(+), 12 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 74910d9290..dfff8a71a6 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -2118,12 +2118,9 @@ dependencies = [ [[package]] name = "once_cell" -version = "1.20.1" +version = "1.19.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "82881c4be219ab5faaf2ad5e5e5ecdff8c66bd7402ca3160975c93b24961afd1" -dependencies = [ - "portable-atomic", -] +checksum = "3fdb12b2476b595f9358c5161aa467c2438859caa136dec86c26fdd2efe17b92" [[package]] name = "oorandom" @@ -2341,12 +2338,6 @@ version = "0.3.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "22686f4785f02a4fcc856d3b3bb19bf6c8160d103f7a99cc258bddd0251dc7f2" -[[package]] -name = "portable-atomic" -version = "1.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cc9c68a3f6da06753e9335d63e27f6b9754dd1920d941135b7ea8224f141adb2" - [[package]] name = "pp-rs" version = "0.2.1" diff --git a/Cargo.toml b/Cargo.toml index 68c29b671b..45c6fa6505 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -104,7 +104,7 @@ nanorand = { version = "0.7", default-features = false, features = ["wyrand"] } noise = { version = "0.8", git = "https://github.com/Razaekel/noise-rs.git", rev = "c6942d4fb70af26db4441edcf41f90fa115333f2" } nv-flip = "0.1" obj = "0.10" -once_cell = "1.20.1" +once_cell = "1.19.0" parking_lot = "0.12.1" pico-args = { version = "0.5.0", features = [ "eq-separator",