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

[naga] Extend snapshot tests for bounds checks.

Browse Source 
Extend the snapshot tests for bounds checking to cover the case where
a runtime-sized array is indexed by a constant.
This commit is contained in:
Jim Blandy 2024-10-01 12:02:24 -07:00
parent 5a6b749335
commit 7283185305
8 changed files with 114 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
naga/tests/in/bounds-check-restrict.wgsl
View File

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

8
naga/tests/in/bounds-check-zero-atomic.wgsl
View File

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

8
naga/tests/in/bounds-check-zero.wgsl
View File

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

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

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

16
naga/tests/out/msl/bounds-check-zero-atomic.msl
View File

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

19
naga/tests/out/msl/bounds-check-zero.msl
View File

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

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

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

20
naga/tests/out/spv/bounds-check-zero.spvasm
View File

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