Skip to content

Commit d11ceba

Browse files
dj2dj2
authored
Add matrix overflow test cases. (#3700)
This CL tests that const-eval of matric operations cause over/under flow. Fixes: #3682
1 parent 563c6d7 commit d11ceba

File tree

3 files changed

+722
-8
lines changed

3 files changed

+722
-8
lines changed

‎src/webgpu/listing_meta.json

+29-8
Original file line numberDiff line numberDiff line change
@@ -2339,18 +2339,39 @@
23392339
"webgpu:shader,validation,expression,call,builtin,unpack4xU8:unsupported:*": { "subcaseMS": 346.801 },
23402340
"webgpu:shader,validation,expression,call,builtin,workgroupUniformLoad:no_atomics:*": { "subcaseMS": 1.000 },
23412341
"webgpu:shader,validation,expression,call,builtin,workgroupUniformLoad:only_in_compute:*": { "subcaseMS": 1.000 },
2342-
"webgpu:shader,validation,expression,matrix,add_sub:invalid:*": { "subcaseMS": 60923.388 },
2343-
"webgpu:shader,validation,expression,matrix,add_sub:with_abstract:*": { "subcaseMS": 17.617 },
2342+
"webgpu:shader,validation,expression,matrix,add_sub:invalid:*": { "subcaseMS": 59542.637 },
2343+
"webgpu:shader,validation,expression,matrix,add_sub:overflow_abstract:*": { "subcaseMS": 23.005 },
2344+
"webgpu:shader,validation,expression,matrix,add_sub:overflow_f16:*": { "subcaseMS": 8.592 },
2345+
"webgpu:shader,validation,expression,matrix,add_sub:overflow_f32:*": { "subcaseMS": 9.352 },
2346+
"webgpu:shader,validation,expression,matrix,add_sub:underflow_abstract:*": { "subcaseMS": 18.032 },
2347+
"webgpu:shader,validation,expression,matrix,add_sub:underflow_f16:*": { "subcaseMS": 5.255 },
2348+
"webgpu:shader,validation,expression,matrix,add_sub:underflow_f32:*": { "subcaseMS": 27.796 },
2349+
"webgpu:shader,validation,expression,matrix,add_sub:with_abstract:*": { "subcaseMS": 4.561 },
23442350
"webgpu:shader,validation,expression,matrix,and_or_xor:invalid:*": { "subcaseMS": 62661.007 },
23452351
"webgpu:shader,validation,expression,matrix,bitwise_shift:invalid:*": { "subcaseMS": 54820.513 },
23462352
"webgpu:shader,validation,expression,matrix,comparison:invalid:*": { "subcaseMS": 60810.479 },
23472353
"webgpu:shader,validation,expression,matrix,div_rem:invalid:*": { "subcaseMS": 62895.929 },
2348-
"webgpu:shader,validation,expression,matrix,mul:f16_and_f32_matrix:*": { "subcaseMS": 59193.817 },
2349-
"webgpu:shader,validation,expression,matrix,mul:f32_and_f16_matrix:*": { "subcaseMS": 1.089 },
2350-
"webgpu:shader,validation,expression,matrix,mul:invalid:*": { "subcaseMS": 58123.532 },
2351-
"webgpu:shader,validation,expression,matrix,mul:mat_by_mat:*": { "subcaseMS": 2928.744 },
2352-
"webgpu:shader,validation,expression,matrix,mul:mat_by_vec:*": { "subcaseMS": 3089.158 },
2353-
"webgpu:shader,validation,expression,matrix,mul:vec_by_mat:*": { "subcaseMS": 3271.763 },
2354+
"webgpu:shader,validation,expression,matrix,mul:f16_and_f32_matrix:*": { "subcaseMS": 0.566 },
2355+
"webgpu:shader,validation,expression,matrix,mul:f32_and_f16_matrix:*": { "subcaseMS": 0.565 },
2356+
"webgpu:shader,validation,expression,matrix,mul:invalid:*": { "subcaseMS": 63267.850 },
2357+
"webgpu:shader,validation,expression,matrix,mul:mat_by_mat:*": { "subcaseMS": 101.751 },
2358+
"webgpu:shader,validation,expression,matrix,mul:mat_by_vec:*": { "subcaseMS": 30.016 },
2359+
"webgpu:shader,validation,expression,matrix,mul:overflow_mat_abstract:*": { "subcaseMS": 28.320 },
2360+
"webgpu:shader,validation,expression,matrix,mul:overflow_mat_abstract_internal:*": { "subcaseMS": 15.850 },
2361+
"webgpu:shader,validation,expression,matrix,mul:overflow_mat_f16:*": { "subcaseMS": 4.597 },
2362+
"webgpu:shader,validation,expression,matrix,mul:overflow_mat_f16_internal:*": { "subcaseMS": 4.099 },
2363+
"webgpu:shader,validation,expression,matrix,mul:overflow_mat_f32:*": { "subcaseMS": 14.859 },
2364+
"webgpu:shader,validation,expression,matrix,mul:overflow_mat_f32_internal:*": { "subcaseMS": 12.024 },
2365+
"webgpu:shader,validation,expression,matrix,mul:overflow_scalar_abstract:*": { "subcaseMS": 20.669 },
2366+
"webgpu:shader,validation,expression,matrix,mul:overflow_scalar_f16:*": { "subcaseMS": 5.298 },
2367+
"webgpu:shader,validation,expression,matrix,mul:overflow_scalar_f32:*": { "subcaseMS": 8.174 },
2368+
"webgpu:shader,validation,expression,matrix,mul:overflow_vec_abstract:*": { "subcaseMS": 13.956 },
2369+
"webgpu:shader,validation,expression,matrix,mul:overflow_vec_abstract_internal:*": { "subcaseMS": 37.976 },
2370+
"webgpu:shader,validation,expression,matrix,mul:overflow_vec_f16:*": { "subcaseMS": 2.920 },
2371+
"webgpu:shader,validation,expression,matrix,mul:overflow_vec_f16_internal:*": { "subcaseMS": 4.689 },
2372+
"webgpu:shader,validation,expression,matrix,mul:overflow_vec_f32:*": { "subcaseMS": 10.747 },
2373+
"webgpu:shader,validation,expression,matrix,mul:overflow_vec_f32_internal:*": { "subcaseMS": 8.354 },
2374+
"webgpu:shader,validation,expression,matrix,mul:vec_by_mat:*": { "subcaseMS": 30.646 },
23542375
"webgpu:shader,validation,expression,overload_resolution:implicit_conversions:*": { "subcaseMS": 519.260 },
23552376
"webgpu:shader,validation,expression,overload_resolution:overload_resolution:*": { "subcaseMS": 436.603 },
23562377
"webgpu:shader,validation,expression,precedence:binary_requires_parentheses:*": { "subcaseMS": 149.921 },

‎src/webgpu/shader/validation/expression/matrix/add_sub.spec.ts

+189
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@ Validation tests for matrix addition and subtraction expressions.
44

55
import { makeTestGroup } from '../../../../../common/framework/test_group.js';
66
import { keysOf } from '../../../../../common/util/data_tables.js';
7+
import { kValue } from '../../../../util/constants.js';
78
import { ShaderValidationTest } from '../../shader_validation_test.js';
89

910
export const g = makeTestGroup(ShaderValidationTest);
@@ -163,3 +164,191 @@ fn main() {
163164

164165
t.expectCompileResult(true, code);
165166
});
167+
168+
g.test('overflow_f32')
169+
.desc(`Validates that f32 add overflows in shader creation`)
170+
.params(u =>
171+
u
172+
.combine('rhs', [kValue.f32.positive.max, 1])
173+
.combine('c', [2, 3, 4] as const)
174+
.combine('r', [2, 3, 4] as const)
175+
)
176+
.fn(t => {
177+
let lhs = `mat${t.params.c}x${t.params.r}f(`;
178+
let rhs = `mat${t.params.c}x${t.params.r}f(`;
179+
for (let i = 0; i < t.params.c; i++) {
180+
for (let k = 0; k < t.params.r; k++) {
181+
lhs += `${kValue.f32.positive.max / 2},`;
182+
rhs += `${t.params.rhs},`;
183+
}
184+
}
185+
rhs += ')';
186+
lhs += ')';
187+
188+
const code = `
189+
@compute @workgroup_size(1)
190+
fn main() {
191+
const foo = ${lhs} + ${rhs};
192+
}
193+
`;
194+
195+
t.expectCompileResult(t.params.rhs === 1, code);
196+
});
197+
198+
g.test('underflow_f32')
199+
.desc(`Validates that f32 add underflows in shader creation`)
200+
.params(u =>
201+
u
202+
.combine('rhs', [kValue.f32.positive.max, 1])
203+
.combine('c', [2, 3, 4] as const)
204+
.combine('r', [2, 3, 4] as const)
205+
)
206+
.fn(t => {
207+
let lhs = `mat${t.params.c}x${t.params.r}f(`;
208+
let rhs = `mat${t.params.c}x${t.params.r}f(`;
209+
for (let i = 0; i < t.params.c; i++) {
210+
for (let k = 0; k < t.params.r; k++) {
211+
lhs += `${kValue.f32.negative.min / 2},`;
212+
rhs += `${t.params.rhs},`;
213+
}
214+
}
215+
rhs += ')';
216+
lhs += ')';
217+
218+
const code = `
219+
@compute @workgroup_size(1)
220+
fn main() {
221+
const foo = ${lhs} - ${rhs};
222+
}
223+
`;
224+
225+
t.expectCompileResult(t.params.rhs === 1, code);
226+
});
227+
228+
g.test('overflow_f16')
229+
.desc(`Validates that f16 add overflows in shader creation`)
230+
.params(u =>
231+
u
232+
.combine('rhs', [kValue.f16.positive.max, 1])
233+
.combine('c', [2, 3, 4] as const)
234+
.combine('r', [2, 3, 4] as const)
235+
)
236+
.beforeAllSubcases(t => {
237+
t.selectDeviceOrSkipTestCase('shader-f16');
238+
})
239+
.fn(t => {
240+
let lhs = `mat${t.params.c}x${t.params.r}h(`;
241+
let rhs = `mat${t.params.c}x${t.params.r}h(`;
242+
for (let i = 0; i < t.params.c; i++) {
243+
for (let k = 0; k < t.params.r; k++) {
244+
lhs += `${kValue.f16.positive.max / 2},`;
245+
rhs += `${t.params.rhs},`;
246+
}
247+
}
248+
rhs += ')';
249+
lhs += ')';
250+
251+
const code = `
252+
enable f16;
253+
@compute @workgroup_size(1)
254+
fn main() {
255+
const foo = ${lhs} + ${rhs};
256+
}
257+
`;
258+
259+
t.expectCompileResult(t.params.rhs === 1, code);
260+
});
261+
262+
g.test('underflow_f16')
263+
.desc(`Validates that f16 add underflows in shader creation`)
264+
.params(u =>
265+
u
266+
.combine('rhs', [kValue.f16.positive.max, 1])
267+
.combine('c', [2, 3, 4] as const)
268+
.combine('r', [2, 3, 4] as const)
269+
)
270+
.beforeAllSubcases(t => {
271+
t.selectDeviceOrSkipTestCase('shader-f16');
272+
})
273+
.fn(t => {
274+
let lhs = `mat${t.params.c}x${t.params.r}h(`;
275+
let rhs = `mat${t.params.c}x${t.params.r}h(`;
276+
for (let i = 0; i < t.params.c; i++) {
277+
for (let k = 0; k < t.params.r; k++) {
278+
lhs += `${kValue.f32.negative.min / 2},`;
279+
rhs += `${t.params.rhs},`;
280+
}
281+
}
282+
rhs += ')';
283+
lhs += ')';
284+
285+
const code = `
286+
enable f16;
287+
@compute @workgroup_size(1)
288+
fn main() {
289+
const foo = ${lhs} - ${rhs};
290+
}
291+
`;
292+
293+
t.expectCompileResult(t.params.rhs === 1, code);
294+
});
295+
296+
g.test('overflow_abstract')
297+
.desc(`Validates that abstract add overflows in shader creation`)
298+
.params(u =>
299+
u
300+
.combine('rhs', [kValue.f64.positive.max, 1])
301+
.combine('c', [2, 3, 4] as const)
302+
.combine('r', [2, 3, 4] as const)
303+
)
304+
.fn(t => {
305+
let lhs = `mat${t.params.c}x${t.params.r}(`;
306+
let rhs = `mat${t.params.c}x${t.params.r}(`;
307+
for (let i = 0; i < t.params.c; i++) {
308+
for (let k = 0; k < t.params.r; k++) {
309+
lhs += `${kValue.f64.positive.max / 2},`;
310+
rhs += `${t.params.rhs},`;
311+
}
312+
}
313+
rhs += ')';
314+
lhs += ')';
315+
316+
const code = `
317+
@compute @workgroup_size(1)
318+
fn main() {
319+
const foo = ${lhs} + ${rhs};
320+
}
321+
`;
322+
323+
t.expectCompileResult(t.params.rhs === 1, code);
324+
});
325+
326+
g.test('underflow_abstract')
327+
.desc(`Validates that abstract add underflows in shader creation`)
328+
.params(u =>
329+
u
330+
.combine('rhs', [kValue.f64.positive.max, 1])
331+
.combine('c', [2, 3, 4] as const)
332+
.combine('r', [2, 3, 4] as const)
333+
)
334+
.fn(t => {
335+
let lhs = `mat${t.params.c}x${t.params.r}(`;
336+
let rhs = `mat${t.params.c}x${t.params.r}(`;
337+
for (let i = 0; i < t.params.c; i++) {
338+
for (let k = 0; k < t.params.r; k++) {
339+
lhs += `${kValue.f64.negative.min / 2},`;
340+
rhs += `${t.params.rhs},`;
341+
}
342+
}
343+
rhs += ')';
344+
lhs += ')';
345+
346+
const code = `
347+
@compute @workgroup_size(1)
348+
fn main() {
349+
const foo = ${lhs} - ${rhs};
350+
}
351+
`;
352+
353+
t.expectCompileResult(t.params.rhs === 1, code);
354+
});

0 commit comments

Comments
 (0)