diff --git a/naga-cli/src/bin/naga.rs b/naga-cli/src/bin/naga.rs
index 97d947973..002c6dd66 100644
--- a/naga-cli/src/bin/naga.rs
+++ b/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()
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index ad3dd69eb..48f862f8b 100644
--- a/naga/src/back/msl/writer.rs
+++ b/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)?;
diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs
index 3011ee4d1..769971d13 100644
--- a/naga/src/back/spv/image.rs
+++ b/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(())
     }
diff --git a/naga/src/proc/index.rs b/naga/src/proc/index.rs
index 48b987ce8..555b08d2c 100644
--- a/naga/src/proc/index.rs
+++ b/naga/src/proc/index.rs
@@ -112,22 +112,16 @@ 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.
     ///
+    /// 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_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.
-    ///
-    /// [`ImageStore`]: crate::Statement::ImageStore
-    #[cfg_attr(feature = "deserialize", serde(default))]
-    pub image_store: BoundsCheckPolicy,
-
     /// How should the generated code handle binding array indexes that are out of bounds.
     #[cfg_attr(feature = "deserialize", serde(default))]
     pub binding_array: BoundsCheckPolicy,
@@ -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
     }
 }
 
diff --git a/naga/tests/in/binding-arrays.param.ron b/naga/tests/in/binding-arrays.param.ron
index 39d6c0366..56a498370 100644
--- a/naga/tests/in/binding-arrays.param.ron
+++ b/naga/tests/in/binding-arrays.param.ron
@@ -42,6 +42,5 @@
 		index: ReadZeroSkipWrite,
 		buffer: ReadZeroSkipWrite,
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	)
 )
diff --git a/naga/tests/in/bounds-check-image-restrict.param.ron b/naga/tests/in/bounds-check-image-restrict.param.ron
index d7ff0f006..19f739906 100644
--- a/naga/tests/in/bounds-check-image-restrict.param.ron
+++ b/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),
diff --git a/naga/tests/in/bounds-check-image-rzsw.param.ron b/naga/tests/in/bounds-check-image-rzsw.param.ron
index b256790e1..e818d7a3b 100644
--- a/naga/tests/in/bounds-check-image-rzsw.param.ron
+++ b/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),
diff --git a/naga/tests/in/pointers.param.ron b/naga/tests/in/pointers.param.ron
index fc4027283..c3b4d8880 100644
--- a/naga/tests/in/pointers.param.ron
+++ b/naga/tests/in/pointers.param.ron
@@ -1,7 +1,6 @@
 (
 	bounds_check_policies: (
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	),
 	spv: (
 		version: (1, 2),
diff --git a/naga/tests/in/policy-mix.param.ron b/naga/tests/in/policy-mix.param.ron
index e5469157e..31e80e4c5 100644
--- a/naga/tests/in/policy-mix.param.ron
+++ b/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),
diff --git a/naga/tests/in/resource-binding-map.param.ron b/naga/tests/in/resource-binding-map.param.ron
index 25e7b054b..a700a33f2 100644
--- a/naga/tests/in/resource-binding-map.param.ron
+++ b/naga/tests/in/resource-binding-map.param.ron
@@ -49,6 +49,5 @@
 		index: ReadZeroSkipWrite,
 		buffer: ReadZeroSkipWrite,
 		image_load: ReadZeroSkipWrite,
-		image_store: ReadZeroSkipWrite,
 	)
 )
diff --git a/naga/tests/out/msl/binding-arrays.msl b/naga/tests/out/msl/binding-arrays.msl
index f3548c9e7..75f787a9f 100644
--- a/naga/tests/out/msl/binding-arrays.msl
+++ b/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));
diff --git a/naga/tests/out/msl/bounds-check-image-restrict.msl b/naga/tests/out/msl/bounds-check-image-restrict.msl
index 6a3c43f0c..138c0f645 100644
--- a/naga/tests/out/msl/bounds-check-image-restrict.msl
+++ b/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;
 }
 
diff --git a/naga/tests/out/msl/bounds-check-image-rzsw.msl b/naga/tests/out/msl/bounds-check-image-rzsw.msl
index 5db0c9df9..f73b8e3e3 100644
--- a/naga/tests/out/msl/bounds-check-image-rzsw.msl
+++ b/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;
 }
 
diff --git a/naga/tests/out/spv/binding-arrays.spvasm b/naga/tests/out/spv/binding-arrays.spvasm
index 143ee269a..af75dca49 100644
--- a/naga/tests/out/spv/binding-arrays.spvasm
+++ b/naga/tests/out/spv/binding-arrays.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 428
+; Bound: 413
 OpCapability Shader
 OpCapability ImageQuery
 OpCapability ShaderNonUniform
@@ -77,8 +77,8 @@ OpDecorate %380 NonUniform
 OpDecorate %381 NonUniform
 OpDecorate %382 NonUniform
 OpDecorate %383 NonUniform
-OpDecorate %405 NonUniform
-OpDecorate %406 NonUniform
+OpDecorate %395 NonUniform
+OpDecorate %396 NonUniform
 %2 = OpTypeVoid
 %3 = OpTypeInt 32 0
 %4 = OpTypeStruct %3
@@ -521,54 +521,30 @@ OpStore %72 %387
 %389 = OpAccessChain  %388  %36 %55
 %390 = OpLoad  %16  %389
 %391 = OpLoad  %22  %72
-%392 = OpImageQuerySize  %64  %390
-%393 = OpULessThan  %157  %65 %392
-%394 = OpAll  %150  %393
-OpSelectionMerge %395 None
-OpBranchConditional %394 %396 %395
-%396 = OpLabel
 OpImageWrite %390 %65 %391
-OpBranch %395
-%395 = OpLabel
-%397 = OpAccessChain  %388  %36 %77
-%398 = OpLoad  %16  %397
-%399 = OpLoad  %22  %72
-%400 = OpImageQuerySize  %64  %398
-%401 = OpULessThan  %157  %65 %400
-%402 = OpAll  %150  %401
-OpSelectionMerge %403 None
-OpBranchConditional %402 %404 %403
-%404 = OpLabel
-OpImageWrite %398 %65 %399
-OpBranch %403
-%403 = OpLabel
-%405 = OpAccessChain  %388  %36 %78
-%406 = OpLoad  %16  %405
-%407 = OpLoad  %22  %72
-%408 = OpImageQuerySize  %64  %406
-%409 = OpULessThan  %157  %65 %408
-%410 = OpAll  %150  %409
-OpSelectionMerge %411 None
-OpBranchConditional %410 %412 %411
-%412 = OpLabel
-OpImageWrite %406 %65 %407
-OpBranch %411
-%411 = OpLabel
-%413 = OpLoad  %23  %68
-%414 = OpLoad  %3  %66
-%415 = OpCompositeConstruct  %23  %414 %414
-%416 = OpIAdd  %23  %413 %415
-%417 = OpConvertUToF  %60  %416
-%418 = OpLoad  %22  %72
-%419 = OpCompositeExtract  %6  %417 0
-%420 = OpCompositeExtract  %6  %417 1
-%421 = OpCompositeExtract  %6  %417 0
-%422 = OpCompositeExtract  %6  %417 1
-%423 = OpCompositeConstruct  %22  %419 %420 %421 %422
-%424 = OpFAdd  %22  %418 %423
-%425 = OpLoad  %6  %70
-%426 = OpCompositeConstruct  %22  %425 %425 %425 %425
-%427 = OpFAdd  %22  %424 %426
-OpStore %50 %427
+%392 = OpAccessChain  %388  %36 %77
+%393 = OpLoad  %16  %392
+%394 = OpLoad  %22  %72
+OpImageWrite %393 %65 %394
+%395 = OpAccessChain  %388  %36 %78
+%396 = OpLoad  %16  %395
+%397 = OpLoad  %22  %72
+OpImageWrite %396 %65 %397
+%398 = OpLoad  %23  %68
+%399 = OpLoad  %3  %66
+%400 = OpCompositeConstruct  %23  %399 %399
+%401 = OpIAdd  %23  %398 %400
+%402 = OpConvertUToF  %60  %401
+%403 = OpLoad  %22  %72
+%404 = OpCompositeExtract  %6  %402 0
+%405 = OpCompositeExtract  %6  %402 1
+%406 = OpCompositeExtract  %6  %402 0
+%407 = OpCompositeExtract  %6  %402 1
+%408 = OpCompositeConstruct  %22  %404 %405 %406 %407
+%409 = OpFAdd  %22  %403 %408
+%410 = OpLoad  %6  %70
+%411 = OpCompositeConstruct  %22  %410 %410 %410 %410
+%412 = OpFAdd  %22  %409 %411
+OpStore %50 %412
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/bounds-check-image-restrict.spvasm b/naga/tests/out/spv/bounds-check-image-restrict.spvasm
index 038685a55..7837602e0 100644
--- a/naga/tests/out/spv/bounds-check-image-restrict.spvasm
+++ b/naga/tests/out/spv/bounds-check-image-restrict.spvasm
@@ -1,15 +1,15 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 299
+; Bound: 280
 OpCapability Shader
 OpCapability Sampled1D
 OpCapability Image1D
 OpCapability ImageQuery
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Fragment %269 "fragment_shader" %267
-OpExecutionMode %269 OriginUpperLeft
+OpEntryPoint Fragment %250 "fragment_shader" %248
+OpExecutionMode %250 OriginUpperLeft
 OpName %21 "image_1d"
 OpName %23 "image_2d"
 OpName %25 "image_2d_array"
@@ -59,21 +59,21 @@ OpName %195 "test_textureLoad_depth_multisampled_2d"
 OpName %208 "coords"
 OpName %209 "value"
 OpName %210 "test_textureStore_1d"
-OpName %218 "coords"
-OpName %219 "value"
-OpName %220 "test_textureStore_2d"
-OpName %229 "coords"
-OpName %230 "array_index"
-OpName %231 "value"
-OpName %232 "test_textureStore_2d_array_u"
-OpName %243 "coords"
-OpName %244 "array_index"
-OpName %245 "value"
-OpName %246 "test_textureStore_2d_array_s"
-OpName %256 "coords"
-OpName %257 "value"
-OpName %258 "test_textureStore_3d"
-OpName %269 "fragment_shader"
+OpName %215 "coords"
+OpName %216 "value"
+OpName %217 "test_textureStore_2d"
+OpName %222 "coords"
+OpName %223 "array_index"
+OpName %224 "value"
+OpName %225 "test_textureStore_2d_array_u"
+OpName %232 "coords"
+OpName %233 "array_index"
+OpName %234 "value"
+OpName %235 "test_textureStore_2d_array_s"
+OpName %241 "coords"
+OpName %242 "value"
+OpName %243 "test_textureStore_3d"
+OpName %250 "fragment_shader"
 OpDecorate %21 DescriptorSet 0
 OpDecorate %21 Binding 0
 OpDecorate %23 DescriptorSet 0
@@ -102,7 +102,7 @@ OpDecorate %41 Binding 10
 OpDecorate %43 NonReadable
 OpDecorate %43 DescriptorSet 0
 OpDecorate %43 Binding 11
-OpDecorate %267 Location 0
+OpDecorate %248 Location 0
 %2 = OpTypeVoid
 %4 = OpTypeFloat 32
 %3 = OpTypeImage %4 1D 0 0 0 1 Unknown
@@ -165,24 +165,20 @@ OpDecorate %267 Location 0
 %187 = OpConstantComposite  %12  %53 %53 %53
 %202 = OpConstantComposite  %8  %53 %53
 %211 = OpTypeFunction %2 %5 %6
-%221 = OpTypeFunction %2 %8 %6
-%225 = OpConstantComposite  %8  %53 %53
-%233 = OpTypeFunction %2 %8 %10 %6
-%239 = OpConstantComposite  %12  %53 %53 %53
-%247 = OpTypeFunction %2 %8 %5 %6
-%252 = OpConstantComposite  %12  %53 %53 %53
-%259 = OpTypeFunction %2 %12 %6
-%263 = OpConstantComposite  %12  %53 %53 %53
-%268 = OpTypePointer Output %6
-%267 = OpVariable  %268  Output
-%270 = OpTypeFunction %2
-%280 = OpConstant  %5  0
-%281 = OpConstantNull  %8
-%282 = OpConstant  %10  0
-%283 = OpConstantNull  %12
-%284 = OpConstantNull  %6
-%285 = OpConstant  %4  0.0
-%286 = OpConstantComposite  %6  %285 %285 %285 %285
+%218 = OpTypeFunction %2 %8 %6
+%226 = OpTypeFunction %2 %8 %10 %6
+%236 = OpTypeFunction %2 %8 %5 %6
+%244 = OpTypeFunction %2 %12 %6
+%249 = OpTypePointer Output %6
+%248 = OpVariable  %249  Output
+%251 = OpTypeFunction %2
+%261 = OpConstant  %5  0
+%262 = OpConstantNull  %8
+%263 = OpConstant  %10  0
+%264 = OpConstantNull  %12
+%265 = OpConstantNull  %6
+%266 = OpConstant  %4  0.0
+%267 = OpConstantComposite  %6  %266 %266 %266 %266
 %48 = OpFunction  %6  None %49
 %46 = OpFunctionParameter  %5
 %47 = OpFunctionParameter  %5
@@ -364,93 +360,78 @@ OpFunctionEnd
 %212 = OpLoad  %17  %37
 OpBranch %213
 %213 = OpLabel
-%214 = OpImageQuerySize  %5  %212
-%215 = OpISub  %5  %214 %53
-%216 = OpExtInst  %5  %1 UMin %208 %215
-OpImageWrite %212 %216 %209
+OpImageWrite %212 %208 %209
 OpReturn
 OpFunctionEnd
-%220 = OpFunction  %2  None %221
-%218 = OpFunctionParameter  %8
-%219 = OpFunctionParameter  %6
-%217 = OpLabel
-%222 = OpLoad  %18  %39
-OpBranch %223
-%223 = OpLabel
-%224 = OpImageQuerySize  %8  %222
-%226 = OpISub  %8  %224 %225
-%227 = OpExtInst  %8  %1 UMin %218 %226
-OpImageWrite %222 %227 %219
+%217 = OpFunction  %2  None %218
+%215 = OpFunctionParameter  %8
+%216 = OpFunctionParameter  %6
+%214 = OpLabel
+%219 = OpLoad  %18  %39
+OpBranch %220
+%220 = OpLabel
+OpImageWrite %219 %215 %216
 OpReturn
 OpFunctionEnd
-%232 = OpFunction  %2  None %233
-%229 = OpFunctionParameter  %8
-%230 = OpFunctionParameter  %10
-%231 = OpFunctionParameter  %6
+%225 = OpFunction  %2  None %226
+%222 = OpFunctionParameter  %8
+%223 = OpFunctionParameter  %10
+%224 = OpFunctionParameter  %6
+%221 = OpLabel
+%227 = OpLoad  %19  %41
+OpBranch %228
 %228 = OpLabel
-%234 = OpLoad  %19  %41
-OpBranch %235
-%235 = OpLabel
-%236 = OpBitcast  %5  %230
-%237 = OpCompositeConstruct  %12  %229 %236
-%238 = OpImageQuerySize  %12  %234
-%240 = OpISub  %12  %238 %239
-%241 = OpExtInst  %12  %1 UMin %237 %240
-OpImageWrite %234 %241 %231
+%229 = OpBitcast  %5  %223
+%230 = OpCompositeConstruct  %12  %222 %229
+OpImageWrite %227 %230 %224
 OpReturn
 OpFunctionEnd
-%246 = OpFunction  %2  None %247
-%243 = OpFunctionParameter  %8
-%244 = OpFunctionParameter  %5
-%245 = OpFunctionParameter  %6
-%242 = OpLabel
-%248 = OpLoad  %19  %41
-OpBranch %249
-%249 = OpLabel
-%250 = OpCompositeConstruct  %12  %243 %244
-%251 = OpImageQuerySize  %12  %248
-%253 = OpISub  %12  %251 %252
-%254 = OpExtInst  %12  %1 UMin %250 %253
-OpImageWrite %248 %254 %245
+%235 = OpFunction  %2  None %236
+%232 = OpFunctionParameter  %8
+%233 = OpFunctionParameter  %5
+%234 = OpFunctionParameter  %6
+%231 = OpLabel
+%237 = OpLoad  %19  %41
+OpBranch %238
+%238 = OpLabel
+%239 = OpCompositeConstruct  %12  %232 %233
+OpImageWrite %237 %239 %234
 OpReturn
 OpFunctionEnd
-%258 = OpFunction  %2  None %259
-%256 = OpFunctionParameter  %12
-%257 = OpFunctionParameter  %6
-%255 = OpLabel
+%243 = OpFunction  %2  None %244
+%241 = OpFunctionParameter  %12
+%242 = OpFunctionParameter  %6
+%240 = OpLabel
+%245 = OpLoad  %20  %43
+OpBranch %246
+%246 = OpLabel
+OpImageWrite %245 %241 %242
+OpReturn
+OpFunctionEnd
+%250 = OpFunction  %2  None %251
+%247 = OpLabel
+%252 = OpLoad  %3  %21
+%253 = OpLoad  %7  %23
+%254 = OpLoad  %9  %25
+%255 = OpLoad  %11  %27
+%256 = OpLoad  %13  %29
+%257 = OpLoad  %17  %37
+%258 = OpLoad  %18  %39
+%259 = OpLoad  %19  %41
 %260 = OpLoad  %20  %43
-OpBranch %261
-%261 = OpLabel
-%262 = OpImageQuerySize  %12  %260
-%264 = OpISub  %12  %262 %263
-%265 = OpExtInst  %12  %1 UMin %256 %264
-OpImageWrite %260 %265 %257
-OpReturn
-OpFunctionEnd
-%269 = OpFunction  %2  None %270
-%266 = OpLabel
-%271 = OpLoad  %3  %21
-%272 = OpLoad  %7  %23
-%273 = OpLoad  %9  %25
-%274 = OpLoad  %11  %27
-%275 = OpLoad  %13  %29
-%276 = OpLoad  %17  %37
-%277 = OpLoad  %18  %39
-%278 = OpLoad  %19  %41
-%279 = OpLoad  %20  %43
-OpBranch %287
-%287 = OpLabel
-%288 = OpFunctionCall  %6  %48 %280 %280
-%289 = OpFunctionCall  %6  %63 %281 %280
-%290 = OpFunctionCall  %6  %79 %281 %282 %280
-%291 = OpFunctionCall  %6  %97 %281 %280 %280
-%292 = OpFunctionCall  %6  %113 %283 %280
-%293 = OpFunctionCall  %6  %128 %281 %280
-%294 = OpFunctionCall  %2  %210 %280 %284
-%295 = OpFunctionCall  %2  %220 %281 %284
-%296 = OpFunctionCall  %2  %232 %281 %282 %284
-%297 = OpFunctionCall  %2  %246 %281 %280 %284
-%298 = OpFunctionCall  %2  %258 %283 %284
-OpStore %267 %286
+OpBranch %268
+%268 = OpLabel
+%269 = OpFunctionCall  %6  %48 %261 %261
+%270 = OpFunctionCall  %6  %63 %262 %261
+%271 = OpFunctionCall  %6  %79 %262 %263 %261
+%272 = OpFunctionCall  %6  %97 %262 %261 %261
+%273 = OpFunctionCall  %6  %113 %264 %261
+%274 = OpFunctionCall  %6  %128 %262 %261
+%275 = OpFunctionCall  %2  %210 %261 %265
+%276 = OpFunctionCall  %2  %217 %262 %265
+%277 = OpFunctionCall  %2  %225 %262 %263 %265
+%278 = OpFunctionCall  %2  %235 %262 %261 %265
+%279 = OpFunctionCall  %2  %243 %264 %265
+OpStore %248 %267
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/bounds-check-image-rzsw.spvasm b/naga/tests/out/spv/bounds-check-image-rzsw.spvasm
index a9eeb4204..9b8c091bb 100644
--- a/naga/tests/out/spv/bounds-check-image-rzsw.spvasm
+++ b/naga/tests/out/spv/bounds-check-image-rzsw.spvasm
@@ -1,15 +1,15 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 326
+; Bound: 302
 OpCapability Shader
 OpCapability Sampled1D
 OpCapability Image1D
 OpCapability ImageQuery
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Fragment %297 "fragment_shader" %295
-OpExecutionMode %297 OriginUpperLeft
+OpEntryPoint Fragment %273 "fragment_shader" %271
+OpExecutionMode %273 OriginUpperLeft
 OpName %21 "image_1d"
 OpName %23 "image_2d"
 OpName %25 "image_2d_array"
@@ -59,21 +59,21 @@ OpName %216 "test_textureLoad_depth_multisampled_2d"
 OpName %231 "coords"
 OpName %232 "value"
 OpName %233 "test_textureStore_1d"
-OpName %242 "coords"
-OpName %243 "value"
-OpName %244 "test_textureStore_2d"
-OpName %254 "coords"
-OpName %255 "array_index"
-OpName %256 "value"
-OpName %257 "test_textureStore_2d_array_u"
-OpName %269 "coords"
-OpName %270 "array_index"
-OpName %271 "value"
-OpName %272 "test_textureStore_2d_array_s"
-OpName %283 "coords"
-OpName %284 "value"
-OpName %285 "test_textureStore_3d"
-OpName %297 "fragment_shader"
+OpName %238 "coords"
+OpName %239 "value"
+OpName %240 "test_textureStore_2d"
+OpName %245 "coords"
+OpName %246 "array_index"
+OpName %247 "value"
+OpName %248 "test_textureStore_2d_array_u"
+OpName %255 "coords"
+OpName %256 "array_index"
+OpName %257 "value"
+OpName %258 "test_textureStore_2d_array_s"
+OpName %264 "coords"
+OpName %265 "value"
+OpName %266 "test_textureStore_3d"
+OpName %273 "fragment_shader"
 OpDecorate %21 DescriptorSet 0
 OpDecorate %21 Binding 0
 OpDecorate %23 DescriptorSet 0
@@ -102,7 +102,7 @@ OpDecorate %41 Binding 10
 OpDecorate %43 NonReadable
 OpDecorate %43 DescriptorSet 0
 OpDecorate %43 Binding 11
-OpDecorate %295 Location 0
+OpDecorate %271 Location 0
 %2 = OpTypeVoid
 %4 = OpTypeFloat 32
 %3 = OpTypeImage %4 1D 0 0 0 1 Unknown
@@ -159,19 +159,19 @@ OpDecorate %295 Location 0
 %177 = OpTypeFunction %4 %8 %10 %5
 %198 = OpTypeFunction %4 %8 %5 %5
 %234 = OpTypeFunction %2 %5 %6
-%245 = OpTypeFunction %2 %8 %6
-%258 = OpTypeFunction %2 %8 %10 %6
-%273 = OpTypeFunction %2 %8 %5 %6
-%286 = OpTypeFunction %2 %12 %6
-%296 = OpTypePointer Output %6
-%295 = OpVariable  %296  Output
-%298 = OpTypeFunction %2
-%308 = OpConstant  %5  0
-%309 = OpConstantNull  %8
-%310 = OpConstant  %10  0
-%311 = OpConstantNull  %12
-%312 = OpConstant  %4  0.0
-%313 = OpConstantComposite  %6  %312 %312 %312 %312
+%241 = OpTypeFunction %2 %8 %6
+%249 = OpTypeFunction %2 %8 %10 %6
+%259 = OpTypeFunction %2 %8 %5 %6
+%267 = OpTypeFunction %2 %12 %6
+%272 = OpTypePointer Output %6
+%271 = OpVariable  %272  Output
+%274 = OpTypeFunction %2
+%284 = OpConstant  %5  0
+%285 = OpConstantNull  %8
+%286 = OpConstant  %10  0
+%287 = OpConstantNull  %12
+%288 = OpConstant  %4  0.0
+%289 = OpConstantComposite  %6  %288 %288 %288 %288
 %48 = OpFunction  %6  None %49
 %46 = OpFunctionParameter  %5
 %47 = OpFunctionParameter  %5
@@ -422,117 +422,78 @@ OpFunctionEnd
 %235 = OpLoad  %17  %37
 OpBranch %236
 %236 = OpLabel
-%237 = OpImageQuerySize  %5  %235
-%238 = OpULessThan  %52  %231 %237
-OpSelectionMerge %239 None
-OpBranchConditional %238 %240 %239
-%240 = OpLabel
 OpImageWrite %235 %231 %232
-OpBranch %239
-%239 = OpLabel
 OpReturn
 OpFunctionEnd
-%244 = OpFunction  %2  None %245
-%242 = OpFunctionParameter  %8
-%243 = OpFunctionParameter  %6
-%241 = OpLabel
-%246 = OpLoad  %18  %39
-OpBranch %247
-%247 = OpLabel
-%248 = OpImageQuerySize  %8  %246
-%249 = OpULessThan  %75  %242 %248
-%250 = OpAll  %52  %249
-OpSelectionMerge %251 None
-OpBranchConditional %250 %252 %251
-%252 = OpLabel
-OpImageWrite %246 %242 %243
+%240 = OpFunction  %2  None %241
+%238 = OpFunctionParameter  %8
+%239 = OpFunctionParameter  %6
+%237 = OpLabel
+%242 = OpLoad  %18  %39
+OpBranch %243
+%243 = OpLabel
+OpImageWrite %242 %238 %239
+OpReturn
+OpFunctionEnd
+%248 = OpFunction  %2  None %249
+%245 = OpFunctionParameter  %8
+%246 = OpFunctionParameter  %10
+%247 = OpFunctionParameter  %6
+%244 = OpLabel
+%250 = OpLoad  %19  %41
 OpBranch %251
 %251 = OpLabel
+%252 = OpBitcast  %5  %246
+%253 = OpCompositeConstruct  %12  %245 %252
+OpImageWrite %250 %253 %247
 OpReturn
 OpFunctionEnd
-%257 = OpFunction  %2  None %258
-%254 = OpFunctionParameter  %8
-%255 = OpFunctionParameter  %10
-%256 = OpFunctionParameter  %6
-%253 = OpLabel
-%259 = OpLoad  %19  %41
-OpBranch %260
-%260 = OpLabel
-%261 = OpBitcast  %5  %255
-%262 = OpCompositeConstruct  %12  %254 %261
-%263 = OpImageQuerySize  %12  %259
-%264 = OpULessThan  %96  %262 %263
-%265 = OpAll  %52  %264
-OpSelectionMerge %266 None
-OpBranchConditional %265 %267 %266
-%267 = OpLabel
-OpImageWrite %259 %262 %256
-OpBranch %266
-%266 = OpLabel
+%258 = OpFunction  %2  None %259
+%255 = OpFunctionParameter  %8
+%256 = OpFunctionParameter  %5
+%257 = OpFunctionParameter  %6
+%254 = OpLabel
+%260 = OpLoad  %19  %41
+OpBranch %261
+%261 = OpLabel
+%262 = OpCompositeConstruct  %12  %255 %256
+OpImageWrite %260 %262 %257
 OpReturn
 OpFunctionEnd
-%272 = OpFunction  %2  None %273
-%269 = OpFunctionParameter  %8
-%270 = OpFunctionParameter  %5
-%271 = OpFunctionParameter  %6
-%268 = OpLabel
-%274 = OpLoad  %19  %41
-OpBranch %275
-%275 = OpLabel
-%276 = OpCompositeConstruct  %12  %269 %270
-%277 = OpImageQuerySize  %12  %274
-%278 = OpULessThan  %96  %276 %277
-%279 = OpAll  %52  %278
-OpSelectionMerge %280 None
-OpBranchConditional %279 %281 %280
-%281 = OpLabel
-OpImageWrite %274 %276 %271
-OpBranch %280
-%280 = OpLabel
+%266 = OpFunction  %2  None %267
+%264 = OpFunctionParameter  %12
+%265 = OpFunctionParameter  %6
+%263 = OpLabel
+%268 = OpLoad  %20  %43
+OpBranch %269
+%269 = OpLabel
+OpImageWrite %268 %264 %265
 OpReturn
 OpFunctionEnd
-%285 = OpFunction  %2  None %286
-%283 = OpFunctionParameter  %12
-%284 = OpFunctionParameter  %6
-%282 = OpLabel
-%287 = OpLoad  %20  %43
-OpBranch %288
-%288 = OpLabel
-%289 = OpImageQuerySize  %12  %287
-%290 = OpULessThan  %96  %283 %289
-%291 = OpAll  %52  %290
-OpSelectionMerge %292 None
-OpBranchConditional %291 %293 %292
-%293 = OpLabel
-OpImageWrite %287 %283 %284
-OpBranch %292
-%292 = OpLabel
-OpReturn
-OpFunctionEnd
-%297 = OpFunction  %2  None %298
-%294 = OpLabel
-%299 = OpLoad  %3  %21
-%300 = OpLoad  %7  %23
-%301 = OpLoad  %9  %25
-%302 = OpLoad  %11  %27
-%303 = OpLoad  %13  %29
-%304 = OpLoad  %17  %37
-%305 = OpLoad  %18  %39
-%306 = OpLoad  %19  %41
-%307 = OpLoad  %20  %43
-OpBranch %314
-%314 = OpLabel
-%315 = OpFunctionCall  %6  %48 %308 %308
-%316 = OpFunctionCall  %6  %66 %309 %308
-%317 = OpFunctionCall  %6  %85 %309 %310 %308
-%318 = OpFunctionCall  %6  %106 %309 %308 %308
-%319 = OpFunctionCall  %6  %124 %311 %308
-%320 = OpFunctionCall  %6  %141 %309 %308
-%321 = OpFunctionCall  %2  %233 %308 %53
-%322 = OpFunctionCall  %2  %244 %309 %53
-%323 = OpFunctionCall  %2  %257 %309 %310 %53
-%324 = OpFunctionCall  %2  %272 %309 %308 %53
-%325 = OpFunctionCall  %2  %285 %311 %53
-OpStore %295 %313
+%273 = OpFunction  %2  None %274
+%270 = OpLabel
+%275 = OpLoad  %3  %21
+%276 = OpLoad  %7  %23
+%277 = OpLoad  %9  %25
+%278 = OpLoad  %11  %27
+%279 = OpLoad  %13  %29
+%280 = OpLoad  %17  %37
+%281 = OpLoad  %18  %39
+%282 = OpLoad  %19  %41
+%283 = OpLoad  %20  %43
+OpBranch %290
+%290 = OpLabel
+%291 = OpFunctionCall  %6  %48 %284 %284
+%292 = OpFunctionCall  %6  %66 %285 %284
+%293 = OpFunctionCall  %6  %85 %285 %286 %284
+%294 = OpFunctionCall  %6  %106 %285 %284 %284
+%295 = OpFunctionCall  %6  %124 %287 %284
+%296 = OpFunctionCall  %6  %141 %285 %284
+%297 = OpFunctionCall  %2  %233 %284 %53
+%298 = OpFunctionCall  %2  %240 %285 %53
+%299 = OpFunctionCall  %2  %248 %285 %286 %53
+%300 = OpFunctionCall  %2  %258 %285 %284 %53
+%301 = OpFunctionCall  %2  %266 %287 %53
+OpStore %271 %289
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs
index 4f187709a..0f8c381b5 100644
--- a/wgpu-hal/src/gles/device.rs
+++ b/wgpu-hal/src/gles/device.rs
@@ -246,7 +246,6 @@ impl super::Device {
             index: BoundsCheckPolicy::Unchecked,
             buffer: BoundsCheckPolicy::Unchecked,
             image_load: image_check,
-            image_store: BoundsCheckPolicy::Unchecked,
             binding_array: BoundsCheckPolicy::Unchecked,
         };
 
diff --git a/wgpu-hal/src/metal/device.rs b/wgpu-hal/src/metal/device.rs
index e108d3820..18b9c2dba 100644
--- a/wgpu-hal/src/metal/device.rs
+++ b/wgpu-hal/src/metal/device.rs
@@ -146,7 +146,6 @@ impl super::Device {
                 index: bounds_check_policy,
                 buffer: bounds_check_policy,
                 image_load: bounds_check_policy,
-                image_store: naga::proc::BoundsCheckPolicy::Unchecked,
                 // TODO: support bounds checks on binding arrays
                 binding_array: naga::proc::BoundsCheckPolicy::Unchecked,
             },
diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs
index 1a89aa807..215c0dd95 100644
--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@@ -1773,7 +1773,6 @@ impl super::Adapter {
                     } else {
                         naga::proc::BoundsCheckPolicy::Restrict
                     },
-                    image_store: naga::proc::BoundsCheckPolicy::Unchecked,
                     // TODO: support bounds checks on binding arrays
                     binding_array: naga::proc::BoundsCheckPolicy::Unchecked,
                 },
diff --git a/wgpu-hal/src/vulkan/device.rs b/wgpu-hal/src/vulkan/device.rs
index 00f6c7a41..2f2e045fd 100644
--- a/wgpu-hal/src/vulkan/device.rs
+++ b/wgpu-hal/src/vulkan/device.rs
@@ -736,7 +736,6 @@ impl super::Device {
                             index: naga::proc::BoundsCheckPolicy::Unchecked,
                             buffer: naga::proc::BoundsCheckPolicy::Unchecked,
                             image_load: naga::proc::BoundsCheckPolicy::Unchecked,
-                            image_store: naga::proc::BoundsCheckPolicy::Unchecked,
                             binding_array: naga::proc::BoundsCheckPolicy::Unchecked,
                         };
                     }
@@ -1678,7 +1677,6 @@ impl crate::Device for super::Device {
                         index: naga::proc::BoundsCheckPolicy::Unchecked,
                         buffer: naga::proc::BoundsCheckPolicy::Unchecked,
                         image_load: naga::proc::BoundsCheckPolicy::Unchecked,
-                        image_store: naga::proc::BoundsCheckPolicy::Unchecked,
                         binding_array: naga::proc::BoundsCheckPolicy::Unchecked,
                     };
                 }