From 72831853053089209fb707132cd0390c0eda83f7 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Tue, 1 Oct 2024 12:02:24 -0700 Subject: [PATCH] [naga] Extend snapshot tests for bounds checks. Extend the snapshot tests for bounds checking to cover the case where a runtime-sized array is indexed by a constant. --- naga/tests/in/bounds-check-restrict.wgsl | 8 ++++++++ naga/tests/in/bounds-check-zero-atomic.wgsl | 8 ++++++++ naga/tests/in/bounds-check-zero.wgsl | 8 ++++++++ naga/tests/out/msl/bounds-check-restrict.msl | 17 ++++++++++++++++ .../out/msl/bounds-check-zero-atomic.msl | 16 +++++++++++++++ naga/tests/out/msl/bounds-check-zero.msl | 19 ++++++++++++++++++ .../out/spv/bounds-check-restrict.spvasm | 20 ++++++++++++++++++- naga/tests/out/spv/bounds-check-zero.spvasm | 20 ++++++++++++++++++- 8 files changed, 114 insertions(+), 2 deletions(-) diff --git a/naga/tests/in/bounds-check-restrict.wgsl b/naga/tests/in/bounds-check-restrict.wgsl index 2b7208355..6690d2fea 100644 --- a/naga/tests/in/bounds-check-restrict.wgsl +++ b/naga/tests/in/bounds-check-restrict.wgsl @@ -70,3 +70,11 @@ fn set_in_bounds(v: f32) { globals.v[3] = v; globals.m[2][3] = v; } + +fn index_dynamic_array_constant_index() -> f32 { + return globals.d[1000]; +} + +fn set_dynamic_array_constant_index(v: f32) { + globals.d[1000] = v; +} diff --git a/naga/tests/in/bounds-check-zero-atomic.wgsl b/naga/tests/in/bounds-check-zero-atomic.wgsl index 004f08a0a..ed927f518 100644 --- a/naga/tests/in/bounds-check-zero-atomic.wgsl +++ b/naga/tests/in/bounds-check-zero-atomic.wgsl @@ -36,3 +36,11 @@ fn exchange_atomic_dynamic_sized_array(i: i32) -> u32 { return atomicExchange(&globals.c[i], 1u); } +fn fetch_add_atomic_dynamic_sized_array_static_index() -> u32 { + return atomicAdd(&globals.c[1000], 1u); +} + +fn exchange_atomic_dynamic_sized_array_static_index() -> u32 { + return atomicExchange(&globals.c[1000], 1u); +} + diff --git a/naga/tests/in/bounds-check-zero.wgsl b/naga/tests/in/bounds-check-zero.wgsl index 010f46ec3..514efc489 100644 --- a/naga/tests/in/bounds-check-zero.wgsl +++ b/naga/tests/in/bounds-check-zero.wgsl @@ -70,3 +70,11 @@ fn set_in_bounds(v: f32) { globals.v[3] = v; globals.m[2][3] = v; } + +fn index_dynamic_array_constant_index() -> f32 { + return globals.d[1000]; +} + +fn set_dynamic_array_constant_index(v: f32) { + globals.d[1000] = v; +} diff --git a/naga/tests/out/msl/bounds-check-restrict.msl b/naga/tests/out/msl/bounds-check-restrict.msl index 0d4143653..a8bf37155 100644 --- a/naga/tests/out/msl/bounds-check-restrict.msl +++ b/naga/tests/out/msl/bounds-check-restrict.msl @@ -163,3 +163,20 @@ void set_in_bounds( globals.m[2].w = v_7; return; } + +float index_dynamic_array_constant_index( + device Globals const& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e3 = globals.d[metal::min(unsigned(1000), (_buffer_sizes.size0 - 112 - 4) / 4)]; + return _e3; +} + +void set_dynamic_array_constant_index( + float v_8, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.d[metal::min(unsigned(1000), (_buffer_sizes.size0 - 112 - 4) / 4)] = v_8; + return; +} diff --git a/naga/tests/out/msl/bounds-check-zero-atomic.msl b/naga/tests/out/msl/bounds-check-zero-atomic.msl index 4a2f0b07d..232ae34c6 100644 --- a/naga/tests/out/msl/bounds-check-zero-atomic.msl +++ b/naga/tests/out/msl/bounds-check-zero-atomic.msl @@ -75,3 +75,19 @@ uint exchange_atomic_dynamic_sized_array( uint _e5 = uint(i_3) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_exchange_explicit(&globals.c[i_3], 1u, metal::memory_order_relaxed) : DefaultConstructible(); return _e5; } + +uint fetch_add_atomic_dynamic_sized_array_static_index( + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + uint _e4 = uint(1000) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_fetch_add_explicit(&globals.c[1000], 1u, metal::memory_order_relaxed) : DefaultConstructible(); + return _e4; +} + +uint exchange_atomic_dynamic_sized_array_static_index( + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + uint _e4 = uint(1000) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_exchange_explicit(&globals.c[1000], 1u, metal::memory_order_relaxed) : DefaultConstructible(); + return _e4; +} diff --git a/naga/tests/out/msl/bounds-check-zero.msl b/naga/tests/out/msl/bounds-check-zero.msl index 7bbdd50d1..8269d4bf7 100644 --- a/naga/tests/out/msl/bounds-check-zero.msl +++ b/naga/tests/out/msl/bounds-check-zero.msl @@ -183,3 +183,22 @@ void set_in_bounds( globals.m[2].w = v_7; return; } + +float index_dynamic_array_constant_index( + device Globals const& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e3 = uint(1000) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4 ? globals.d[1000] : DefaultConstructible(); + return _e3; +} + +void set_dynamic_array_constant_index( + float v_8, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + if (uint(1000) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4) { + globals.d[1000] = v_8; + } + return; +} diff --git a/naga/tests/out/spv/bounds-check-restrict.spvasm b/naga/tests/out/spv/bounds-check-restrict.spvasm index c7cf675a1..3987ab521 100644 --- a/naga/tests/out/spv/bounds-check-restrict.spvasm +++ b/naga/tests/out/spv/bounds-check-restrict.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 163 +; Bound: 174 OpCapability Shader OpCapability Linkage OpExtension "SPV_KHR_storage_buffer_storage_class" @@ -52,6 +52,7 @@ OpDecorate %12 Binding 0 %129 = OpTypeFunction %2 %11 %7 %138 = OpTypeFunction %2 %11 %11 %3 %158 = OpTypeFunction %2 %3 +%166 = OpConstant %6 1000 %16 = OpFunction %3 None %17 %15 = OpFunctionParameter %11 %14 = OpLabel @@ -232,4 +233,21 @@ OpStore %161 %156 %162 = OpAccessChain %43 %12 %62 %62 %35 OpStore %162 %156 OpReturn +OpFunctionEnd +%164 = OpFunction %3 None %91 +%163 = OpLabel +OpBranch %165 +%165 = OpLabel +%167 = OpAccessChain %20 %12 %35 %166 +%168 = OpLoad %3 %167 +OpReturnValue %168 +OpFunctionEnd +%171 = OpFunction %2 None %158 +%170 = OpFunctionParameter %3 +%169 = OpLabel +OpBranch %172 +%172 = OpLabel +%173 = OpAccessChain %20 %12 %35 %166 +OpStore %173 %170 +OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/bounds-check-zero.spvasm b/naga/tests/out/spv/bounds-check-zero.spvasm index 2bb81261e..fb34f5d81 100644 --- a/naga/tests/out/spv/bounds-check-zero.spvasm +++ b/naga/tests/out/spv/bounds-check-zero.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 200 +; Bound: 211 OpCapability Shader OpCapability Linkage OpExtension "SPV_KHR_storage_buffer_storage_class" @@ -56,6 +56,7 @@ OpDecorate %12 Binding 0 %159 = OpTypeFunction %2 %11 %7 %170 = OpTypeFunction %2 %11 %11 %3 %195 = OpTypeFunction %2 %3 +%203 = OpConstant %6 1000 %16 = OpFunction %3 None %17 %15 = OpFunctionParameter %11 %14 = OpLabel @@ -308,4 +309,21 @@ OpStore %198 %193 %199 = OpAccessChain %48 %12 %76 %76 %37 OpStore %199 %193 OpReturn +OpFunctionEnd +%201 = OpFunction %3 None %115 +%200 = OpLabel +OpBranch %202 +%202 = OpLabel +%204 = OpAccessChain %20 %12 %37 %203 +%205 = OpLoad %3 %204 +OpReturnValue %205 +OpFunctionEnd +%208 = OpFunction %2 None %195 +%207 = OpFunctionParameter %3 +%206 = OpLabel +OpBranch %209 +%209 = OpLabel +%210 = OpAccessChain %20 %12 %37 %203 +OpStore %210 %207 +OpReturn OpFunctionEnd \ No newline at end of file