diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index a8a103e6d0..9755dfa334 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -32,6 +32,44 @@ const RAY_QUERY_FIELD_INTERSECTION: &str = "intersection"; const RAY_QUERY_FIELD_READY: &str = "ready"; const RAY_QUERY_FUN_MAP_INTERSECTION: &str = "_map_intersection_type"; +/// A custom little error type used by `checked_put_expression` to indicate that +/// it was halfway through writing a bounds check and then realised that it +/// needed to use an out-of-bounds local for the failure case. If you get this +/// error, you need to write it yourself. +/// +/// OOB locals should only ever be needed for function arguments, since every +/// other place where an expression goes doesn't allow pointers: the expression +/// always needs to be an index, or an image, or a boolean, etc, as well as a +/// handful of places that accept most types but explicitly exclude pointers +/// (e.g. return values). +/// +/// The one exception to this is `Statement::Emit`, but the code that writes +/// those has a special check that stops it from emitting pointers. This is why +/// the regular `put_expression` just returns an 'invalid module' error if this +/// happens: this is only a problem in function arguments, no need to clutter up +/// everything else. +/// +/// `checked_put_insert` can't just write the OOB local itself because it might +/// not be able to resolve the right type of OOB local that it needs: although +/// it'll always be able to resolve the `TypeInner` of the bounds-checked +/// expression, it needs a `Handle`, and `Type`s have names. That means +/// that if it can't guess the right name for the type, it can't get a handle, +/// and then it can't look up the name of the OOB local. +/// +/// Note that the only cases where it can't resolve the expression to a +/// `Handle` is when indexing into a matrix or a vector, yielding a vector +/// or a scalar respectively. In WGSL, I don't think these types can end up +/// having names, and so this isn't necessary when that's the source language; +/// however, I believe that in SPIR-V they can be assigned names, so it's +/// necessary for that. +#[derive(Debug, Clone, Copy)] +struct NeedsOobLocal { + /// Any extra text that needs to added after the bounds-checked expression. + /// + /// This is always just either "" or ")". + end: &'static str, +} + /// Write the Metal name for a Naga numeric type: scalar, vector, or matrix. /// /// The `sizes` slice determines whether this function writes a @@ -544,6 +582,32 @@ impl<'a> ExpressionContext<'a> { index::access_needs_check(base, index, self.module, self.function, self.info) } + /// Returns an iterator over the chain of `Access` and `AccessIndex` + /// expressions starting from `chain`. + /// + /// They're yielded as `(base, index)` pairs, where `base` is the expression + /// being indexed into and `index` is the index being used. + /// + /// The index is `None` if `base` is a struct, since you never need bounds + /// checks for accessing struct fields. + /// + /// If `chain` isn't an `Access` or `AccessIndex` expression, this just + /// yields nothing. + fn access_chain( + &self, + chain: Handle, + ) -> impl Iterator, Option)> + '_ { + index::access_chain(chain, self.module, self.function, self.info) + } + + /// Returns all the types which we need out-of-bounds locals for; that is, + /// all of the types which the code might attempt to get an out-of-bounds + /// pointer to, in which case we yield a pointer to the out-of-bounds local + /// of the correct type. + fn oob_locals(&self) -> FastHashSet> { + index::oob_locals(self.module, self.function, self.info) + } + fn get_packed_vec_kind( &self, expr_handle: Handle, @@ -625,6 +689,78 @@ impl Writer { Ok(()) } + /// Writes the local variables of the given function, as well as any extra + /// out-of-bounds locals that are needed. + /// + /// The names of the OOB locals are also added to `self.names` at the same + /// time. + fn put_locals( + &mut self, + module: &crate::Module, + origin: FunctionOrigin, + fun_info: &valid::FunctionInfo, + ) -> BackendResult { + let fun = match origin { + FunctionOrigin::Handle(handle) => &module.functions[handle], + FunctionOrigin::EntryPoint(idx) => &module.entry_points[usize::from(idx)].function, + }; + + let oob_locals = index::oob_locals(module, fun, fun_info); + for &ty in oob_locals.iter() { + let name_key = match origin { + FunctionOrigin::Handle(handle) => NameKey::FunctionOobLocal(handle, ty), + FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointOobLocal(idx, ty), + }; + self.names.insert(name_key, self.namer.call("oob")); + } + + for (name_key, ty, init) in fun + .local_variables + .iter() + .map(|(local_handle, local)| { + let name_key = match origin { + FunctionOrigin::Handle(handle) => NameKey::FunctionLocal(handle, local_handle), + FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointLocal(idx, local_handle), + }; + (name_key, local.ty, local.init) + }) + .chain(oob_locals.iter().map(|&ty| { + let name_key = match origin { + FunctionOrigin::Handle(handle) => NameKey::FunctionOobLocal(handle, ty), + FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointOobLocal(idx, ty), + }; + (name_key, ty, None) + })) + { + let ty_name = TypeContext { + handle: ty, + gctx: module.to_ctx(), + names: &self.names, + access: crate::StorageAccess::empty(), + binding: None, + first_time: false, + }; + write!( + self.out, + "{}{} {}", + back::INDENT, + ty_name, + self.names[&name_key] + )?; + match init { + Some(value) => { + write!(self.out, " = ")?; + self.put_const_expression(value, module)?; + } + None => { + write!(self.out, " = {{}}")?; + } + }; + writeln!(self.out, ";")?; + } + Ok(()) + } + fn put_level_of_detail( &mut self, level: LevelOfDetail, @@ -1299,12 +1435,31 @@ impl Writer { /// /// - Pass `false` if it is an operand of a `?:` operator, a `[]`, or really /// almost anything else. + #[track_caller] fn put_expression( &mut self, expr_handle: Handle, context: &ExpressionContext, is_scoped: bool, ) -> BackendResult { + self.checked_put_expression(expr_handle, context, is_scoped)? + .map_err(|_| Error::Validation) + } + + /// A version of `put_expression` which surfaces the additional potential + /// error of needing to insert an out-of-bounds local. + /// + /// OOB locals are only ever needed when evaluating function arguments, + /// which is why the regular `put_expression` silently returns an + /// 'invalid module' error if that error ever comes up: if an OOB local is + /// ever needed anywhere else, the module must be invalid, and so there's no + /// need to clutter up the rest of the code with explicit checks for that. + fn checked_put_expression( + &mut self, + expr_handle: Handle, + context: &ExpressionContext, + is_scoped: bool, + ) -> Result, Error> { // Add to the set in order to track the stack size. #[cfg(test)] #[allow(trivial_casts)] @@ -1313,7 +1468,7 @@ impl Writer { if let Some(name) = self.named_expressions.get(&expr_handle) { write!(self.out, "{name}")?; - return Ok(()); + return Ok(Ok(())); } let expression = &context.function.expressions[expr_handle]; @@ -1348,7 +1503,18 @@ impl Writer { { write!(self.out, " ? ")?; self.put_access_chain(expr_handle, policy, context)?; - write!(self.out, " : DefaultConstructible()")?; + write!(self.out, " : ")?; + + if context.resolve_type(base).pointer_space().is_some() { + // We can't just use `DefaultConstructible` if this is a pointer, so punt it to + // the caller to insert an out-of-bounds local instead. + // See the docs for `NeedsOobLocal` for more details. + return Ok(Err(NeedsOobLocal { + end: if is_scoped { "" } else { ")" }, + })); + } + + write!(self.out, "DefaultConstructible()")?; if !is_scoped { write!(self.out, ")")?; @@ -1694,7 +1860,9 @@ impl Writer { .. } => "dot", crate::TypeInner::Vector { size, .. } => { - return self.put_dot_product(arg, arg1.unwrap(), size as usize, context) + return self + .put_dot_product(arg, arg1.unwrap(), size as usize, context) + .map(Ok) } _ => unreachable!( "Correct TypeInner for dot product should be already validated" @@ -1927,7 +2095,7 @@ impl Writer { write!(self.out, "}}")?; } } - Ok(()) + Ok(Ok(())) } /// Used by expressions like Swizzle and Binary since they need packed_vec3's to be casted to a vec3 @@ -1995,7 +2163,7 @@ impl Writer { #[allow(unused_variables)] fn put_bounds_checks( &mut self, - mut chain: Handle, + chain: Handle, context: &ExpressionContext, level: back::Level, prefix: &'static str, @@ -2003,29 +2171,8 @@ impl Writer { let mut check_written = false; // Iterate over the access chain, handling each expression. - loop { - // Produce a `GuardedIndex`, so we can shared code between the - // `Access` and `AccessIndex` cases. - let (base, guarded_index) = match context.function.expressions[chain] { - crate::Expression::Access { base, index } => { - (base, Some(index::GuardedIndex::Expression(index))) - } - crate::Expression::AccessIndex { base, index } => { - // Don't try to check indices into structs. Validation already took - // care of them, and index::needs_guard doesn't handle that case. - let mut base_inner = context.resolve_type(base); - if let crate::TypeInner::Pointer { base, .. } = *base_inner { - base_inner = &context.module.types[base].inner; - } - match *base_inner { - crate::TypeInner::Struct { .. } => (base, None), - _ => (base, Some(index::GuardedIndex::Known(index))), - } - } - _ => break, - }; - - if let Some(index) = guarded_index { + for (base, index) in context.access_chain(chain) { + if let Some(index) = index { if let Some(length) = context.access_needs_check(base, index) { if check_written { write!(self.out, " && ")?; @@ -2053,8 +2200,6 @@ impl Writer { } } } - - chain = base } Ok(check_written) @@ -2737,11 +2882,34 @@ impl Writer { let fun_name = &self.names[&NameKey::Function(function)]; write!(self.out, "{fun_name}(")?; // first, write down the actual arguments - for (i, &handle) in arguments.iter().enumerate() { + for (i, (info, &handle)) in context.expression.module.functions[function] + .arguments + .iter() + .zip(arguments) + .enumerate() + { if i != 0 { write!(self.out, ", ")?; } - self.put_expression(handle, &context.expression, true)?; + if let Err(NeedsOobLocal { end }) = + self.checked_put_expression(handle, &context.expression, true)? + { + let crate::TypeInner::Pointer { base, .. } = + context.expression.module.types[info.ty].inner + else { + unreachable!() + }; + let name_key = match context.expression.origin { + FunctionOrigin::Handle(handle) => { + NameKey::FunctionOobLocal(handle, base) + } + FunctionOrigin::EntryPoint(idx) => { + NameKey::EntryPointOobLocal(idx, base) + } + }; + let name = &self.names[&name_key]; + write!(self.out, "{name}{end}")?; + } } // follow-up with any global resources used let mut separate = !arguments.is_empty(); @@ -3437,28 +3605,7 @@ impl Writer { writeln!(self.out, ") {{")?; - for (local_handle, local) in fun.local_variables.iter() { - let ty_name = TypeContext { - handle: local.ty, - gctx: module.to_ctx(), - names: &self.names, - access: crate::StorageAccess::empty(), - binding: None, - first_time: false, - }; - let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)]; - write!(self.out, "{}{} {}", back::INDENT, ty_name, local_name)?; - match local.init { - Some(value) => { - write!(self.out, " = ")?; - self.put_const_expression(value, module)?; - } - None => { - write!(self.out, " = {{}}")?; - } - }; - writeln!(self.out, ";")?; - } + self.put_locals(module, FunctionOrigin::Handle(fun_handle), fun_info)?; let guarded_indices = index::find_checked_indexes(module, fun, fun_info, options.bounds_check_policies); @@ -3995,28 +4142,7 @@ impl Writer { // Finally, declare all the local variables that we need //TODO: we can postpone this till the relevant expressions are emitted - for (local_handle, local) in fun.local_variables.iter() { - let name = &self.names[&NameKey::EntryPointLocal(ep_index as _, local_handle)]; - let ty_name = TypeContext { - handle: local.ty, - gctx: module.to_ctx(), - names: &self.names, - access: crate::StorageAccess::empty(), - binding: None, - first_time: false, - }; - write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?; - match local.init { - Some(value) => { - write!(self.out, " = ")?; - self.put_const_expression(value, module)?; - } - None => { - write!(self.out, " = {{}}")?; - } - }; - writeln!(self.out, ";")?; - } + self.put_locals(module, FunctionOrigin::EntryPoint(ep_index as _), fun_info)?; let guarded_indices = index::find_checked_indexes(module, fun, fun_info, options.bounds_check_policies); diff --git a/src/proc/index.rs b/src/proc/index.rs index af3221c0fe..f915578b26 100644 --- a/src/proc/index.rs +++ b/src/proc/index.rs @@ -2,8 +2,9 @@ Definitions for index bounds checking. */ -use crate::{valid, Handle, UniqueArena}; +use crate::{valid, FastHashSet, Handle, UniqueArena}; use bit_set::BitSet; +use std::iter::{self, zip}; /// How should code generated by Naga do bounds checks? /// @@ -339,6 +340,99 @@ pub fn access_needs_check( Some(length) } +/// Returns an iterator over the chain of `Access` and `AccessIndex` +/// expressions starting from `chain`. +/// +/// They're yielded as `(base, index)` pairs, where `base` is the expression +/// being indexed into and `index` is the index being used. +/// +/// The index is `None` if `base` is a struct, since you never need bounds +/// checks for accessing struct fields. +/// +/// If `chain` isn't an `Access` or `AccessIndex` expression, this just +/// yields nothing. +pub fn access_chain<'a>( + mut chain: Handle, + module: &'a crate::Module, + function: &'a crate::Function, + info: &'a valid::FunctionInfo, +) -> impl Iterator, Option)> + 'a { + iter::from_fn(move || { + let (base, index) = match function.expressions[chain] { + crate::Expression::Access { base, index } => { + (base, Some(GuardedIndex::Expression(index))) + } + crate::Expression::AccessIndex { base, index } => { + // Don't try to check indices into structs. Validation already took + // care of them, and needs_guard doesn't handle that case. + let mut base_inner = info[base].ty.inner_with(&module.types); + if let crate::TypeInner::Pointer { base, .. } = *base_inner { + base_inner = &module.types[base].inner; + } + match *base_inner { + crate::TypeInner::Struct { .. } => (base, None), + _ => (base, Some(GuardedIndex::Known(index))), + } + } + _ => return None, + }; + chain = base; + Some((base, index)) + }) +} + +/// Returns all the types which we need out-of-bounds locals for; that is, +/// all of the types which the code might attempt to get an out-of-bounds +/// pointer to, in which case we yield a pointer to the out-of-bounds local +/// of the correct type. +pub fn oob_locals( + module: &crate::Module, + function: &crate::Function, + info: &valid::FunctionInfo, +) -> FastHashSet> { + let mut result = FastHashSet::default(); + for statement in &function.body { + // The only situation in which we end up actually needing to create an + // out-of-bounds pointer is when passing one to a function. + // + // This is because pointers are never baked; so they're just inlined everywhere + // they're used. That means that loads can just return 0, and stores can just do + // nothing; functions are the only case where you actually *have* to produce a + // pointer. + if let crate::Statement::Call { + function: callee, + ref arguments, + .. + } = *statement + { + // Now go through the arguments of the function looking for pointers which need bounds checks. + for (arg_info, &arg) in zip(&module.functions[callee].arguments, arguments) { + match module.types[arg_info.ty].inner { + crate::TypeInner::ValuePointer { .. } => { + // `ValuePointer`s should only ever be used when resolving the types of + // expressions, since the arena can no longer be modified at that point; things + // in the arena should always use proper `Pointer`s. + unreachable!("`ValuePointer` found in arena") + } + crate::TypeInner::Pointer { base, .. } => { + if access_chain(arg, module, function, info).any(|(base, index)| { + index + .and_then(|index| { + access_needs_check(base, index, module, function, info) + }) + .is_some() + }) { + result.insert(base); + } + } + _ => continue, + }; + } + } + } + result +} + impl GuardedIndex { /// Make a `GuardedIndex::Known` from a `GuardedIndex::Expression` if possible. /// diff --git a/src/proc/namer.rs b/src/proc/namer.rs index 2f262d7d5f..7be00d2a32 100644 --- a/src/proc/namer.rs +++ b/src/proc/namer.rs @@ -14,9 +14,11 @@ pub enum NameKey { Function(Handle), FunctionArgument(Handle, u32), FunctionLocal(Handle, Handle), + FunctionOobLocal(Handle, Handle), EntryPoint(EntryPointIndex), EntryPointLocal(EntryPointIndex, Handle), EntryPointArgument(EntryPointIndex, u32), + EntryPointOobLocal(EntryPointIndex, Handle), } /// This processor assigns names to all the things in a module diff --git a/tests/in/access.param.ron b/tests/in/access.param.ron index 5cd8e79a48..4af789b9a9 100644 --- a/tests/in/access.param.ron +++ b/tests/in/access.param.ron @@ -1,4 +1,11 @@ ( + bounds_check_policies: ( + index: ReadZeroSkipWrite, + buffer: ReadZeroSkipWrite, + image_load: ReadZeroSkipWrite, + image_store: ReadZeroSkipWrite, + binding_array: ReadZeroSkipWrite, + ), spv: ( version: (1, 1), debug: true, diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index c9097f913b..9bab86dc1d 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -60,6 +60,14 @@ fn test_matrix_within_struct_accesses() { t.m[0][idx] = 20.0; t.m[idx][1] = 30.0; t.m[idx][idx] = 40.0; + + // passing pointers to a function + // FIXME: these are currently commented out because getting pointers to + // vector/matrix elements is broken in Metal and HLSL. + // let pl0 = read_from_private(&t.m[0][1]); + // let pl1 = read_from_private(&t.m[0][idx]); + // let pl2 = read_from_private(&t.m[idx][1]); + // let pl3 = read_from_private(&t.m[idx][idx]); } struct MatCx2InArray { @@ -97,12 +105,24 @@ fn test_matrix_within_array_within_struct_accesses() { t.am[0][0][idx] = 20.0; t.am[0][idx][1] = 30.0; t.am[0][idx][idx] = 40.0; + + // passing pointers to a function + // FIXME: these are currently commented out because getting pointers to + // vector/matrix elements is broken in Metal and HLSL. + // let pl0 = read_from_private(&t.am[0][0][1]); + // let pl1 = read_from_private(&t.am[0][0][idx]); + // let pl2 = read_from_private(&t.am[0][idx][1]); + // let pl3 = read_from_private(&t.am[0][idx][idx]); } fn read_from_private(foo: ptr) -> f32 { return *foo; } +fn read_i32_from_private(foo: ptr) -> i32 { + return *foo; +} + fn test_arr_as_arg(a: array, 5>) -> f32 { return a[4][9]; } @@ -133,6 +153,7 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var c2 = array(a, i32(b), 3, 4, 5); c2[vi + 1u] = 42; let value = c2[vi]; + let value_again = read_i32_from_private(&c2[vi]); test_arr_as_arg(array, 5>()); diff --git a/tests/out/analysis/access.info.ron b/tests/out/analysis/access.info.ron index 62e898d3c0..347e20e0ba 100644 --- a/tests/out/analysis/access.info.ron +++ b/tests/out/analysis/access.info.ron @@ -23,6 +23,7 @@ ("DATA | SIZED | COPY | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), ("SIZED | COPY | ARGUMENT"), + ("SIZED | COPY | ARGUMENT"), ("DATA | SIZED | COPY | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), ("DATA | SIZED | COPY | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), @@ -2873,7 +2874,46 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(25), + ty: Handle(24), + ), + ( + uniformity: ( + non_uniform_result: Some(1), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(3), + ), + ], + sampling: [], + ), + ( + flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE"), + uniformity: ( + non_uniform_result: Some(1), + requirements: (""), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + (""), + (""), + (""), + (""), + (""), + (""), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(1), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(26), ), ( uniformity: ( @@ -2894,7 +2934,7 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(24), + ty: Handle(25), ), ( uniformity: ( @@ -2945,7 +2985,7 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(31), + ty: Handle(32), ), ( uniformity: ( @@ -2987,7 +3027,7 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(33), + ty: Handle(34), ), ( uniformity: ( @@ -3046,7 +3086,7 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(32), + ty: Handle(33), ), ], sampling: [], @@ -3076,7 +3116,7 @@ non_uniform_result: Some(1), requirements: (""), ), - ref_count: 2, + ref_count: 3, assignable_global: None, ty: Handle(1), ), @@ -3567,17 +3607,17 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(28), + ty: Handle(29), ), ( uniformity: ( non_uniform_result: Some(41), requirements: (""), ), - ref_count: 3, + ref_count: 4, assignable_global: None, ty: Value(Pointer( - base: 28, + base: 29, space: Function, )), ), @@ -3647,6 +3687,27 @@ assignable_global: None, ty: Handle(3), ), + ( + uniformity: ( + non_uniform_result: Some(41), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Pointer( + base: 3, + space: Function, + )), + ), + ( + uniformity: ( + non_uniform_result: Some(1), + requirements: (""), + ), + ref_count: 0, + assignable_global: None, + ty: Handle(3), + ), ( uniformity: ( non_uniform_result: None, @@ -3654,7 +3715,7 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(25), + ty: Handle(26), ), ( uniformity: ( @@ -3723,7 +3784,7 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(26), + ty: Handle(27), ), ], sampling: [], @@ -4260,7 +4321,7 @@ ), ref_count: 1, assignable_global: None, - ty: Handle(32), + ty: Handle(33), ), ( uniformity: ( @@ -4270,7 +4331,7 @@ ref_count: 2, assignable_global: None, ty: Value(Pointer( - base: 32, + base: 33, space: Function, )), ), diff --git a/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/tests/out/glsl/access.assign_through_ptr.Compute.glsl index ab5bd9a3fb..ce5a575a68 100644 --- a/tests/out/glsl/access.assign_through_ptr.Compute.glsl +++ b/tests/out/glsl/access.assign_through_ptr.Compute.glsl @@ -27,6 +27,11 @@ float read_from_private(inout float foo_1) { return _e1; } +int read_i32_from_private(inout int foo_2) { + int _e1 = foo_2; + return _e1; +} + float test_arr_as_arg(float a[5][10]) { return a[4][9]; } @@ -36,8 +41,8 @@ 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; } diff --git a/tests/out/glsl/access.foo_frag.Fragment.glsl b/tests/out/glsl/access.foo_frag.Fragment.glsl index 3d52fa56b0..4788fee5c6 100644 --- a/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -35,6 +35,11 @@ float read_from_private(inout float foo_1) { return _e1; } +int read_i32_from_private(inout int foo_2) { + int _e1 = foo_2; + return _e1; +} + float test_arr_as_arg(float a[5][10]) { return a[4][9]; } @@ -44,8 +49,8 @@ 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; } diff --git a/tests/out/glsl/access.foo_vert.Vertex.glsl b/tests/out/glsl/access.foo_vert.Vertex.glsl index a25eda03bb..35781661e3 100644 --- a/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -112,6 +112,11 @@ float read_from_private(inout float foo_1) { return _e1; } +int read_i32_from_private(inout int foo_2) { + int _e1 = foo_2; + return _e1; +} + float test_arr_as_arg(float a[5][10]) { return a[4][9]; } @@ -121,8 +126,8 @@ 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; } @@ -144,7 +149,8 @@ void main() { c2_ = int[5](a_1, int(b), 3, 4, 5); c2_[(vi + 1u)] = 42; int value = c2_[vi]; - float _e48 = 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))); + int _e48 = read_i32_from_private(c2_[vi]); + float _e50 = 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))); gl_Position = vec4((_matrix * vec4(ivec4(value))), 2.0); gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w); return; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 5fdd97fbcd..203a6f320f 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -207,6 +207,12 @@ float read_from_private(inout float foo_1) return _expr1; } +int read_i32_from_private(inout int foo_2) +{ + int _expr1 = foo_2; + return _expr1; +} + float test_arr_as_arg(float a[5][10]) { return a[4][9]; @@ -224,9 +230,9 @@ 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; } @@ -268,7 +274,8 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position c2_ = Constructarray5_int_(a_1, int(b), 3, 4, 5); c2_[(vi + 1u)] = 42; int value = c2_[vi]; - const float _e48 = test_arr_as_arg((float[5][10])0); + const int _e48 = read_i32_from_private(c2_[vi]); + const float _e50 = test_arr_as_arg((float[5][10])0); return float4(mul(float4((value).xxxx), _matrix), 2.0); } diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index 7447249127..5363112fef 100644 --- a/tests/out/ir/access.ron +++ b/tests/out/ir/access.ron @@ -251,6 +251,13 @@ space: Function, ), ), + ( + name: None, + inner: Pointer( + base: 3, + space: Function, + ), + ), ( name: None, inner: Array( @@ -262,7 +269,7 @@ ( name: None, inner: Array( - base: 24, + base: 25, size: Constant(5), stride: 40, ), @@ -318,7 +325,7 @@ ( name: None, inner: Array( - base: 26, + base: 27, size: Constant(2), stride: 16, ), @@ -326,7 +333,7 @@ ( name: None, inner: Pointer( - base: 32, + base: 33, space: Function, ), ), @@ -1608,12 +1615,45 @@ ), ], ), + ( + name: Some("read_i32_from_private"), + arguments: [ + ( + name: Some("foo"), + ty: 24, + binding: None, + ), + ], + result: Some(( + ty: 3, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + Load( + pointer: 1, + ), + ], + named_expressions: { + 1: "foo", + }, + body: [ + Emit(( + start: 1, + end: 2, + )), + Return( + value: Some(2), + ), + ], + ), ( name: Some("test_arr_as_arg"), arguments: [ ( name: Some("a"), - ty: 25, + ty: 26, binding: None, ), ], @@ -1657,7 +1697,7 @@ arguments: [ ( name: Some("p"), - ty: 31, + ty: 32, binding: None, ), ], @@ -1685,7 +1725,7 @@ arguments: [ ( name: Some("foo"), - ty: 33, + ty: 34, binding: None, ), ], @@ -1704,7 +1744,7 @@ value: 4, ), Compose( - ty: 32, + ty: 33, components: [ 3, 5, @@ -1749,7 +1789,7 @@ ), ], result: Some(( - ty: 26, + ty: 27, binding: Some(BuiltIn(Position( invariant: false, ))), @@ -1762,7 +1802,7 @@ ), ( name: Some("c2"), - ty: 28, + ty: 29, init: None, ), ], @@ -1863,7 +1903,7 @@ Literal(I32(4)), Literal(I32(5)), Compose( - ty: 28, + ty: 29, components: [ 27, 36, @@ -1891,28 +1931,33 @@ Load( pointer: 46, ), - ZeroValue(25), + Access( + base: 41, + index: 1, + ), CallResult(4), + ZeroValue(26), + CallResult(5), Splat( size: Quad, value: 47, ), As( - expr: 50, + expr: 52, kind: Float, convert: Some(4), ), Binary( op: Multiply, left: 8, - right: 51, + right: 53, ), Literal(F32(2.0)), Compose( - ty: 26, + ty: 27, components: [ - 52, - 53, + 54, + 55, ], ), ], @@ -1928,6 +1973,7 @@ 34: "data_pointer", 35: "foo_value", 47: "value", + 49: "value_again", }, body: [ Store( @@ -2019,6 +2065,10 @@ start: 45, end: 47, )), + Emit(( + start: 47, + end: 48, + )), Call( function: 4, arguments: [ @@ -2026,16 +2076,23 @@ ], result: Some(49), ), + Call( + function: 5, + arguments: [ + 50, + ], + result: Some(51), + ), Emit(( - start: 49, - end: 52, + start: 51, + end: 54, )), Emit(( - start: 53, - end: 54, + start: 55, + end: 56, )), Return( - value: Some(54), + value: Some(56), ), ], ), @@ -2049,7 +2106,7 @@ name: Some("foo_frag"), arguments: [], result: Some(( - ty: 26, + ty: 27, binding: Some(Location( location: 0, interpolation: Some(Perspective), @@ -2244,7 +2301,7 @@ local_variables: [ ( name: Some("arr"), - ty: 32, + ty: 33, init: None, ), ], @@ -2260,7 +2317,7 @@ value: 3, ), Compose( - ty: 32, + ty: 33, components: [ 2, 4, @@ -2284,14 +2341,14 @@ value: 5, ), Call( - function: 5, + function: 6, arguments: [ 7, ], result: None, ), Call( - function: 6, + function: 7, arguments: [ 6, ], diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index e5d875dd19..b87a5a418b 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -3,6 +3,12 @@ #include using metal::uint; +struct DefaultConstructible { + template + operator T() && { + return T {}; + } +}; struct _mslBufferSizes { uint size1; @@ -45,16 +51,16 @@ struct type_15 { struct MatCx2InArray { type_15 am; }; -struct type_18 { +struct type_19 { float inner[10]; }; -struct type_19 { - type_18 inner[5]; +struct type_20 { + type_19 inner[5]; }; -struct type_22 { +struct type_23 { int inner[5]; }; -struct type_26 { +struct type_27 { metal::float4 inner[2]; }; @@ -69,30 +75,38 @@ void test_matrix_within_struct_accesses( metal::float3x2 l0_ = baz.m; metal::float2 l1_ = baz.m[0]; int _e15 = idx; - metal::float2 l2_ = baz.m[_e15]; + metal::float2 l2_ = uint(_e15) < 3 ? baz.m[_e15] : DefaultConstructible(); float l3_ = baz.m[0].y; int _e29 = idx; - float l4_ = baz.m[0][_e29]; + float l4_ = uint(_e29) < 2 ? baz.m[0][_e29] : DefaultConstructible(); int _e34 = idx; - float l5_ = baz.m[_e34].y; + float l5_ = uint(_e34) < 3 ? baz.m[_e34].y : DefaultConstructible(); int _e41 = idx; int _e43 = idx; - float l6_ = baz.m[_e41][_e43]; + float l6_ = uint(_e43) < 2 && uint(_e41) < 3 ? baz.m[_e41][_e43] : DefaultConstructible(); t = Baz {metal::float3x2(metal::float2(1.0), metal::float2(2.0), metal::float2(3.0))}; int _e56 = idx; idx = _e56 + 1; t.m = metal::float3x2(metal::float2(6.0), metal::float2(5.0), metal::float2(4.0)); t.m[0] = metal::float2(9.0); int _e72 = idx; - t.m[_e72] = metal::float2(90.0); + if (uint(_e72) < 3) { + t.m[_e72] = metal::float2(90.0); + } t.m[0].y = 10.0; int _e85 = idx; - t.m[0][_e85] = 20.0; + if (uint(_e85) < 2) { + t.m[0][_e85] = 20.0; + } int _e89 = idx; - t.m[_e89].y = 30.0; + if (uint(_e89) < 3) { + t.m[_e89].y = 30.0; + } int _e95 = idx; int _e97 = idx; - t.m[_e95][_e97] = 40.0; + if (uint(_e97) < 2 && uint(_e95) < 3) { + t.m[_e95][_e97] = 40.0; + } return; } @@ -108,15 +122,15 @@ void test_matrix_within_array_within_struct_accesses( metal::float4x2 l1_1 = nested_mat_cx2_.am.inner[0]; metal::float2 l2_1 = nested_mat_cx2_.am.inner[0][0]; int _e24 = idx_1; - metal::float2 l3_1 = nested_mat_cx2_.am.inner[0][_e24]; + metal::float2 l3_1 = uint(_e24) < 4 ? nested_mat_cx2_.am.inner[0][_e24] : DefaultConstructible(); float l4_1 = nested_mat_cx2_.am.inner[0][0].y; int _e42 = idx_1; - float l5_1 = nested_mat_cx2_.am.inner[0][0][_e42]; + float l5_1 = uint(_e42) < 2 ? nested_mat_cx2_.am.inner[0][0][_e42] : DefaultConstructible(); int _e49 = idx_1; - float l6_1 = nested_mat_cx2_.am.inner[0][_e49].y; + float l6_1 = uint(_e49) < 4 ? nested_mat_cx2_.am.inner[0][_e49].y : DefaultConstructible(); int _e58 = idx_1; int _e60 = idx_1; - float l7_ = nested_mat_cx2_.am.inner[0][_e58][_e60]; + float l7_ = uint(_e60) < 2 && uint(_e58) < 4 ? nested_mat_cx2_.am.inner[0][_e58][_e60] : DefaultConstructible(); t_1 = MatCx2InArray {type_15 {}}; int _e67 = idx_1; idx_1 = _e67 + 1; @@ -124,15 +138,23 @@ void test_matrix_within_array_within_struct_accesses( t_1.am.inner[0] = metal::float4x2(metal::float2(8.0), metal::float2(7.0), metal::float2(6.0), metal::float2(5.0)); t_1.am.inner[0][0] = metal::float2(9.0); int _e93 = idx_1; - t_1.am.inner[0][_e93] = metal::float2(90.0); + if (uint(_e93) < 4) { + t_1.am.inner[0][_e93] = metal::float2(90.0); + } t_1.am.inner[0][0].y = 10.0; int _e110 = idx_1; - t_1.am.inner[0][0][_e110] = 20.0; + if (uint(_e110) < 2) { + t_1.am.inner[0][0][_e110] = 20.0; + } int _e116 = idx_1; - t_1.am.inner[0][_e116].y = 30.0; + if (uint(_e116) < 4) { + t_1.am.inner[0][_e116].y = 30.0; + } int _e124 = idx_1; int _e126 = idx_1; - t_1.am.inner[0][_e124][_e126] = 40.0; + if (uint(_e126) < 2 && uint(_e124) < 4) { + t_1.am.inner[0][_e124][_e126] = 40.0; + } return; } @@ -143,8 +165,15 @@ float read_from_private( return _e1; } +int read_i32_from_private( + thread int& foo_2 +) { + int _e1 = foo_2; + return _e1; +} + float test_arr_as_arg( - type_19 a + type_20 a ) { return a.inner[4].inner[9]; } @@ -157,9 +186,9 @@ void assign_through_ptr_fn( } void assign_array_through_ptr_fn( - thread type_26& foo_2 + thread type_27& foo_3 ) { - foo_2 = type_26 {metal::float4(1.0), metal::float4(2.0)}; + foo_3 = type_27 {metal::float4(1.0), metal::float4(2.0)}; return; } @@ -177,7 +206,8 @@ vertex foo_vertOutput foo_vert( , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { float foo = {}; - type_22 c2_ = {}; + type_23 c2_ = {}; + int oob = {}; foo = 0.0; float baz_1 = foo; foo = 1.0; @@ -186,13 +216,18 @@ vertex foo_vertOutput foo_vert( metal::float4x3 _matrix = bar._matrix; type_9 arr_1 = bar.arr; float b = bar._matrix[3].x; - int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; + uint _e23 = (1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u; + int a_1 = uint(_e23) < 1 + (_buffer_sizes.size1 - 160 - 8) / 8 ? bar.data[_e23].value : DefaultConstructible(); metal::int2 c = qux; float _e34 = read_from_private(foo); - c2_ = type_22 {a_1, static_cast(b), 3, 4, 5}; - c2_.inner[vi + 1u] = 42; - int value = c2_.inner[vi]; - float _e48 = test_arr_as_arg(type_19 {}); + c2_ = type_23 {a_1, static_cast(b), 3, 4, 5}; + uint _e42 = vi + 1u; + if (uint(_e42) < 5) { + c2_.inner[_e42] = 42; + } + int value = uint(vi) < 5 ? c2_.inner[vi] : DefaultConstructible(); + int _e48 = read_i32_from_private(uint(vi) < 5 ? c2_.inner[vi] : oob); + float _e50 = test_arr_as_arg(type_20 {}); return foo_vertOutput { metal::float4(_matrix * static_cast(metal::int4(value)), 2.0) }; } @@ -208,7 +243,9 @@ fragment foo_fragOutput foo_frag( bar._matrix[1].z = 1.0; bar._matrix = metal::float4x3(metal::float3(0.0), metal::float3(1.0), metal::float3(2.0), metal::float3(3.0)); bar.arr = type_9 {metal::uint2(0u), metal::uint2(1u)}; - bar.data[1].value = 1; + if (uint(1) < 1 + (_buffer_sizes.size1 - 160 - 8) / 8) { + bar.data[1].value = 1; + } qux = metal::int2 {}; return foo_fragOutput { metal::float4(0.0) }; } @@ -222,8 +259,8 @@ kernel void assign_through_ptr( val = {}; } metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - type_26 arr = {}; - arr = type_26 {metal::float4(6.0), metal::float4(7.0)}; + type_27 arr = {}; + arr = type_27 {metal::float4(6.0), metal::float4(7.0)}; assign_through_ptr_fn(val); assign_array_through_ptr_fn(arr); return; diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index f7be2338c5..31fe6d43d4 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,16 +1,16 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 324 +; Bound: 330 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %237 "foo_vert" %232 %235 -OpEntryPoint Fragment %284 "foo_frag" %283 -OpEntryPoint GLCompute %304 "assign_through_ptr" %307 -OpExecutionMode %284 OriginUpperLeft -OpExecutionMode %304 LocalSize 1 1 1 +OpEntryPoint Vertex %243 "foo_vert" %238 %241 +OpEntryPoint Fragment %290 "foo_frag" %289 +OpEntryPoint GLCompute %310 "assign_through_ptr" %313 +OpExecutionMode %290 OriginUpperLeft +OpExecutionMode %310 LocalSize 1 1 1 OpMemberName %6 0 "a" OpMemberName %6 1 "b" OpMemberName %6 2 "c" @@ -28,13 +28,13 @@ OpMemberName %22 0 "m" OpName %22 "Baz" OpMemberName %26 0 "am" OpName %26 "MatCx2InArray" -OpName %45 "global_const" -OpName %47 "bar" -OpName %49 "baz" -OpName %52 "qux" -OpName %55 "nested_mat_cx2" -OpName %58 "val" -OpName %59 "idx" +OpName %46 "global_const" +OpName %48 "bar" +OpName %50 "baz" +OpName %53 "qux" +OpName %56 "nested_mat_cx2" +OpName %59 "val" +OpName %60 "idx" OpName %62 "t" OpName %66 "test_matrix_within_struct_accesses" OpName %137 "idx" @@ -42,19 +42,21 @@ OpName %138 "t" OpName %142 "test_matrix_within_array_within_struct_accesses" OpName %200 "foo" OpName %201 "read_from_private" -OpName %206 "a" -OpName %207 "test_arr_as_arg" -OpName %213 "p" -OpName %214 "assign_through_ptr_fn" -OpName %219 "foo" -OpName %220 "assign_array_through_ptr_fn" -OpName %226 "foo" -OpName %228 "c2" -OpName %232 "vi" -OpName %237 "foo_vert" -OpName %284 "foo_frag" -OpName %301 "arr" -OpName %304 "assign_through_ptr" +OpName %206 "foo" +OpName %207 "read_i32_from_private" +OpName %212 "a" +OpName %213 "test_arr_as_arg" +OpName %219 "p" +OpName %220 "assign_through_ptr_fn" +OpName %225 "foo" +OpName %226 "assign_array_through_ptr_fn" +OpName %232 "foo" +OpName %234 "c2" +OpName %238 "vi" +OpName %243 "foo_vert" +OpName %290 "foo_frag" +OpName %307 "arr" +OpName %310 "assign_through_ptr" OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %6 1 Offset 16 OpMemberDecorate %6 2 Offset 28 @@ -80,29 +82,29 @@ OpDecorate %25 ArrayStride 32 OpMemberDecorate %26 0 Offset 0 OpMemberDecorate %26 0 ColMajor OpMemberDecorate %26 0 MatrixStride 8 -OpDecorate %28 ArrayStride 4 -OpDecorate %29 ArrayStride 40 -OpDecorate %33 ArrayStride 4 -OpDecorate %36 ArrayStride 16 -OpDecorate %47 DescriptorSet 0 -OpDecorate %47 Binding 0 +OpDecorate %29 ArrayStride 4 +OpDecorate %30 ArrayStride 40 +OpDecorate %34 ArrayStride 4 +OpDecorate %37 ArrayStride 16 +OpDecorate %48 DescriptorSet 0 +OpDecorate %48 Binding 0 OpDecorate %20 Block -OpDecorate %49 DescriptorSet 0 -OpDecorate %49 Binding 1 -OpDecorate %50 Block -OpMemberDecorate %50 0 Offset 0 -OpDecorate %52 DescriptorSet 0 -OpDecorate %52 Binding 2 -OpDecorate %53 Block -OpMemberDecorate %53 0 Offset 0 -OpDecorate %55 DescriptorSet 0 -OpDecorate %55 Binding 3 -OpDecorate %56 Block -OpMemberDecorate %56 0 Offset 0 -OpDecorate %232 BuiltIn VertexIndex -OpDecorate %235 BuiltIn Position -OpDecorate %283 Location 0 -OpDecorate %307 BuiltIn LocalInvocationId +OpDecorate %50 DescriptorSet 0 +OpDecorate %50 Binding 1 +OpDecorate %51 Block +OpMemberDecorate %51 0 Offset 0 +OpDecorate %53 DescriptorSet 0 +OpDecorate %53 Binding 2 +OpDecorate %54 Block +OpMemberDecorate %54 0 Offset 0 +OpDecorate %56 DescriptorSet 0 +OpDecorate %56 Binding 3 +OpDecorate %57 Block +OpMemberDecorate %57 0 Offset 0 +OpDecorate %238 BuiltIn VertexIndex +OpDecorate %241 BuiltIn Position +OpDecorate %289 Location 0 +OpDecorate %313 BuiltIn LocalInvocationId %2 = OpTypeVoid %3 = OpTypeInt 32 0 %4 = OpTypeVector %3 3 @@ -129,38 +131,38 @@ OpDecorate %307 BuiltIn LocalInvocationId %25 = OpTypeArray %24 %14 %26 = OpTypeStruct %25 %27 = OpTypePointer Function %10 -%28 = OpTypeArray %10 %16 -%30 = OpConstant %3 5 -%29 = OpTypeArray %28 %30 -%31 = OpTypeVector %10 4 -%32 = OpTypePointer StorageBuffer %5 -%33 = OpTypeArray %5 %30 -%34 = OpTypeVector %5 4 -%35 = OpTypePointer Workgroup %3 -%36 = OpTypeArray %31 %14 -%37 = OpTypePointer Function %36 -%38 = OpConstant %3 0 -%39 = OpConstantComposite %4 %38 %38 %38 -%40 = OpConstant %5 0 -%41 = OpConstantComposite %6 %38 %39 %40 -%42 = OpConstant %5 2 -%43 = OpConstant %5 10 -%44 = OpConstant %5 5 -%46 = OpTypePointer Private %6 -%45 = OpVariable %46 Private %41 -%48 = OpTypePointer StorageBuffer %20 -%47 = OpVariable %48 StorageBuffer -%50 = OpTypeStruct %22 -%51 = OpTypePointer Uniform %50 -%49 = OpVariable %51 Uniform -%53 = OpTypeStruct %23 -%54 = OpTypePointer StorageBuffer %53 -%52 = OpVariable %54 StorageBuffer -%56 = OpTypeStruct %26 -%57 = OpTypePointer Uniform %56 -%55 = OpVariable %57 Uniform -%58 = OpVariable %35 Workgroup -%60 = OpTypePointer Function %5 +%28 = OpTypePointer Function %5 +%29 = OpTypeArray %10 %16 +%31 = OpConstant %3 5 +%30 = OpTypeArray %29 %31 +%32 = OpTypeVector %10 4 +%33 = OpTypePointer StorageBuffer %5 +%34 = OpTypeArray %5 %31 +%35 = OpTypeVector %5 4 +%36 = OpTypePointer Workgroup %3 +%37 = OpTypeArray %32 %14 +%38 = OpTypePointer Function %37 +%39 = OpConstant %3 0 +%40 = OpConstantComposite %4 %39 %39 %39 +%41 = OpConstant %5 0 +%42 = OpConstantComposite %6 %39 %40 %41 +%43 = OpConstant %5 2 +%44 = OpConstant %5 10 +%45 = OpConstant %5 5 +%47 = OpTypePointer Private %6 +%46 = OpVariable %47 Private %42 +%49 = OpTypePointer StorageBuffer %20 +%48 = OpVariable %49 StorageBuffer +%51 = OpTypeStruct %22 +%52 = OpTypePointer Uniform %51 +%50 = OpVariable %52 Uniform +%54 = OpTypeStruct %23 +%55 = OpTypePointer StorageBuffer %54 +%53 = OpVariable %55 StorageBuffer +%57 = OpTypeStruct %26 +%58 = OpTypePointer Uniform %57 +%56 = OpVariable %58 Uniform +%59 = OpVariable %36 Workgroup %61 = OpConstantNull %5 %63 = OpTypePointer Function %22 %64 = OpConstantNull %22 @@ -197,70 +199,71 @@ OpDecorate %307 BuiltIn LocalInvocationId %177 = OpTypePointer Function %25 %179 = OpTypePointer Function %24 %202 = OpTypeFunction %10 %27 -%208 = OpTypeFunction %10 %29 -%215 = OpTypeFunction %2 %35 -%216 = OpConstant %3 42 -%221 = OpTypeFunction %2 %37 -%227 = OpConstantNull %10 -%229 = OpTypePointer Function %33 -%230 = OpConstantNull %33 -%233 = OpTypePointer Input %3 -%232 = OpVariable %233 Input -%236 = OpTypePointer Output %31 -%235 = OpVariable %236 Output -%239 = OpTypePointer StorageBuffer %23 -%242 = OpConstant %10 0.0 -%243 = OpConstant %3 3 -%244 = OpConstant %5 3 -%245 = OpConstant %5 4 -%246 = OpConstant %5 42 -%247 = OpConstantNull %29 -%252 = OpTypePointer StorageBuffer %8 -%255 = OpTypePointer StorageBuffer %18 -%256 = OpConstant %3 4 -%259 = OpTypePointer StorageBuffer %9 -%260 = OpTypePointer StorageBuffer %10 -%263 = OpTypePointer StorageBuffer %19 -%266 = OpTypePointer StorageBuffer %7 -%283 = OpVariable %236 Output -%286 = OpConstantNull %23 -%302 = OpConstantNull %36 -%306 = OpConstantNull %3 -%308 = OpTypePointer Input %4 -%307 = OpVariable %308 Input -%310 = OpConstantNull %4 -%312 = OpTypeBool -%311 = OpTypeVector %312 3 -%317 = OpConstant %3 264 +%208 = OpTypeFunction %5 %28 +%214 = OpTypeFunction %10 %30 +%221 = OpTypeFunction %2 %36 +%222 = OpConstant %3 42 +%227 = OpTypeFunction %2 %38 +%233 = OpConstantNull %10 +%235 = OpTypePointer Function %34 +%236 = OpConstantNull %34 +%239 = OpTypePointer Input %3 +%238 = OpVariable %239 Input +%242 = OpTypePointer Output %32 +%241 = OpVariable %242 Output +%245 = OpTypePointer StorageBuffer %23 +%248 = OpConstant %10 0.0 +%249 = OpConstant %3 3 +%250 = OpConstant %5 3 +%251 = OpConstant %5 4 +%252 = OpConstant %5 42 +%253 = OpConstantNull %30 +%258 = OpTypePointer StorageBuffer %8 +%261 = OpTypePointer StorageBuffer %18 +%262 = OpConstant %3 4 +%265 = OpTypePointer StorageBuffer %9 +%266 = OpTypePointer StorageBuffer %10 +%269 = OpTypePointer StorageBuffer %19 +%272 = OpTypePointer StorageBuffer %7 +%289 = OpVariable %242 Output +%292 = OpConstantNull %23 +%308 = OpConstantNull %37 +%312 = OpConstantNull %3 +%314 = OpTypePointer Input %4 +%313 = OpVariable %314 Input +%316 = OpConstantNull %4 +%318 = OpTypeBool +%317 = OpTypeVector %318 3 +%323 = OpConstant %3 264 %66 = OpFunction %2 None %67 %65 = OpLabel -%59 = OpVariable %60 Function %61 +%60 = OpVariable %28 Function %61 %62 = OpVariable %63 Function %64 -%69 = OpAccessChain %68 %49 %38 +%69 = OpAccessChain %68 %50 %39 OpBranch %83 %83 = OpLabel -OpStore %59 %70 -%84 = OpLoad %5 %59 +OpStore %60 %70 +%84 = OpLoad %5 %60 %85 = OpISub %5 %84 %70 -OpStore %59 %85 -%87 = OpAccessChain %86 %69 %38 +OpStore %60 %85 +%87 = OpAccessChain %86 %69 %39 %88 = OpLoad %21 %87 -%90 = OpAccessChain %89 %69 %38 %38 +%90 = OpAccessChain %89 %69 %39 %39 %91 = OpLoad %12 %90 -%92 = OpLoad %5 %59 -%93 = OpAccessChain %89 %69 %38 %92 +%92 = OpLoad %5 %60 +%93 = OpAccessChain %89 %69 %39 %92 %94 = OpLoad %12 %93 -%97 = OpAccessChain %95 %69 %38 %38 %96 +%97 = OpAccessChain %95 %69 %39 %39 %96 %98 = OpLoad %10 %97 -%99 = OpLoad %5 %59 -%100 = OpAccessChain %95 %69 %38 %38 %99 +%99 = OpLoad %5 %60 +%100 = OpAccessChain %95 %69 %39 %39 %99 %101 = OpLoad %10 %100 -%102 = OpLoad %5 %59 -%103 = OpAccessChain %95 %69 %38 %102 %96 +%102 = OpLoad %5 %60 +%103 = OpAccessChain %95 %69 %39 %102 %96 %104 = OpLoad %10 %103 -%105 = OpLoad %5 %59 -%106 = OpLoad %5 %59 -%107 = OpAccessChain %95 %69 %38 %105 %106 +%105 = OpLoad %5 %60 +%106 = OpLoad %5 %60 +%107 = OpAccessChain %95 %69 %39 %105 %106 %108 = OpLoad %10 %107 %109 = OpCompositeConstruct %12 %71 %71 %110 = OpCompositeConstruct %12 %72 %72 @@ -268,100 +271,100 @@ OpStore %59 %85 %112 = OpCompositeConstruct %21 %109 %110 %111 %113 = OpCompositeConstruct %22 %112 OpStore %62 %113 -%114 = OpLoad %5 %59 +%114 = OpLoad %5 %60 %115 = OpIAdd %5 %114 %70 -OpStore %59 %115 +OpStore %60 %115 %117 = OpCompositeConstruct %12 %74 %74 %118 = OpCompositeConstruct %12 %75 %75 %119 = OpCompositeConstruct %12 %76 %76 %120 = OpCompositeConstruct %21 %117 %118 %119 -%121 = OpAccessChain %116 %62 %38 +%121 = OpAccessChain %116 %62 %39 OpStore %121 %120 %123 = OpCompositeConstruct %12 %77 %77 -%124 = OpAccessChain %122 %62 %38 %38 +%124 = OpAccessChain %122 %62 %39 %39 OpStore %124 %123 -%125 = OpLoad %5 %59 +%125 = OpLoad %5 %60 %126 = OpCompositeConstruct %12 %78 %78 -%127 = OpAccessChain %122 %62 %38 %125 +%127 = OpAccessChain %122 %62 %39 %125 OpStore %127 %126 -%129 = OpAccessChain %128 %62 %38 %38 %96 +%129 = OpAccessChain %128 %62 %39 %39 %96 OpStore %129 %79 -%130 = OpLoad %5 %59 -%131 = OpAccessChain %128 %62 %38 %38 %130 +%130 = OpLoad %5 %60 +%131 = OpAccessChain %128 %62 %39 %39 %130 OpStore %131 %80 -%132 = OpLoad %5 %59 -%133 = OpAccessChain %128 %62 %38 %132 %96 +%132 = OpLoad %5 %60 +%133 = OpAccessChain %128 %62 %39 %132 %96 OpStore %133 %81 -%134 = OpLoad %5 %59 -%135 = OpLoad %5 %59 -%136 = OpAccessChain %128 %62 %38 %134 %135 +%134 = OpLoad %5 %60 +%135 = OpLoad %5 %60 +%136 = OpAccessChain %128 %62 %39 %134 %135 OpStore %136 %82 OpReturn OpFunctionEnd %142 = OpFunction %2 None %67 %141 = OpLabel -%137 = OpVariable %60 Function %61 +%137 = OpVariable %28 Function %61 %138 = OpVariable %139 Function %140 -%144 = OpAccessChain %143 %55 %38 +%144 = OpAccessChain %143 %56 %39 OpBranch %148 %148 = OpLabel OpStore %137 %70 %149 = OpLoad %5 %137 %150 = OpISub %5 %149 %70 OpStore %137 %150 -%152 = OpAccessChain %151 %144 %38 +%152 = OpAccessChain %151 %144 %39 %153 = OpLoad %25 %152 -%155 = OpAccessChain %154 %144 %38 %38 +%155 = OpAccessChain %154 %144 %39 %39 %156 = OpLoad %24 %155 -%157 = OpAccessChain %89 %144 %38 %38 %38 +%157 = OpAccessChain %89 %144 %39 %39 %39 %158 = OpLoad %12 %157 %159 = OpLoad %5 %137 -%160 = OpAccessChain %89 %144 %38 %38 %159 +%160 = OpAccessChain %89 %144 %39 %39 %159 %161 = OpLoad %12 %160 -%162 = OpAccessChain %95 %144 %38 %38 %38 %96 +%162 = OpAccessChain %95 %144 %39 %39 %39 %96 %163 = OpLoad %10 %162 %164 = OpLoad %5 %137 -%165 = OpAccessChain %95 %144 %38 %38 %38 %164 +%165 = OpAccessChain %95 %144 %39 %39 %39 %164 %166 = OpLoad %10 %165 %167 = OpLoad %5 %137 -%168 = OpAccessChain %95 %144 %38 %38 %167 %96 +%168 = OpAccessChain %95 %144 %39 %39 %167 %96 %169 = OpLoad %10 %168 %170 = OpLoad %5 %137 %171 = OpLoad %5 %137 -%172 = OpAccessChain %95 %144 %38 %38 %170 %171 +%172 = OpAccessChain %95 %144 %39 %39 %170 %171 %173 = OpLoad %10 %172 %174 = OpCompositeConstruct %26 %145 OpStore %138 %174 %175 = OpLoad %5 %137 %176 = OpIAdd %5 %175 %70 OpStore %137 %176 -%178 = OpAccessChain %177 %138 %38 +%178 = OpAccessChain %177 %138 %39 OpStore %178 %145 %180 = OpCompositeConstruct %12 %146 %146 %181 = OpCompositeConstruct %12 %147 %147 %182 = OpCompositeConstruct %12 %74 %74 %183 = OpCompositeConstruct %12 %75 %75 %184 = OpCompositeConstruct %24 %180 %181 %182 %183 -%185 = OpAccessChain %179 %138 %38 %38 +%185 = OpAccessChain %179 %138 %39 %39 OpStore %185 %184 %186 = OpCompositeConstruct %12 %77 %77 -%187 = OpAccessChain %122 %138 %38 %38 %38 +%187 = OpAccessChain %122 %138 %39 %39 %39 OpStore %187 %186 %188 = OpLoad %5 %137 %189 = OpCompositeConstruct %12 %78 %78 -%190 = OpAccessChain %122 %138 %38 %38 %188 +%190 = OpAccessChain %122 %138 %39 %39 %188 OpStore %190 %189 -%191 = OpAccessChain %128 %138 %38 %38 %38 %96 +%191 = OpAccessChain %128 %138 %39 %39 %39 %96 OpStore %191 %79 %192 = OpLoad %5 %137 -%193 = OpAccessChain %128 %138 %38 %38 %38 %192 +%193 = OpAccessChain %128 %138 %39 %39 %39 %192 OpStore %193 %80 %194 = OpLoad %5 %137 -%195 = OpAccessChain %128 %138 %38 %38 %194 %96 +%195 = OpAccessChain %128 %138 %39 %39 %194 %96 OpStore %195 %81 %196 = OpLoad %5 %137 %197 = OpLoad %5 %137 -%198 = OpAccessChain %128 %138 %38 %38 %196 %197 +%198 = OpAccessChain %128 %138 %39 %39 %196 %197 OpStore %198 %82 OpReturn OpFunctionEnd @@ -373,125 +376,133 @@ OpBranch %203 %204 = OpLoad %10 %200 OpReturnValue %204 OpFunctionEnd -%207 = OpFunction %10 None %208 -%206 = OpFunctionParameter %29 +%207 = OpFunction %5 None %208 +%206 = OpFunctionParameter %28 %205 = OpLabel OpBranch %209 %209 = OpLabel -%210 = OpCompositeExtract %28 %206 4 -%211 = OpCompositeExtract %10 %210 9 -OpReturnValue %211 +%210 = OpLoad %5 %206 +OpReturnValue %210 OpFunctionEnd -%214 = OpFunction %2 None %215 -%213 = OpFunctionParameter %35 -%212 = OpLabel -OpBranch %217 -%217 = OpLabel -OpStore %213 %216 -OpReturn +%213 = OpFunction %10 None %214 +%212 = OpFunctionParameter %30 +%211 = OpLabel +OpBranch %215 +%215 = OpLabel +%216 = OpCompositeExtract %29 %212 4 +%217 = OpCompositeExtract %10 %216 9 +OpReturnValue %217 OpFunctionEnd %220 = OpFunction %2 None %221 -%219 = OpFunctionParameter %37 +%219 = OpFunctionParameter %36 %218 = OpLabel -OpBranch %222 -%222 = OpLabel -%223 = OpCompositeConstruct %31 %71 %71 %71 %71 -%224 = OpCompositeConstruct %31 %72 %72 %72 %72 -%225 = OpCompositeConstruct %36 %223 %224 -OpStore %219 %225 +OpBranch %223 +%223 = OpLabel +OpStore %219 %222 +OpReturn +OpFunctionEnd +%226 = OpFunction %2 None %227 +%225 = OpFunctionParameter %38 +%224 = OpLabel +OpBranch %228 +%228 = OpLabel +%229 = OpCompositeConstruct %32 %71 %71 %71 %71 +%230 = OpCompositeConstruct %32 %72 %72 %72 %72 +%231 = OpCompositeConstruct %37 %229 %230 +OpStore %225 %231 OpReturn OpFunctionEnd -%237 = OpFunction %2 None %67 -%231 = OpLabel -%226 = OpVariable %27 Function %227 -%228 = OpVariable %229 Function %230 -%234 = OpLoad %3 %232 -%238 = OpAccessChain %68 %49 %38 -%240 = OpAccessChain %239 %52 %38 -%241 = OpAccessChain %143 %55 %38 -OpBranch %248 -%248 = OpLabel -OpStore %226 %242 -%249 = OpLoad %10 %226 -OpStore %226 %71 -%250 = OpFunctionCall %2 %66 -%251 = OpFunctionCall %2 %142 -%253 = OpAccessChain %252 %47 %38 -%254 = OpLoad %8 %253 -%257 = OpAccessChain %255 %47 %256 -%258 = OpLoad %18 %257 -%261 = OpAccessChain %260 %47 %38 %243 %38 -%262 = OpLoad %10 %261 -%264 = OpArrayLength %3 %47 5 -%265 = OpISub %3 %264 %14 -%267 = OpAccessChain %32 %47 %30 %265 %38 -%268 = OpLoad %5 %267 -%269 = OpLoad %23 %240 -%270 = OpFunctionCall %10 %201 %226 -%271 = OpConvertFToS %5 %262 -%272 = OpCompositeConstruct %33 %268 %271 %244 %245 %44 -OpStore %228 %272 -%273 = OpIAdd %3 %234 %96 -%274 = OpAccessChain %60 %228 %273 -OpStore %274 %246 -%275 = OpAccessChain %60 %228 %234 -%276 = OpLoad %5 %275 -%277 = OpFunctionCall %10 %207 %247 -%278 = OpCompositeConstruct %34 %276 %276 %276 %276 -%279 = OpConvertSToF %31 %278 -%280 = OpMatrixTimesVector %9 %254 %279 -%281 = OpCompositeConstruct %31 %280 %72 -OpStore %235 %281 +%243 = OpFunction %2 None %67 +%237 = OpLabel +%232 = OpVariable %27 Function %233 +%234 = OpVariable %235 Function %236 +%240 = OpLoad %3 %238 +%244 = OpAccessChain %68 %50 %39 +%246 = OpAccessChain %245 %53 %39 +%247 = OpAccessChain %143 %56 %39 +OpBranch %254 +%254 = OpLabel +OpStore %232 %248 +%255 = OpLoad %10 %232 +OpStore %232 %71 +%256 = OpFunctionCall %2 %66 +%257 = OpFunctionCall %2 %142 +%259 = OpAccessChain %258 %48 %39 +%260 = OpLoad %8 %259 +%263 = OpAccessChain %261 %48 %262 +%264 = OpLoad %18 %263 +%267 = OpAccessChain %266 %48 %39 %249 %39 +%268 = OpLoad %10 %267 +%270 = OpArrayLength %3 %48 5 +%271 = OpISub %3 %270 %14 +%273 = OpAccessChain %33 %48 %31 %271 %39 +%274 = OpLoad %5 %273 +%275 = OpLoad %23 %246 +%276 = OpFunctionCall %10 %201 %232 +%277 = OpConvertFToS %5 %268 +%278 = OpCompositeConstruct %34 %274 %277 %250 %251 %45 +OpStore %234 %278 +%279 = OpIAdd %3 %240 %96 +%280 = OpAccessChain %28 %234 %279 +OpStore %280 %252 +%281 = OpAccessChain %28 %234 %240 +%282 = OpLoad %5 %281 +%283 = OpFunctionCall %10 %213 %253 +%284 = OpCompositeConstruct %35 %282 %282 %282 %282 +%285 = OpConvertSToF %32 %284 +%286 = OpMatrixTimesVector %9 %260 %285 +%287 = OpCompositeConstruct %32 %286 %72 +OpStore %241 %287 OpReturn OpFunctionEnd -%284 = OpFunction %2 None %67 -%282 = OpLabel -%285 = OpAccessChain %239 %52 %38 -OpBranch %287 -%287 = OpLabel -%288 = OpAccessChain %260 %47 %38 %96 %14 -OpStore %288 %71 -%289 = OpCompositeConstruct %9 %242 %242 %242 -%290 = OpCompositeConstruct %9 %71 %71 %71 -%291 = OpCompositeConstruct %9 %72 %72 %72 -%292 = OpCompositeConstruct %9 %73 %73 %73 -%293 = OpCompositeConstruct %8 %289 %290 %291 %292 -%294 = OpAccessChain %252 %47 %38 -OpStore %294 %293 -%295 = OpCompositeConstruct %17 %38 %38 -%296 = OpCompositeConstruct %17 %96 %96 -%297 = OpCompositeConstruct %18 %295 %296 -%298 = OpAccessChain %255 %47 %256 -OpStore %298 %297 -%299 = OpAccessChain %32 %47 %30 %96 %38 -OpStore %299 %70 -OpStore %285 %286 -%300 = OpCompositeConstruct %31 %242 %242 %242 %242 -OpStore %283 %300 +%290 = OpFunction %2 None %67 +%288 = OpLabel +%291 = OpAccessChain %245 %53 %39 +OpBranch %293 +%293 = OpLabel +%294 = OpAccessChain %266 %48 %39 %96 %14 +OpStore %294 %71 +%295 = OpCompositeConstruct %9 %248 %248 %248 +%296 = OpCompositeConstruct %9 %71 %71 %71 +%297 = OpCompositeConstruct %9 %72 %72 %72 +%298 = OpCompositeConstruct %9 %73 %73 %73 +%299 = OpCompositeConstruct %8 %295 %296 %297 %298 +%300 = OpAccessChain %258 %48 %39 +OpStore %300 %299 +%301 = OpCompositeConstruct %17 %39 %39 +%302 = OpCompositeConstruct %17 %96 %96 +%303 = OpCompositeConstruct %18 %301 %302 +%304 = OpAccessChain %261 %48 %262 +OpStore %304 %303 +%305 = OpAccessChain %33 %48 %31 %96 %39 +OpStore %305 %70 +OpStore %291 %292 +%306 = OpCompositeConstruct %32 %248 %248 %248 %248 +OpStore %289 %306 OpReturn OpFunctionEnd -%304 = OpFunction %2 None %67 -%303 = OpLabel -%301 = OpVariable %37 Function %302 -OpBranch %305 -%305 = OpLabel -%309 = OpLoad %4 %307 -%313 = OpIEqual %311 %309 %310 -%314 = OpAll %312 %313 -OpSelectionMerge %315 None -OpBranchConditional %314 %316 %315 -%316 = OpLabel -OpStore %58 %306 -OpBranch %315 -%315 = OpLabel -OpControlBarrier %14 %14 %317 -OpBranch %318 -%318 = OpLabel -%319 = OpCompositeConstruct %31 %74 %74 %74 %74 -%320 = OpCompositeConstruct %31 %147 %147 %147 %147 -%321 = OpCompositeConstruct %36 %319 %320 -OpStore %301 %321 -%322 = OpFunctionCall %2 %214 %58 -%323 = OpFunctionCall %2 %220 %301 +%310 = OpFunction %2 None %67 +%309 = OpLabel +%307 = OpVariable %38 Function %308 +OpBranch %311 +%311 = OpLabel +%315 = OpLoad %4 %313 +%319 = OpIEqual %317 %315 %316 +%320 = OpAll %318 %319 +OpSelectionMerge %321 None +OpBranchConditional %320 %322 %321 +%322 = OpLabel +OpStore %59 %312 +OpBranch %321 +%321 = OpLabel +OpControlBarrier %14 %14 %323 +OpBranch %324 +%324 = OpLabel +%325 = OpCompositeConstruct %32 %74 %74 %74 %74 +%326 = OpCompositeConstruct %32 %147 %147 %147 %147 +%327 = OpCompositeConstruct %37 %325 %326 +OpStore %307 %327 +%328 = OpFunctionCall %2 %220 %59 +%329 = OpFunctionCall %2 %226 %307 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 77ebe1ee18..19de9a40f4 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -117,6 +117,11 @@ fn read_from_private(foo_1: ptr) -> f32 { return _e1; } +fn read_i32_from_private(foo_2: ptr) -> i32 { + let _e1 = (*foo_2); + return _e1; +} + fn test_arr_as_arg(a: array, 5>) -> f32 { return a[4][9]; } @@ -126,8 +131,8 @@ fn assign_through_ptr_fn(p: ptr) { return; } -fn assign_array_through_ptr_fn(foo_2: ptr, 2>>) { - (*foo_2) = array, 2>(vec4(1.0), vec4(2.0)); +fn assign_array_through_ptr_fn(foo_3: ptr, 2>>) { + (*foo_3) = array, 2>(vec4(1.0), vec4(2.0)); return; } @@ -151,7 +156,8 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { c2_ = array(a_1, i32(b), 3, 4, 5); c2_[(vi + 1u)] = 42; let value = c2_[vi]; - let _e48 = test_arr_as_arg(array, 5>()); + let _e48 = read_i32_from_private((&c2_[vi])); + let _e50 = test_arr_as_arg(array, 5>()); return vec4((_matrix * vec4(vec4(value))), 2.0); } diff --git a/tests/snapshots.rs b/tests/snapshots.rs index e1e1144139..31f4033880 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -519,8 +519,8 @@ fn convert_wgsl() { ), ( "access", - Targets::SPIRV - | Targets::METAL + // Targets::SPIRV | + Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL