remove BoundsCheckPolicies.image_store

naga-cli/src/bin/naga.rs

@@ -38,13 +38,6 @@ struct Args {
     #[argh(option)]
     image_load_bounds_check_policy: Option<BoundsCheckPolicyArg>,
 
-    /// what policy to use for texture stores bounds checking.
-    ///
-    /// Possible values are the same as for `index-bounds-check-policy`. If
-    /// omitted, defaults to the index bounds check policy.
-    #[argh(option)]
-    image_store_bounds_check_policy: Option<BoundsCheckPolicyArg>,
-
     /// directory to dump the SPIR-V block context dump to
     #[argh(option)]
     block_ctx_dir: Option<String>,
@@ -409,10 +402,6 @@ fn run() -> anyhow::Result<()> {
         Some(arg) => arg.0,
         None => params.bounds_check_policies.index,
     };
-    params.bounds_check_policies.image_store = match args.image_store_bounds_check_policy {
-        Some(arg) => arg.0,
-        None => params.bounds_check_policies.index,
-    };
     params.overrides = args
         .overrides
         .iter()

naga/src/back/msl/writer.rs

@@ -1063,43 +1063,6 @@ impl<W: Write> Writer<W> {
         address: &TexelAddress,
         value: Handle<crate::Expression>,
         context: &StatementContext,
-    ) -> BackendResult {
-        match context.expression.policies.image_store {
-            proc::BoundsCheckPolicy::Restrict => {
-                // We don't have a restricted level value, because we don't
-                // support writes to mipmapped textures.
-                debug_assert!(address.level.is_none());
-
-                write!(self.out, "{level}")?;
-                self.put_expression(image, &context.expression, false)?;
-                write!(self.out, ".write(")?;
-                self.put_expression(value, &context.expression, true)?;
-                write!(self.out, ", ")?;
-                self.put_restricted_texel_address(image, address, &context.expression)?;
-                writeln!(self.out, ");")?;
-            }
-            proc::BoundsCheckPolicy::ReadZeroSkipWrite => {
-                write!(self.out, "{level}if (")?;
-                self.put_image_access_bounds_check(image, address, &context.expression)?;
-                writeln!(self.out, ") {{")?;
-                self.put_unchecked_image_store(level.next(), image, address, value, context)?;
-                writeln!(self.out, "{level}}}")?;
-            }
-            proc::BoundsCheckPolicy::Unchecked => {
-                self.put_unchecked_image_store(level, image, address, value, context)?;
-            }
-        }
-
-        Ok(())
-    }
-
-    fn put_unchecked_image_store(
-        &mut self,
-        level: back::Level,
-        image: Handle<crate::Expression>,
-        address: &TexelAddress,
-        value: Handle<crate::Expression>,
-        context: &StatementContext,
     ) -> BackendResult {
         write!(self.out, "{level}")?;
         self.put_expression(image, &context.expression, false)?;

naga/src/back/spv/image.rs

@@ -1178,32 +1178,13 @@ impl<'w> BlockContext<'w> {
             _ => {}
         }
 
-        match self.writer.bounds_check_policies.image_store {
-            crate::proc::BoundsCheckPolicy::Restrict => {
-                let (coords, _, _) =
-                    self.write_restricted_coordinates(image_id, coordinates, None, None, block)?;
-                write.generate(&mut self.writer.id_gen, coords, None, None, block);
-            }
-            crate::proc::BoundsCheckPolicy::ReadZeroSkipWrite => {
-                self.write_conditional_image_access(
-                    image_id,
-                    coordinates,
-                    None,
-                    None,
-                    block,
-                    &write,
-                )?;
-            }
-            crate::proc::BoundsCheckPolicy::Unchecked => {
-                write.generate(
-                    &mut self.writer.id_gen,
-                    coordinates.value_id,
-                    None,
-                    None,
-                    block,
-                );
-            }
-        }
+        write.generate(
+            &mut self.writer.id_gen,
+            coordinates.value_id,
+            None,
+            None,
+            block,
+        );
 
         Ok(())
     }

naga/src/proc/index.rs

@@ -112,21 +112,15 @@ pub struct BoundsCheckPolicies {
     /// This controls the behavior of [`ImageLoad`] expressions when a coordinate,
     /// texture array index, level of detail, or multisampled sample number is out of range.
     ///
-    /// [`ImageLoad`]: crate::Expression::ImageLoad
-    #[cfg_attr(feature = "deserialize", serde(default))]
-    pub image_load: BoundsCheckPolicy,
-
-    /// How should the generated code handle image texel stores that are out
-    /// of range?
-    ///
-    /// This controls the behavior of [`ImageStore`] statements when a coordinate,
-    /// texture array index, level of detail, or multisampled sample number is out of range.
-    ///
-    /// This policy should't be needed since all backends should ignore OOB writes.
+    /// There is no corresponding policy for [`ImageStore`] statements. All the
+    /// platforms we support already discard out-of-bounds image stores,
+    /// effectively implementing the "skip write" part of [`ReadZeroSkipWrite`].
     ///
+    /// [`ImageLoad`]: crate::Expression::ImageLoad
     /// [`ImageStore`]: crate::Statement::ImageStore
+    /// [`ReadZeroSkipWrite`]: BoundsCheckPolicy::ReadZeroSkipWrite
     #[cfg_attr(feature = "deserialize", serde(default))]
-    pub image_store: BoundsCheckPolicy,
+    pub image_load: BoundsCheckPolicy,
 
     /// How should the generated code handle binding array indexes that are out of bounds.
     #[cfg_attr(feature = "deserialize", serde(default))]
@@ -173,10 +167,7 @@ impl BoundsCheckPolicies {
 
     /// Return `true` if any of `self`'s policies are `policy`.
     pub fn contains(&self, policy: BoundsCheckPolicy) -> bool {
-        self.index == policy
-            || self.buffer == policy
-            || self.image_load == policy
-            || self.image_store == policy
+        self.index == policy || self.buffer == policy || self.image_load == policy
     }
 }
 

naga/tests/in/binding-arrays.param.ron

@@ -42,6 +42,5 @@
 		index: ReadZeroSkipWrite,
 		buffer: ReadZeroSkipWrite,
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	)
 )

naga/tests/in/bounds-check-image-restrict.param.ron

@@ -1,7 +1,6 @@
 (
 	bounds_check_policies: (
 		image_load: Restrict,
-		image_store: Restrict,
 	),
 	spv: (
 		version: (1, 1),

naga/tests/in/bounds-check-image-rzsw.param.ron

@@ -1,7 +1,6 @@
 (
 	bounds_check_policies: (
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	),
 	spv: (
 		version: (1, 1),

naga/tests/in/pointers.param.ron

@@ -1,7 +1,6 @@
 (
 	bounds_check_policies: (
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	),
 	spv: (
 		version: (1, 2),

naga/tests/in/policy-mix.param.ron

@@ -3,7 +3,6 @@
 		index: Restrict,
 		buffer: Unchecked,
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	),
 	spv: (
 		version: (1, 1),

naga/tests/in/resource-binding-map.param.ron

@@ -49,6 +49,5 @@
 		index: ReadZeroSkipWrite,
 		buffer: ReadZeroSkipWrite,
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	)
 )

naga/tests/out/msl/binding-arrays.msl

@@ -150,17 +150,11 @@ fragment main_Output main_(
     metal::float4 _e278 = v4_;
     v4_ = _e278 + _e277;
     metal::float4 _e282 = v4_;
-    if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[0].get_width(), texture_array_storage[0].get_height()))) {
-        texture_array_storage[0].write(_e282, metal::uint2(pix));
-    }
+    texture_array_storage[0].write(_e282, metal::uint2(pix));
     metal::float4 _e285 = v4_;
-    if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[uniform_index].get_width(), texture_array_storage[uniform_index].get_height()))) {
-        texture_array_storage[uniform_index].write(_e285, metal::uint2(pix));
-    }
+    texture_array_storage[uniform_index].write(_e285, metal::uint2(pix));
     metal::float4 _e288 = v4_;
-    if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[non_uniform_index].get_width(), texture_array_storage[non_uniform_index].get_height()))) {
-        texture_array_storage[non_uniform_index].write(_e288, metal::uint2(pix));
-    }
+    texture_array_storage[non_uniform_index].write(_e288, metal::uint2(pix));
     metal::uint2 _e289 = u2_;
     uint _e290 = u1_;
     metal::float2 v2_ = static_cast<metal::float2>(_e289 + metal::uint2(_e290));

naga/tests/out/msl/bounds-check-image-restrict.msl

@@ -111,7 +111,7 @@ void test_textureStore_1d(
     metal::float4 value,
     metal::texture1d<float, metal::access::write> image_storage_1d
 ) {
-    image_storage_1d.write(value, metal::min(uint(coords_10), image_storage_1d.get_width() - 1));
+    image_storage_1d.write(value, uint(coords_10));
     return;
 }
 
@@ -120,7 +120,7 @@ void test_textureStore_2d(
     metal::float4 value_1,
     metal::texture2d<float, metal::access::write> image_storage_2d
 ) {
-    image_storage_2d.write(value_1, metal::min(metal::uint2(coords_11), metal::uint2(image_storage_2d.get_width(), image_storage_2d.get_height()) - 1));
+    image_storage_2d.write(value_1, metal::uint2(coords_11));
     return;
 }
 
@@ -130,7 +130,7 @@ void test_textureStore_2d_array_u(
     metal::float4 value_2,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    image_storage_2d_array.write(value_2, metal::min(metal::uint2(coords_12), metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()) - 1), metal::min(uint(array_index), image_storage_2d_array.get_array_size() - 1));
+    image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index);
     return;
 }
 
@@ -140,7 +140,7 @@ void test_textureStore_2d_array_s(
     metal::float4 value_3,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    image_storage_2d_array.write(value_3, metal::min(metal::uint2(coords_13), metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()) - 1), metal::min(uint(array_index_1), image_storage_2d_array.get_array_size() - 1));
+    image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1);
     return;
 }
 
@@ -149,7 +149,7 @@ void test_textureStore_3d(
     metal::float4 value_4,
     metal::texture3d<float, metal::access::write> image_storage_3d
 ) {
-    image_storage_3d.write(value_4, metal::min(metal::uint3(coords_14), metal::uint3(image_storage_3d.get_width(), image_storage_3d.get_height(), image_storage_3d.get_depth()) - 1));
+    image_storage_3d.write(value_4, metal::uint3(coords_14));
     return;
 }
 

naga/tests/out/msl/bounds-check-image-rzsw.msl

@@ -110,9 +110,7 @@ void test_textureStore_1d(
     metal::float4 value,
     metal::texture1d<float, metal::access::write> image_storage_1d
 ) {
-    if (uint(coords_10) < image_storage_1d.get_width()) {
-        image_storage_1d.write(value, uint(coords_10));
-    }
+    image_storage_1d.write(value, uint(coords_10));
     return;
 }
 
@@ -121,9 +119,7 @@ void test_textureStore_2d(
     metal::float4 value_1,
     metal::texture2d<float, metal::access::write> image_storage_2d
 ) {
-    if (metal::all(metal::uint2(coords_11) < metal::uint2(image_storage_2d.get_width(), image_storage_2d.get_height()))) {
-        image_storage_2d.write(value_1, metal::uint2(coords_11));
-    }
+    image_storage_2d.write(value_1, metal::uint2(coords_11));
     return;
 }
 
@@ -133,9 +129,7 @@ void test_textureStore_2d_array_u(
     metal::float4 value_2,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    if (uint(array_index) < image_storage_2d_array.get_array_size() && metal::all(metal::uint2(coords_12) < metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()))) {
-        image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index);
-    }
+    image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index);
     return;
 }
 
@@ -145,9 +139,7 @@ void test_textureStore_2d_array_s(
     metal::float4 value_3,
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
-    if (uint(array_index_1) < image_storage_2d_array.get_array_size() && metal::all(metal::uint2(coords_13) < metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()))) {
-        image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1);
-    }
+    image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1);
     return;
 }
 
@@ -156,9 +148,7 @@ void test_textureStore_3d(
     metal::float4 value_4,
     metal::texture3d<float, metal::access::write> image_storage_3d
 ) {
-    if (metal::all(metal::uint3(coords_14) < metal::uint3(image_storage_3d.get_width(), image_storage_3d.get_height(), image_storage_3d.get_depth()))) {
-        image_storage_3d.write(value_4, metal::uint3(coords_14));
-    }
+    image_storage_3d.write(value_4, metal::uint3(coords_14));
     return;
 }