From a7afb562761e7039cc66f0545d9eec127df3b562 Mon Sep 17 00:00:00 2001 From: Andy Leiserson Date: Tue, 11 Mar 2025 18:27:02 -0700 Subject: [PATCH] [msl-out] Fix `ReadZeroSkipWrite` bounds check mode for pointer arguments Fixes #4541 -- Co-authored-by: Liam Murphy Co-Authored-By: Erich Gubler --- naga/src/back/msl/mod.rs | 14 ++ naga/src/back/msl/writer.rs | 69 ++++++++- naga/src/proc/index.rs | 59 +++++++- naga/src/proc/namer.rs | 10 ++ .../wgsl/pointer-function-arg-restrict.toml | 4 + .../wgsl/pointer-function-arg-restrict.wgsl | 61 ++++++++ .../in/wgsl/pointer-function-arg-rzsw.toml | 4 + .../in/wgsl/pointer-function-arg-rzsw.wgsl | 61 ++++++++ naga/tests/in/wgsl/pointer-function-arg.toml | 1 + naga/tests/in/wgsl/pointer-function-arg.wgsl | 64 ++++++++ ...gsl-pointer-function-arg.main.Compute.glsl | 80 ++++++++++ .../out/hlsl/wgsl-pointer-function-arg.hlsl | 93 ++++++++++++ .../out/hlsl/wgsl-pointer-function-arg.ron | 12 ++ .../wgsl-pointer-function-arg-restrict.msl | 123 ++++++++++++++++ .../msl/wgsl-pointer-function-arg-rzsw.msl | 139 ++++++++++++++++++ .../out/msl/wgsl-pointer-function-arg.msl | 128 ++++++++++++++++ .../out/wgsl/wgsl-pointer-function-arg.wgsl | 76 ++++++++++ 17 files changed, 994 insertions(+), 4 deletions(-) create mode 100644 naga/tests/in/wgsl/pointer-function-arg-restrict.toml create mode 100644 naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl create mode 100644 naga/tests/in/wgsl/pointer-function-arg-rzsw.toml create mode 100644 naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl create mode 100644 naga/tests/in/wgsl/pointer-function-arg.toml create mode 100644 naga/tests/in/wgsl/pointer-function-arg.wgsl create mode 100644 naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl create mode 100644 naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl create mode 100644 naga/tests/out/hlsl/wgsl-pointer-function-arg.ron create mode 100644 naga/tests/out/msl/wgsl-pointer-function-arg-restrict.msl create mode 100644 naga/tests/out/msl/wgsl-pointer-function-arg-rzsw.msl create mode 100644 naga/tests/out/msl/wgsl-pointer-function-arg.msl create mode 100644 naga/tests/out/wgsl/wgsl-pointer-function-arg.wgsl diff --git a/naga/src/back/msl/mod.rs b/naga/src/back/msl/mod.rs index 376662a7d0..01ac1ac419 100644 --- a/naga/src/back/msl/mod.rs +++ b/naga/src/back/msl/mod.rs @@ -29,6 +29,20 @@ holding the result. [msl]: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf [all-atom]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS +## Pointer-typed bounds-checked expressions and OOB locals + +MSL (unlike HLSL and GLSL) has native support for pointer-typed function +arguments. When the [`BoundsCheckPolicy`] is `ReadZeroSkipWrite` and an +out-of-bounds index expression is used for such an argument, our strategy is to +pass a pointer to a dummy variable. These dummy variables are called "OOB +locals". We emit at most one OOB local per function for each type, since all +expressions producing a result of that type can share the same OOB local. (Note +that the OOB local mechanism is not actually implementing "skip write", nor even +"read zero" in some cases of read-after-write, but doing so would require +additional effort and the difference is unlikely to matter.) + +[`BoundsCheckPolicy`]: crate::proc::BoundsCheckPolicy + */ use alloc::{ diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 260dd8fd2f..721595a375 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -612,6 +612,17 @@ trait NameKeyExt { FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointLocal(idx, local_handle), } } + + /// Return the name key for a local variable used by ReadZeroSkipWrite bounds-check + /// policy when it needs to produce a pointer-typed result for an OOB access. These + /// are unique per accessed type, so the second argument is a type handle. See docs + /// for [`crate::back::msl`]. + fn oob_local_for_type(origin: FunctionOrigin, ty: Handle) -> NameKey { + match origin { + FunctionOrigin::Handle(handle) => NameKey::FunctionOobLocal(handle, ty), + FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointOobLocal(idx, ty), + } + } } impl NameKeyExt for NameKey {} @@ -722,6 +733,11 @@ impl<'a> ExpressionContext<'a> { index::bounds_check_iter(chain, self.module, self.function, self.info) } + /// See docs for [`proc::index::oob_local_types`]. + fn oob_local_types(&self) -> FastHashSet> { + index::oob_local_types(self.module, self.function, self.info, self.policies) + } + fn get_packed_vec_kind(&self, expr_handle: Handle) -> Option { match self.function.expressions[expr_handle] { crate::Expression::AccessIndex { base, index } => { @@ -929,8 +945,18 @@ impl Writer { Ok(()) } - /// Writes the local variables of the given function. + /// 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, context: &ExpressionContext) -> BackendResult { + let oob_local_types = context.oob_local_types(); + for &ty in oob_local_types.iter() { + let name_key = NameKey::oob_local_for_type(context.origin, ty); + self.names.insert(name_key, self.namer.call("oob")); + } + for (name_key, ty, init) in context .function .local_variables @@ -939,6 +965,10 @@ impl Writer { let name_key = NameKey::local(context.origin, local_handle); (name_key, local.ty, local.init) }) + .chain(oob_local_types.iter().map(|&ty| { + let name_key = NameKey::oob_local_for_type(context.origin, ty); + (name_key, ty, None) + })) { let ty_name = TypeContext { handle: ty, @@ -1761,7 +1791,42 @@ 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. + // Instead, we create a dummy local variable to serve as pointer + // target if the access is out of bounds. + let result_ty = context.info[expr_handle] + .ty + .inner_with(&context.module.types) + .pointer_base_type(); + let result_ty_handle = match result_ty { + Some(TypeResolution::Handle(handle)) => handle, + Some(TypeResolution::Value(_)) => { + // As long as the result of a pointer access expression is + // passed to a function or stored in a let binding, the + // type will be in the arena. If additional uses of + // pointers become valid, this assumption might no longer + // hold. Note that the LHS of a load or store doesn't + // take this path -- there is dedicated code in `put_load` + // and `put_store`. + unreachable!( + "Expected type {result_ty:?} of access through pointer type {base:?} to be in the arena", + ); + } + None => { + unreachable!( + "Expected access through pointer type {base:?} to return a pointer, but got {result_ty:?}", + ) + } + }; + let name_key = + NameKey::oob_local_for_type(context.origin, result_ty_handle); + self.out.write_str(&self.names[&name_key])?; + } else { + write!(self.out, "DefaultConstructible()")?; + } if !is_scoped { write!(self.out, ")")?; diff --git a/naga/src/proc/index.rs b/naga/src/proc/index.rs index 5508145734..8e76c85241 100644 --- a/naga/src/proc/index.rs +++ b/naga/src/proc/index.rs @@ -2,10 +2,10 @@ Definitions for index bounds checking. */ -use core::iter; +use core::iter::{self, zip}; use crate::arena::{Handle, HandleSet, UniqueArena}; -use crate::valid; +use crate::{valid, FastHashSet}; /// How should code generated by Naga do bounds checks? /// @@ -389,6 +389,61 @@ pub(crate) fn bounds_check_iter<'a>( }) } +/// 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_local_types( + module: &crate::Module, + function: &crate::Function, + info: &valid::FunctionInfo, + policies: BoundsCheckPolicies, +) -> FastHashSet> { + let mut result = FastHashSet::default(); + + if policies.index != BoundsCheckPolicy::ReadZeroSkipWrite { + return result; + } + + 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; 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 bounds_check_iter(arg, module, function, info) + .next() + .is_some() + { + result.insert(base); + } + } + _ => continue, + }; + } + } + } + result +} + impl GuardedIndex { /// Make a `GuardedIndex::Known` from a `GuardedIndex::Expression` if possible. /// diff --git a/naga/src/proc/namer.rs b/naga/src/proc/namer.rs index 6b831cce79..05678242f8 100644 --- a/naga/src/proc/namer.rs +++ b/naga/src/proc/namer.rs @@ -21,9 +21,19 @@ pub enum NameKey { Function(Handle), FunctionArgument(Handle, u32), FunctionLocal(Handle, Handle), + + /// A local variable used by ReadZeroSkipWrite bounds-check policy + /// when it needs to produce a pointer-typed result for an OOB access. + /// These are unique per accessed type, so the second element is a + /// type handle. See docs for [`crate::back::msl`]. + FunctionOobLocal(Handle, Handle), + EntryPoint(EntryPointIndex), EntryPointLocal(EntryPointIndex, Handle), EntryPointArgument(EntryPointIndex, u32), + + /// Entry point version of `FunctionOobLocal`. + EntryPointOobLocal(EntryPointIndex, Handle), } /// This processor assigns names to all the things in a module diff --git a/naga/tests/in/wgsl/pointer-function-arg-restrict.toml b/naga/tests/in/wgsl/pointer-function-arg-restrict.toml new file mode 100644 index 0000000000..bc4112cddf --- /dev/null +++ b/naga/tests/in/wgsl/pointer-function-arg-restrict.toml @@ -0,0 +1,4 @@ +targets = "METAL" + +[bounds_check_policies] +index = "Restrict" diff --git a/naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl b/naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl new file mode 100644 index 0000000000..08ac388354 --- /dev/null +++ b/naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl @@ -0,0 +1,61 @@ +fn takes_ptr(p: ptr) {} +fn takes_array_ptr(p: ptr>) {} +fn takes_vec_ptr(p: ptr>) {} +fn takes_mat_ptr(p: ptr>) {} + +fn local_var(i: u32) { + var arr = array(1, 2, 3, 4); + takes_ptr(&arr[i]); + takes_array_ptr(&arr); + +} + +fn mat_vec_ptrs( + pv: ptr, 4>>, + pm: ptr, 4>>, + i: u32, +) { + takes_vec_ptr(&pv[i]); + takes_mat_ptr(&pm[i]); +} + +fn argument(v: ptr>, i: u32) { + takes_ptr(&v[i]); +} + +fn argument_nested_x2(v: ptr, 4>>, i: u32, j: u32) { + takes_ptr(&v[i][j]); + + // Mixing compile and runtime bounds checks + takes_ptr(&v[i][0]); + takes_ptr(&v[0][j]); + + takes_array_ptr(&v[i]); +} + +fn argument_nested_x3(v: ptr, 4>, 4>>, i: u32, j: u32) { + takes_ptr(&v[i][0][j]); + takes_ptr(&v[i][j][0]); + takes_ptr(&v[0][i][j]); +} + +fn index_from_self(v: ptr>, i: u32) { + takes_ptr(&v[v[i]]); +} + +fn local_var_from_arg(a: array, i: u32) { + var b = a; + takes_ptr(&b[i]); +} + +fn let_binding(a: ptr>, i: u32) { + let p0 = &a[i]; + takes_ptr(p0); + + let p1 = &a[0]; + takes_ptr(p1); +} + +// Runtime-sized arrays can only appear in storage buffers, while (in the base +// language) pointers can only appear in function or private space, so there +// is no interaction to test. diff --git a/naga/tests/in/wgsl/pointer-function-arg-rzsw.toml b/naga/tests/in/wgsl/pointer-function-arg-rzsw.toml new file mode 100644 index 0000000000..44773b7d04 --- /dev/null +++ b/naga/tests/in/wgsl/pointer-function-arg-rzsw.toml @@ -0,0 +1,4 @@ +targets = "METAL" + +[bounds_check_policies] +index = "ReadZeroSkipWrite" diff --git a/naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl b/naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl new file mode 100644 index 0000000000..08ac388354 --- /dev/null +++ b/naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl @@ -0,0 +1,61 @@ +fn takes_ptr(p: ptr) {} +fn takes_array_ptr(p: ptr>) {} +fn takes_vec_ptr(p: ptr>) {} +fn takes_mat_ptr(p: ptr>) {} + +fn local_var(i: u32) { + var arr = array(1, 2, 3, 4); + takes_ptr(&arr[i]); + takes_array_ptr(&arr); + +} + +fn mat_vec_ptrs( + pv: ptr, 4>>, + pm: ptr, 4>>, + i: u32, +) { + takes_vec_ptr(&pv[i]); + takes_mat_ptr(&pm[i]); +} + +fn argument(v: ptr>, i: u32) { + takes_ptr(&v[i]); +} + +fn argument_nested_x2(v: ptr, 4>>, i: u32, j: u32) { + takes_ptr(&v[i][j]); + + // Mixing compile and runtime bounds checks + takes_ptr(&v[i][0]); + takes_ptr(&v[0][j]); + + takes_array_ptr(&v[i]); +} + +fn argument_nested_x3(v: ptr, 4>, 4>>, i: u32, j: u32) { + takes_ptr(&v[i][0][j]); + takes_ptr(&v[i][j][0]); + takes_ptr(&v[0][i][j]); +} + +fn index_from_self(v: ptr>, i: u32) { + takes_ptr(&v[v[i]]); +} + +fn local_var_from_arg(a: array, i: u32) { + var b = a; + takes_ptr(&b[i]); +} + +fn let_binding(a: ptr>, i: u32) { + let p0 = &a[i]; + takes_ptr(p0); + + let p1 = &a[0]; + takes_ptr(p1); +} + +// Runtime-sized arrays can only appear in storage buffers, while (in the base +// language) pointers can only appear in function or private space, so there +// is no interaction to test. diff --git a/naga/tests/in/wgsl/pointer-function-arg.toml b/naga/tests/in/wgsl/pointer-function-arg.toml new file mode 100644 index 0000000000..e74ee3b97b --- /dev/null +++ b/naga/tests/in/wgsl/pointer-function-arg.toml @@ -0,0 +1 @@ +targets = "METAL | GLSL | HLSL | WGSL" diff --git a/naga/tests/in/wgsl/pointer-function-arg.wgsl b/naga/tests/in/wgsl/pointer-function-arg.wgsl new file mode 100644 index 0000000000..606835eef7 --- /dev/null +++ b/naga/tests/in/wgsl/pointer-function-arg.wgsl @@ -0,0 +1,64 @@ +@compute @workgroup_size(1) +fn main() {} + +fn takes_ptr(p: ptr) {} +fn takes_array_ptr(p: ptr>) {} +fn takes_vec_ptr(p: ptr>) {} +fn takes_mat_ptr(p: ptr>) {} + +fn local_var(i: u32) { + var arr = array(1, 2, 3, 4); + takes_ptr(&arr[i]); + takes_array_ptr(&arr); + +} + +fn mat_vec_ptrs( + pv: ptr, 4>>, + pm: ptr, 4>>, + i: u32, +) { + takes_vec_ptr(&pv[i]); + takes_mat_ptr(&pm[i]); +} + +fn argument(v: ptr>, i: u32) { + takes_ptr(&v[i]); +} + +fn argument_nested_x2(v: ptr, 4>>, i: u32, j: u32) { + takes_ptr(&v[i][j]); + + // Mixing compile and runtime bounds checks + takes_ptr(&v[i][0]); + takes_ptr(&v[0][j]); + + takes_array_ptr(&v[i]); +} + +fn argument_nested_x3(v: ptr, 4>, 4>>, i: u32, j: u32) { + takes_ptr(&v[i][0][j]); + takes_ptr(&v[i][j][0]); + takes_ptr(&v[0][i][j]); +} + +fn index_from_self(v: ptr>, i: u32) { + takes_ptr(&v[v[i]]); +} + +fn local_var_from_arg(a: array, i: u32) { + var b = a; + takes_ptr(&b[i]); +} + +fn let_binding(a: ptr>, i: u32) { + let p0 = &a[i]; + takes_ptr(p0); + + let p1 = &a[0]; + takes_ptr(p1); +} + +// Runtime-sized arrays can only appear in storage buffers, while (in the base +// language) pointers can only appear in function or private space, so there +// is no interaction to test. diff --git a/naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl b/naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl new file mode 100644 index 0000000000..a3ffe7952e --- /dev/null +++ b/naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl @@ -0,0 +1,80 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + + +void takes_ptr(inout int p) { + return; +} + +void takes_array_ptr(inout int p_1[4]) { + return; +} + +void takes_vec_ptr(inout ivec2 p_2) { + return; +} + +void takes_mat_ptr(inout mat2x2 p_3) { + return; +} + +void local_var(uint i) { + int arr[4] = int[4](1, 2, 3, 4); + takes_ptr(arr[i]); + takes_array_ptr(arr); + return; +} + +void mat_vec_ptrs(inout ivec2 pv[4], inout mat2x2 pm[4], uint i_1) { + takes_vec_ptr(pv[i_1]); + takes_mat_ptr(pm[i_1]); + return; +} + +void argument(inout int v[4], uint i_2) { + takes_ptr(v[i_2]); + return; +} + +void argument_nested_x2_(inout int v_1[4][4], uint i_3, uint j) { + takes_ptr(v_1[i_3][j]); + takes_ptr(v_1[i_3][0]); + takes_ptr(v_1[0][j]); + takes_array_ptr(v_1[i_3]); + return; +} + +void argument_nested_x3_(inout int v_2[4][4][4], uint i_4, uint j_1) { + takes_ptr(v_2[i_4][0][j_1]); + takes_ptr(v_2[i_4][j_1][0]); + takes_ptr(v_2[0][i_4][j_1]); + return; +} + +void index_from_self(inout int v_3[4], uint i_5) { + int _e3 = v_3[i_5]; + takes_ptr(v_3[_e3]); + return; +} + +void local_var_from_arg(int a[4], uint i_6) { + int b[4] = int[4](0, 0, 0, 0); + b = a; + takes_ptr(b[i_6]); + return; +} + +void let_binding(inout int a_1[4], uint i_7) { + takes_ptr(a_1[i_7]); + takes_ptr(a_1[0]); + return; +} + +void main() { + return; +} + diff --git a/naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl b/naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl new file mode 100644 index 0000000000..1f4e76eeae --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl @@ -0,0 +1,93 @@ +void takes_ptr(inout int p) +{ + return; +} + +void takes_array_ptr(inout int p_1[4]) +{ + return; +} + +void takes_vec_ptr(inout int2 p_2) +{ + return; +} + +void takes_mat_ptr(inout float2x2 p_3) +{ + return; +} + +typedef int ret_Constructarray4_int_[4]; +ret_Constructarray4_int_ Constructarray4_int_(int arg0, int arg1, int arg2, int arg3) { + int ret[4] = { arg0, arg1, arg2, arg3 }; + return ret; +} + +void local_var(uint i) +{ + int arr[4] = Constructarray4_int_(int(1), int(2), int(3), int(4)); + + takes_ptr(arr[min(uint(i), 3u)]); + takes_array_ptr(arr); + return; +} + +void mat_vec_ptrs(inout int2 pv[4], inout float2x2 pm[4], uint i_1) +{ + takes_vec_ptr(pv[min(uint(i_1), 3u)]); + takes_mat_ptr(pm[min(uint(i_1), 3u)]); + return; +} + +void argument(inout int v[4], uint i_2) +{ + takes_ptr(v[min(uint(i_2), 3u)]); + return; +} + +void argument_nested_x2_(inout int v_1[4][4], uint i_3, uint j) +{ + takes_ptr(v_1[min(uint(i_3), 3u)][min(uint(j), 3u)]); + takes_ptr(v_1[min(uint(i_3), 3u)][0]); + takes_ptr(v_1[0][min(uint(j), 3u)]); + takes_array_ptr(v_1[min(uint(i_3), 3u)]); + return; +} + +void argument_nested_x3_(inout int v_2[4][4][4], uint i_4, uint j_1) +{ + takes_ptr(v_2[min(uint(i_4), 3u)][0][min(uint(j_1), 3u)]); + takes_ptr(v_2[min(uint(i_4), 3u)][min(uint(j_1), 3u)][0]); + takes_ptr(v_2[0][min(uint(i_4), 3u)][min(uint(j_1), 3u)]); + return; +} + +void index_from_self(inout int v_3[4], uint i_5) +{ + int _e3 = v_3[min(uint(i_5), 3u)]; + takes_ptr(v_3[min(uint(_e3), 3u)]); + return; +} + +void local_var_from_arg(int a[4], uint i_6) +{ + int b[4] = (int[4])0; + + b = a; + takes_ptr(b[min(uint(i_6), 3u)]); + return; +} + +void let_binding(inout int a_1[4], uint i_7) +{ + takes_ptr(a_1[min(uint(i_7), 3u)]); + takes_ptr(a_1[0]); + return; +} + +[numthreads(1, 1, 1)] +void main() +{ + return; +} diff --git a/naga/tests/out/hlsl/wgsl-pointer-function-arg.ron b/naga/tests/out/hlsl/wgsl-pointer-function-arg.ron new file mode 100644 index 0000000000..a07b03300b --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-pointer-function-arg.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/msl/wgsl-pointer-function-arg-restrict.msl b/naga/tests/out/msl/wgsl-pointer-function-arg-restrict.msl new file mode 100644 index 0000000000..66990660e0 --- /dev/null +++ b/naga/tests/out/msl/wgsl-pointer-function-arg-restrict.msl @@ -0,0 +1,123 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_2 { + int inner[4]; +}; +struct type_9 { + metal::int2 inner[4]; +}; +struct type_11 { + metal::float2x2 inner[4]; +}; +struct type_13 { + type_2 inner[4]; +}; +struct type_15 { + type_13 inner[4]; +}; + +void takes_ptr( + thread int& p +) { + return; +} + +void takes_array_ptr( + thread type_2& p_1 +) { + return; +} + +void takes_vec_ptr( + thread metal::int2& p_2 +) { + return; +} + +void takes_mat_ptr( + thread metal::float2x2& p_3 +) { + return; +} + +void local_var( + uint i +) { + type_2 arr = type_2 {1, 2, 3, 4}; + takes_ptr(arr.inner[metal::min(unsigned(i), 3u)]); + takes_array_ptr(arr); + return; +} + +void mat_vec_ptrs( + thread type_9& pv, + thread type_11& pm, + uint i_1 +) { + takes_vec_ptr(pv.inner[metal::min(unsigned(i_1), 3u)]); + takes_mat_ptr(pm.inner[metal::min(unsigned(i_1), 3u)]); + return; +} + +void argument( + thread type_2& v, + uint i_2 +) { + takes_ptr(v.inner[metal::min(unsigned(i_2), 3u)]); + return; +} + +void argument_nested_x2_( + thread type_13& v_1, + uint i_3, + uint j +) { + takes_ptr(v_1.inner[metal::min(unsigned(i_3), 3u)].inner[metal::min(unsigned(j), 3u)]); + takes_ptr(v_1.inner[metal::min(unsigned(i_3), 3u)].inner[0]); + takes_ptr(v_1.inner[0].inner[metal::min(unsigned(j), 3u)]); + takes_array_ptr(v_1.inner[metal::min(unsigned(i_3), 3u)]); + return; +} + +void argument_nested_x3_( + thread type_15& v_2, + uint i_4, + uint j_1 +) { + takes_ptr(v_2.inner[metal::min(unsigned(i_4), 3u)].inner[0].inner[metal::min(unsigned(j_1), 3u)]); + takes_ptr(v_2.inner[metal::min(unsigned(i_4), 3u)].inner[metal::min(unsigned(j_1), 3u)].inner[0]); + takes_ptr(v_2.inner[0].inner[metal::min(unsigned(i_4), 3u)].inner[metal::min(unsigned(j_1), 3u)]); + return; +} + +void index_from_self( + thread type_2& v_3, + uint i_5 +) { + int _e3 = v_3.inner[metal::min(unsigned(i_5), 3u)]; + takes_ptr(v_3.inner[metal::min(unsigned(_e3), 3u)]); + return; +} + +void local_var_from_arg( + type_2 a, + uint i_6 +) { + type_2 b = {}; + b = a; + takes_ptr(b.inner[metal::min(unsigned(i_6), 3u)]); + return; +} + +void let_binding( + thread type_2& a_1, + uint i_7 +) { + takes_ptr(a_1.inner[metal::min(unsigned(i_7), 3u)]); + takes_ptr(a_1.inner[0]); + return; +} diff --git a/naga/tests/out/msl/wgsl-pointer-function-arg-rzsw.msl b/naga/tests/out/msl/wgsl-pointer-function-arg-rzsw.msl new file mode 100644 index 0000000000..8b76ec2505 --- /dev/null +++ b/naga/tests/out/msl/wgsl-pointer-function-arg-rzsw.msl @@ -0,0 +1,139 @@ +// language: metal1.0 +#include +#include + +using metal::uint; +struct DefaultConstructible { + template + operator T() && { + return T {}; + } +}; + +struct type_2 { + int inner[4]; +}; +struct type_9 { + metal::int2 inner[4]; +}; +struct type_11 { + metal::float2x2 inner[4]; +}; +struct type_13 { + type_2 inner[4]; +}; +struct type_15 { + type_13 inner[4]; +}; + +void takes_ptr( + thread int& p +) { + return; +} + +void takes_array_ptr( + thread type_2& p_1 +) { + return; +} + +void takes_vec_ptr( + thread metal::int2& p_2 +) { + return; +} + +void takes_mat_ptr( + thread metal::float2x2& p_3 +) { + return; +} + +void local_var( + uint i +) { + type_2 arr = type_2 {1, 2, 3, 4}; + int oob = {}; + takes_ptr(uint(i) < 4 ? arr.inner[i] : oob); + takes_array_ptr(arr); + return; +} + +void mat_vec_ptrs( + thread type_9& pv, + thread type_11& pm, + uint i_1 +) { + metal::int2 oob_1 = {}; + metal::float2x2 oob_2 = {}; + takes_vec_ptr(uint(i_1) < 4 ? pv.inner[i_1] : oob_1); + takes_mat_ptr(uint(i_1) < 4 ? pm.inner[i_1] : oob_2); + return; +} + +void argument( + thread type_2& v, + uint i_2 +) { + int oob_3 = {}; + takes_ptr(uint(i_2) < 4 ? v.inner[i_2] : oob_3); + return; +} + +void argument_nested_x2_( + thread type_13& v_1, + uint i_3, + uint j +) { + int oob_4 = {}; + type_2 oob_5 = {}; + takes_ptr(uint(j) < 4 && uint(i_3) < 4 ? v_1.inner[i_3].inner[j] : oob_4); + takes_ptr(uint(i_3) < 4 ? v_1.inner[i_3].inner[0] : oob_4); + takes_ptr(uint(j) < 4 ? v_1.inner[0].inner[j] : oob_4); + takes_array_ptr(uint(i_3) < 4 ? v_1.inner[i_3] : oob_5); + return; +} + +void argument_nested_x3_( + thread type_15& v_2, + uint i_4, + uint j_1 +) { + int oob_6 = {}; + takes_ptr(uint(j_1) < 4 && uint(i_4) < 4 ? v_2.inner[i_4].inner[0].inner[j_1] : oob_6); + takes_ptr(uint(j_1) < 4 && uint(i_4) < 4 ? v_2.inner[i_4].inner[j_1].inner[0] : oob_6); + takes_ptr(uint(j_1) < 4 && uint(i_4) < 4 ? v_2.inner[0].inner[i_4].inner[j_1] : oob_6); + return; +} + +void index_from_self( + thread type_2& v_3, + uint i_5 +) { + int oob_7 = {}; + int _e3 = uint(i_5) < 4 ? v_3.inner[i_5] : DefaultConstructible(); + takes_ptr(uint(_e3) < 4 ? v_3.inner[_e3] : oob_7); + return; +} + +void local_var_from_arg( + type_2 a, + uint i_6 +) { + type_2 b = {}; + int oob_8 = {}; + b = a; + takes_ptr(uint(i_6) < 4 ? b.inner[i_6] : oob_8); + return; +} + +void let_binding( + thread type_2& a_1, + uint i_7 +) { + int oob_9 = {}; + takes_ptr(uint(i_7) < 4 ? a_1.inner[i_7] : oob_9); + takes_ptr(a_1.inner[0]); + return; +} diff --git a/naga/tests/out/msl/wgsl-pointer-function-arg.msl b/naga/tests/out/msl/wgsl-pointer-function-arg.msl new file mode 100644 index 0000000000..c0d1d264fa --- /dev/null +++ b/naga/tests/out/msl/wgsl-pointer-function-arg.msl @@ -0,0 +1,128 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_2 { + int inner[4]; +}; +struct type_9 { + metal::int2 inner[4]; +}; +struct type_11 { + metal::float2x2 inner[4]; +}; +struct type_13 { + type_2 inner[4]; +}; +struct type_15 { + type_13 inner[4]; +}; + +void takes_ptr( + thread int& p +) { + return; +} + +void takes_array_ptr( + thread type_2& p_1 +) { + return; +} + +void takes_vec_ptr( + thread metal::int2& p_2 +) { + return; +} + +void takes_mat_ptr( + thread metal::float2x2& p_3 +) { + return; +} + +void local_var( + uint i +) { + type_2 arr = type_2 {1, 2, 3, 4}; + takes_ptr(arr.inner[i]); + takes_array_ptr(arr); + return; +} + +void mat_vec_ptrs( + thread type_9& pv, + thread type_11& pm, + uint i_1 +) { + takes_vec_ptr(pv.inner[i_1]); + takes_mat_ptr(pm.inner[i_1]); + return; +} + +void argument( + thread type_2& v, + uint i_2 +) { + takes_ptr(v.inner[i_2]); + return; +} + +void argument_nested_x2_( + thread type_13& v_1, + uint i_3, + uint j +) { + takes_ptr(v_1.inner[i_3].inner[j]); + takes_ptr(v_1.inner[i_3].inner[0]); + takes_ptr(v_1.inner[0].inner[j]); + takes_array_ptr(v_1.inner[i_3]); + return; +} + +void argument_nested_x3_( + thread type_15& v_2, + uint i_4, + uint j_1 +) { + takes_ptr(v_2.inner[i_4].inner[0].inner[j_1]); + takes_ptr(v_2.inner[i_4].inner[j_1].inner[0]); + takes_ptr(v_2.inner[0].inner[i_4].inner[j_1]); + return; +} + +void index_from_self( + thread type_2& v_3, + uint i_5 +) { + int _e3 = v_3.inner[i_5]; + takes_ptr(v_3.inner[_e3]); + return; +} + +void local_var_from_arg( + type_2 a, + uint i_6 +) { + type_2 b = {}; + b = a; + takes_ptr(b.inner[i_6]); + return; +} + +void let_binding( + thread type_2& a_1, + uint i_7 +) { + takes_ptr(a_1.inner[i_7]); + takes_ptr(a_1.inner[0]); + return; +} + +kernel void main_( +) { + return; +} diff --git a/naga/tests/out/wgsl/wgsl-pointer-function-arg.wgsl b/naga/tests/out/wgsl/wgsl-pointer-function-arg.wgsl new file mode 100644 index 0000000000..20a2349fcb --- /dev/null +++ b/naga/tests/out/wgsl/wgsl-pointer-function-arg.wgsl @@ -0,0 +1,76 @@ +fn takes_ptr(p: ptr) { + return; +} + +fn takes_array_ptr(p_1: ptr>) { + return; +} + +fn takes_vec_ptr(p_2: ptr>) { + return; +} + +fn takes_mat_ptr(p_3: ptr>) { + return; +} + +fn local_var(i: u32) { + var arr: array = array(1i, 2i, 3i, 4i); + + takes_ptr((&arr[i])); + takes_array_ptr((&arr)); + return; +} + +fn mat_vec_ptrs(pv: ptr, 4>>, pm: ptr, 4>>, i_1: u32) { + takes_vec_ptr((&(*pv)[i_1])); + takes_mat_ptr((&(*pm)[i_1])); + return; +} + +fn argument(v: ptr>, i_2: u32) { + takes_ptr((&(*v)[i_2])); + return; +} + +fn argument_nested_x2_(v_1: ptr, 4>>, i_3: u32, j: u32) { + takes_ptr((&(*v_1)[i_3][j])); + takes_ptr((&(*v_1)[i_3][0])); + takes_ptr((&(*v_1)[0][j])); + takes_array_ptr((&(*v_1)[i_3])); + return; +} + +fn argument_nested_x3_(v_2: ptr, 4>, 4>>, i_4: u32, j_1: u32) { + takes_ptr((&(*v_2)[i_4][0][j_1])); + takes_ptr((&(*v_2)[i_4][j_1][0])); + takes_ptr((&(*v_2)[0][i_4][j_1])); + return; +} + +fn index_from_self(v_3: ptr>, i_5: u32) { + let _e3 = (*v_3)[i_5]; + takes_ptr((&(*v_3)[_e3])); + return; +} + +fn local_var_from_arg(a: array, i_6: u32) { + var b: array; + + b = a; + takes_ptr((&b[i_6])); + return; +} + +fn let_binding(a_1: ptr>, i_7: u32) { + let p0_ = (&(*a_1)[i_7]); + takes_ptr(p0_); + let p1_ = (&(*a_1)[0]); + takes_ptr(p1_); + return; +} + +@compute @workgroup_size(1, 1, 1) +fn main() { + return; +}