nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-25 00:03:29 +00:00
Code Issues Packages Projects Releases Wiki Activity

remove BoundsCheckPolicies.image_store

Browse Source
This commit is contained in:
teoxoy 2024-06-26 10:00:28 +02:00 committed by Jim Blandy
parent 339ecf37da
commit ccd6d2ca48
20 changed files with 244 additions and 429 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
naga-cli/src/bin/naga.rs
View File

@ -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()

37
naga/src/back/msl/writer.rs
View File

@ -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)?;

33
naga/src/back/spv/image.rs
View File

@ -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(())
}

23
naga/src/proc/index.rs
View File

@ -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
}
}

1
naga/tests/in/binding-arrays.param.ron
View File

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

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

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

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

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

1
naga/tests/in/pointers.param.ron
View File

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

1
naga/tests/in/policy-mix.param.ron
View File

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

1
naga/tests/in/resource-binding-map.param.ron
View File

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

12
naga/tests/out/msl/binding-arrays.msl
View File

@ -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));

10
naga/tests/out/msl/bounds-check-image-restrict.msl
View File

@ -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;
}

20
naga/tests/out/msl/bounds-check-image-rzsw.msl
View File

@ -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;
}

78
naga/tests/out/spv/binding-arrays.spvasm
View File

@ -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

211
naga/tests/out/spv/bounds-check-image-restrict.spvasm
View File

@ -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

227
naga/tests/out/spv/bounds-check-image-rzsw.spvasm
View File

@ -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

1
wgpu-hal/src/gles/device.rs
View File

@ -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,
};

1
wgpu-hal/src/metal/device.rs
View File

@ -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,
},

1
wgpu-hal/src/vulkan/adapter.rs
View File

@ -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,
},

2
wgpu-hal/src/vulkan/device.rs
View File

@ -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,
};
}