From 7f64498d6e1f259c337481a8ae09a74c790ef44e Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Wed, 2 Nov 2022 19:42:18 -0400 Subject: [PATCH] Test `clamp` and Generalize Shader Tests (#3167) * Sample numeric tests * Generalize shader tests --- wgpu/tests/shader/mod.rs | 133 +++++++++++++++++++++++--- wgpu/tests/shader/numeric_builtins.rs | 56 +++++++++++ wgpu/tests/shader/shader_test.wgsl | 6 +- wgpu/tests/shader/struct_layout.rs | 122 +++++++++++------------ 4 files changed, 235 insertions(+), 82 deletions(-) create mode 100644 wgpu/tests/shader/numeric_builtins.rs diff --git a/wgpu/tests/shader/mod.rs b/wgpu/tests/shader/mod.rs index 3b7ea7faaa..cec5f9f33b 100644 --- a/wgpu/tests/shader/mod.rs +++ b/wgpu/tests/shader/mod.rs @@ -4,7 +4,7 @@ //! shader is run on the input buffer which generates an output buffer. This //! buffer is then read and compared to a given output. -use std::borrow::Cow; +use std::{borrow::Cow, fmt::Debug}; use wgpu::{ Backends, BindGroupDescriptor, BindGroupEntry, BindGroupLayoutDescriptor, BindGroupLayoutEntry, @@ -15,6 +15,7 @@ use wgpu::{ use crate::common::TestingContext; +mod numeric_builtins; mod struct_layout; #[derive(Clone, Copy, PartialEq)] @@ -40,21 +41,128 @@ struct ShaderTest { name: String, /// This text will be the body of the `Input` struct. Replaces "{{input_members}}" /// in the shader_test shader. - input_members: String, + custom_struct_members: String, /// This text will be the body of the compute shader. Replaces "{{body}}" /// in the shader_test shader. body: String, + /// This text will be the input type of the compute shader. Replaces "{{input_type}}". + /// + /// Defaults to "CustomStruct" + input_type: String, + /// This text will be the output type of the compute shader. Replaces "{{output_type}}". + /// + /// Defaults to "array". + output_type: String, /// List of values will be written to the input buffer. input_values: Vec, - /// List of expected outputs from the shader. - output_values: Vec, + /// List of lists of valid expected outputs from the shader. + output_values: Vec>, + /// Function which compares the output values to the resulting values and + /// prints a message on failure. + /// + /// Defaults [`Self::default_comparison_function`]. + output_comparison_fn: fn(&str, &[u32], &[Vec]) -> bool, /// Value to pre-initialize the output buffer to. Often u32::MAX so /// that writing a 0 looks different than not writing a value at all. + /// + /// Defaults to u32::MAX. output_initialization: u32, /// Which backends this test will fail on. If the test passes on this /// backend when it shouldn't, an assert will be raised. + /// + /// Defaults to Backends::empty(). failures: Backends, } +impl ShaderTest { + fn default_comparison_function( + test_name: &str, + actual_values: &[u32], + expected_values: &[Vec], + ) -> bool { + let cast_actual = bytemuck::cast_slice::(actual_values); + + // When printing the error message, we want to trim `cast_actual` to the length + // of the longest set of expected values. This tracks that value. + let mut max_relevant_value_count = 0; + + for expected in expected_values { + let cast_expected = bytemuck::cast_slice::(expected); + + // We shorten the actual to the length of the expected. + if &cast_actual[0..cast_expected.len()] == cast_expected { + return true; + } + + max_relevant_value_count = max_relevant_value_count.max(cast_expected.len()); + } + + // We haven't found a match, lets print an error. + + eprint!( + "Inner test failure. Actual {:?}. Expected", + &cast_actual[0..max_relevant_value_count] + ); + + if expected_values.len() != 1 { + eprint!(" one of: "); + } else { + eprint!(": "); + } + + for (idx, expected) in expected_values.iter().enumerate() { + let cast_expected = bytemuck::cast_slice::(expected); + eprint!("{cast_expected:?}"); + if idx + 1 != expected_values.len() { + eprint!(" "); + } + } + + eprintln!(". Test {test_name}"); + + false + } + + fn new( + name: String, + custom_struct_members: String, + body: String, + input_values: &[I], + output_values: &[O], + ) -> Self { + Self { + name, + custom_struct_members, + body, + input_type: String::from("CustomStruct"), + output_type: String::from("array"), + input_values: bytemuck::cast_slice(input_values).to_vec(), + output_values: vec![bytemuck::cast_slice(output_values).to_vec()], + output_comparison_fn: Self::default_comparison_function::, + output_initialization: u32::MAX, + failures: Backends::empty(), + } + } + + /// Add another set of possible outputs. If any of the given + /// output values are seen it's considered a success (i.e. this is OR, not AND). + /// + /// Assumes that this type O is the same as the O provided to new. + fn extra_output_values( + mut self, + output_values: &[O], + ) -> Self { + self.output_values + .push(bytemuck::cast_slice(output_values).to_vec()); + + self + } + + fn failures(mut self, failures: Backends) -> Self { + self.failures = failures; + + self + } +} const MAX_BUFFER_SIZE: u64 = 128; @@ -160,9 +268,13 @@ fn shader_input_output_test( // -- Building shader + pipeline -- + // This isn't terribly efficient but the string is short and it's a test. + // The body and input members are the longest part, so do them last. let mut processed = source .replace("{{storage_type}}", storage_type.as_str()) - .replace("{{input_members}}", &test.input_members) + .replace("{{input_type}}", &test.input_type) + .replace("{{output_type}}", &test.output_type) + .replace("{{input_members}}", &test.custom_struct_members) .replace("{{body}}", &test.body); // Add the bindings for all inputs besides push constants. @@ -239,17 +351,8 @@ fn shader_input_output_test( // -- Check results -- - let left = &typed[..test.output_values.len()]; - let right = test.output_values; - let failure = left != right; + let failure = !(test.output_comparison_fn)(&test_name, typed, &test.output_values); // We don't immediately panic to let all tests execute - if failure { - eprintln!( - "Inner test failure. Actual {:?}. Expected {:?}. Test {test_name}", - left.to_vec(), - right.to_vec(), - ); - } if failure != test .failures diff --git a/wgpu/tests/shader/numeric_builtins.rs b/wgpu/tests/shader/numeric_builtins.rs new file mode 100644 index 0000000000..83b278cfbf --- /dev/null +++ b/wgpu/tests/shader/numeric_builtins.rs @@ -0,0 +1,56 @@ +use wgpu::{DownlevelFlags, Limits}; + +use crate::{ + common::{initialize_test, TestParameters}, + shader::{shader_input_output_test, InputStorageType, ShaderTest}, +}; + +fn create_numeric_builtin_test() -> Vec { + let mut tests = Vec::new(); + + #[rustfmt::skip] + let clamp_values: &[(f32, f32, f32, &[f32])] = &[ + // value - low - high - valid outputs + + // normal clamps + ( 20.0, 0.0, 10.0, &[10.0]), + ( -10.0, 0.0, 10.0, &[0.0]), + ( 5.0, 0.0, 10.0, &[5.0]), + + // med-of-three or min/max + ( 3.0, 2.0, 1.0, &[1.0, 2.0]), + ]; + + for &(input, low, high, output) in clamp_values { + let mut test = ShaderTest::new( + format!("clamp({input}, 0.0, 10.0) == {output:?})"), + String::from("value: f32, low: f32, high: f32"), + String::from("output[0] = bitcast(clamp(input.value, input.low, input.high));"), + &[input, low, high], + &[output[0]], + ); + for &extra in &output[1..] { + test = test.extra_output_values(&[extra]); + } + + tests.push(test); + } + + tests +} + +#[test] +fn numeric_builtins() { + initialize_test( + TestParameters::default() + .downlevel_flags(DownlevelFlags::COMPUTE_SHADERS) + .limits(Limits::downlevel_defaults()), + |ctx| { + shader_input_output_test( + ctx, + InputStorageType::Storage, + create_numeric_builtin_test(), + ); + }, + ); +} diff --git a/wgpu/tests/shader/shader_test.wgsl b/wgpu/tests/shader/shader_test.wgsl index 9bb591037a..efe8692bd5 100644 --- a/wgpu/tests/shader/shader_test.wgsl +++ b/wgpu/tests/shader/shader_test.wgsl @@ -1,12 +1,12 @@ -struct InputStruct { +struct CustomStruct { {{input_members}} } {{input_bindings}} -var<{{storage_type}}> input: InputStruct; +var<{{storage_type}}> input: {{input_type}}; @group(0) @binding(1) -var output: array; +var output: {{output_type}}; @compute @workgroup_size(1) fn cs_main() { diff --git a/wgpu/tests/shader/struct_layout.rs b/wgpu/tests/shader/struct_layout.rs index 62d9bc418c..2250143b5f 100644 --- a/wgpu/tests/shader/struct_layout.rs +++ b/wgpu/tests/shader/struct_layout.rs @@ -9,7 +9,6 @@ use crate::{ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec { let input_values: Vec<_> = (0..(MAX_BUFFER_SIZE as u32 / 4)).collect(); - let output_initialization = u32::MAX; let mut tests = Vec::new(); @@ -35,25 +34,21 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec writeln!(loaded, "output[{idx}] = bitcast(loaded.{component});").unwrap(); } - tests.push(ShaderTest { - name: format!("vec{components}<{ty}> - direct"), - input_members: input_members.clone(), - body: direct, - input_values: input_values.clone(), - output_values: (0..components as u32).collect(), - output_initialization, - failures: Backends::empty(), - }); + tests.push(ShaderTest::new( + format!("vec{components}<{ty}> - direct"), + input_members.clone(), + direct, + &input_values, + &(0..components as u32).collect::>(), + )); - tests.push(ShaderTest { - name: format!("vec{components}<{ty}> - loaded"), - input_members, - body: loaded, - input_values: input_values.clone(), - output_values: (0..components as u32).collect(), - output_initialization, - failures: Backends::empty(), - }); + tests.push(ShaderTest::new( + format!("vec{components}<{ty}> - loaded"), + input_members.clone(), + loaded, + &input_values, + &(0..components as u32).collect::>(), + )); } } @@ -113,35 +108,38 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec Backends::empty() }; - tests.push(ShaderTest { - name: format!("{ty} - direct"), - input_members: input_members.clone(), - body: direct, - input_values: input_values.clone(), - output_values: output_values.clone(), - output_initialization, - failures, - }); + tests.push( + ShaderTest::new( + format!("{ty} - direct"), + input_members.clone(), + direct, + &input_values, + &output_values, + ) + .failures(failures), + ); - tests.push(ShaderTest { - name: format!("{ty} - vector loaded"), - input_members: input_members.clone(), - body: vector_loaded, - input_values: input_values.clone(), - output_values: output_values.clone(), - output_initialization, - failures, - }); + tests.push( + ShaderTest::new( + format!("{ty} - vector loaded"), + input_members.clone(), + vector_loaded, + &input_values, + &output_values, + ) + .failures(failures), + ); - tests.push(ShaderTest { - name: format!("{ty} - fully loaded"), - input_members, - body: fully_loaded, - input_values: input_values.clone(), - output_values, - output_initialization, - failures, - }); + tests.push( + ShaderTest::new( + format!("{ty} - fully loaded"), + input_members.clone(), + fully_loaded, + &input_values, + &output_values, + ) + .failures(failures), + ); } } @@ -150,15 +148,13 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec let members = format!("_vec: vec3<{ty}>,\nscalar: {ty},"); let direct = String::from("output[0] = bitcast(input.scalar);"); - tests.push(ShaderTest { - name: format!("vec3<{ty}>, {ty} alignment"), - input_members: members, - body: direct, - input_values: input_values.clone(), - output_values: vec![3], - output_initialization, - failures: Backends::empty(), - }); + tests.push(ShaderTest::new( + format!("vec3<{ty}>, {ty} alignment"), + members, + direct, + &input_values, + &[3], + )); } // Mat3 alignment tests @@ -167,15 +163,13 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec let members = format!("_mat: mat{columns}x3,\nscalar: {ty},"); let direct = String::from("output[0] = bitcast(input.scalar);"); - tests.push(ShaderTest { - name: format!("mat{columns}x3, {ty} alignment"), - input_members: members, - body: direct, - input_values: input_values.clone(), - output_values: vec![columns * 4], - output_initialization, - failures: Backends::empty(), - }); + tests.push(ShaderTest::new( + format!("mat{columns}x3, {ty} alignment"), + members, + direct, + &input_values, + &[columns * 4], + )); } }