From c911aa7970a473b0c2887f099c0096340a6cc3a8 Mon Sep 17 00:00:00 2001 From: Jamie Nicol Date: Wed, 20 Nov 2024 16:13:54 +0000 Subject: [PATCH] [naga wgsl-in] Attempt automatic conversion for arguments to user defined function calls When lowering arguments for a user-defined function call, avoid concretizing the argument types. Instead make use of the existing `try_automatic_conversions()` machinery to attempt to convert each argument to the type expected by the function. This is straightforward as user-defined functions only have a single overload. This additionally changes an argument type in the test parse_pointers() from `ptr` to `ptr`. The former is invalid code which is indeed caught by the validator, but the test only asserts that parsing succeeds, not validation. With this patch, this error is now caught during parsing which caused the test to fail. --- CHANGELOG.md | 1 + naga/src/front/wgsl/lower/mod.rs | 22 ++- naga/src/front/wgsl/tests.rs | 24 ++- .../in/abstract-types-function-calls.wgsl | 38 +++++ .../hlsl/abstract-types-function-calls.ron | 8 + .../out/msl/abstract-types-function-calls.msl | 103 +++++++++++++ .../spv/abstract-types-function-calls.spvasm | 145 ++++++++++++++++++ .../wgsl/abstract-types-function-calls.wgsl | 64 ++++++++ naga/tests/snapshots.rs | 4 + 9 files changed, 407 insertions(+), 2 deletions(-) create mode 100644 naga/tests/in/abstract-types-function-calls.wgsl create mode 100644 naga/tests/out/hlsl/abstract-types-function-calls.ron create mode 100644 naga/tests/out/msl/abstract-types-function-calls.msl create mode 100644 naga/tests/out/spv/abstract-types-function-calls.spvasm create mode 100644 naga/tests/out/wgsl/abstract-types-function-calls.wgsl diff --git a/CHANGELOG.md b/CHANGELOG.md index bb86b1b53d..ad446f6de8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -100,6 +100,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148] - Fix textureNumLevels in the GLSL backend. By @magcius in [#6483](https://github.com/gfx-rs/wgpu/pull/6483). - Implement `quantizeToF16()` for WGSL frontend, and WGSL, SPIR-V, HLSL, MSL, and GLSL backends. By @jamienicol in [#6519](https://github.com/gfx-rs/wgpu/pull/6519). - Add support for GLSL `usampler*` and `isampler*`. By @DavidPeicho in [#6513](https://github.com/gfx-rs/wgpu/pull/6513). +- Implement type inference for abstract arguments to user-defined functions. By @jamienicol in [#6577](https://github.com/gfx-rs/wgpu/pull/6577). #### General diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 111d4f18d9..4b334b915f 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -2142,7 +2142,27 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { Some(&LoweredGlobalDecl::Function(function)) => { let arguments = arguments .iter() - .map(|&arg| self.expression(arg, ctx)) + .enumerate() + .map(|(i, &arg)| { + // Try to convert abstract values to the known argument types + match ctx.module.functions[function] + .arguments + .get(i) + .map(|arg| arg.ty) + { + Some(arg_ty) => { + let expr = self.expression_for_abstract(arg, ctx)?; + ctx.try_automatic_conversions( + expr, + &crate::proc::TypeResolution::Handle(arg_ty), + ctx.ast_expressions.get_span(arg), + ) + } + // Wrong number of arguments... just concretize the type here + // and let the validator report the error. + None => self.expression(arg, ctx), + } + }) .collect::, _>>()?; let has_result = ctx.module.functions[function].result.is_some(); diff --git a/naga/src/front/wgsl/tests.rs b/naga/src/front/wgsl/tests.rs index cc3d858317..7c13ca5894 100644 --- a/naga/src/front/wgsl/tests.rs +++ b/naga/src/front/wgsl/tests.rs @@ -84,6 +84,28 @@ fn parse_type_cast() { .is_err()); } +#[test] +fn parse_type_coercion() { + parse_str( + " + fn foo(bar: f32) {} + fn main() { + foo(0); + } + ", + ) + .unwrap(); + assert!(parse_str( + " + fn foo(bar: i32) {} + fn main() { + foo(0.0); + } + ", + ) + .is_err()); +} + #[test] fn parse_struct() { parse_str( @@ -461,7 +483,7 @@ fn binary_expression_mixed_scalar_and_vector_operands() { #[test] fn parse_pointers() { parse_str( - "fn foo(a: ptr) -> f32 { return *a; } + "fn foo(a: ptr) -> f32 { return *a; } fn bar() { var x: f32 = 1.0; let px = &x; diff --git a/naga/tests/in/abstract-types-function-calls.wgsl b/naga/tests/in/abstract-types-function-calls.wgsl new file mode 100644 index 0000000000..9c3c28ef06 --- /dev/null +++ b/naga/tests/in/abstract-types-function-calls.wgsl @@ -0,0 +1,38 @@ +fn func_f(a: f32) {} +fn func_i(a: i32) {} +fn func_u(a: u32) {} + +fn func_vf(a: vec2) {} +fn func_vi(a: vec2) {} +fn func_vu(a: vec2) {} + +fn func_mf(a: mat2x2) {} + +fn func_af(a: array) {} +fn func_ai(a: array) {} +fn func_au(a: array) {} + +fn func_f_i(a: f32, b: i32) {} + +fn main() { + func_f(0.0); + func_f(0); + func_i(0); + func_u(0); + + func_vf(vec2(0.0)); + func_vf(vec2(0)); + func_vi(vec2(0)); + func_vu(vec2(0)); + + func_mf(mat2x2(vec2(0.0), vec2(0.0))); + func_mf(mat2x2(vec2(0), vec2(0))); + + func_af(array(0.0, 0.0)); + func_af(array(0, 0)); + func_ai(array(0, 0)); + func_au(array(0, 0)); + + func_f_i(0.0, 0); + func_f_i(0, 0); +} diff --git a/naga/tests/out/hlsl/abstract-types-function-calls.ron b/naga/tests/out/hlsl/abstract-types-function-calls.ron new file mode 100644 index 0000000000..4d056ac29b --- /dev/null +++ b/naga/tests/out/hlsl/abstract-types-function-calls.ron @@ -0,0 +1,8 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ], +) diff --git a/naga/tests/out/msl/abstract-types-function-calls.msl b/naga/tests/out/msl/abstract-types-function-calls.msl new file mode 100644 index 0000000000..bf62d87a54 --- /dev/null +++ b/naga/tests/out/msl/abstract-types-function-calls.msl @@ -0,0 +1,103 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_7 { + float inner[2]; +}; +struct type_8 { + int inner[2]; +}; +struct type_9 { + uint inner[2]; +}; + +void func_f( + float a +) { + return; +} + +void func_i( + int a_1 +) { + return; +} + +void func_u( + uint a_2 +) { + return; +} + +void func_vf( + metal::float2 a_3 +) { + return; +} + +void func_vi( + metal::int2 a_4 +) { + return; +} + +void func_vu( + metal::uint2 a_5 +) { + return; +} + +void func_mf( + metal::float2x2 a_6 +) { + return; +} + +void func_af( + type_7 a_7 +) { + return; +} + +void func_ai( + type_8 a_8 +) { + return; +} + +void func_au( + type_9 a_9 +) { + return; +} + +void func_f_i( + float a_10, + int b +) { + return; +} + +void main_( +) { + func_f(0.0); + func_f(0.0); + func_i(0); + func_u(0u); + func_vf(metal::float2(0.0)); + func_vf(metal::float2(0.0)); + func_vi(metal::int2(0)); + func_vu(metal::uint2(0u)); + func_mf(metal::float2x2(metal::float2(0.0), metal::float2(0.0))); + func_mf(metal::float2x2(metal::float2(0.0), metal::float2(0.0))); + func_af(type_7 {0.0, 0.0}); + func_af(type_7 {0.0, 0.0}); + func_ai(type_8 {0, 0}); + func_au(type_9 {0u, 0u}); + func_f_i(0.0, 0); + func_f_i(0.0, 0); + return; +} diff --git a/naga/tests/out/spv/abstract-types-function-calls.spvasm b/naga/tests/out/spv/abstract-types-function-calls.spvasm new file mode 100644 index 0000000000..58e064bb4e --- /dev/null +++ b/naga/tests/out/spv/abstract-types-function-calls.spvasm @@ -0,0 +1,145 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 100 +OpCapability Shader +OpCapability Linkage +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpDecorate %10 ArrayStride 4 +OpDecorate %12 ArrayStride 4 +OpDecorate %13 ArrayStride 4 +%2 = OpTypeVoid +%3 = OpTypeFloat 32 +%4 = OpTypeInt 32 1 +%5 = OpTypeInt 32 0 +%6 = OpTypeVector %3 2 +%7 = OpTypeVector %4 2 +%8 = OpTypeVector %5 2 +%9 = OpTypeMatrix %6 2 +%11 = OpConstant %5 2 +%10 = OpTypeArray %3 %11 +%12 = OpTypeArray %4 %11 +%13 = OpTypeArray %5 %11 +%17 = OpTypeFunction %2 %3 +%22 = OpTypeFunction %2 %4 +%27 = OpTypeFunction %2 %5 +%32 = OpTypeFunction %2 %6 +%37 = OpTypeFunction %2 %7 +%42 = OpTypeFunction %2 %8 +%47 = OpTypeFunction %2 %9 +%52 = OpTypeFunction %2 %10 +%57 = OpTypeFunction %2 %12 +%62 = OpTypeFunction %2 %13 +%68 = OpTypeFunction %2 %3 %4 +%72 = OpTypeFunction %2 +%73 = OpConstant %3 0.0 +%74 = OpConstant %4 0 +%75 = OpConstant %5 0 +%76 = OpConstantComposite %6 %73 %73 +%77 = OpConstantComposite %7 %74 %74 +%78 = OpConstantComposite %8 %75 %75 +%79 = OpConstantComposite %9 %76 %76 +%80 = OpConstantComposite %10 %73 %73 +%81 = OpConstantComposite %12 %74 %74 +%82 = OpConstantComposite %13 %75 %75 +%16 = OpFunction %2 None %17 +%15 = OpFunctionParameter %3 +%14 = OpLabel +OpBranch %18 +%18 = OpLabel +OpReturn +OpFunctionEnd +%21 = OpFunction %2 None %22 +%20 = OpFunctionParameter %4 +%19 = OpLabel +OpBranch %23 +%23 = OpLabel +OpReturn +OpFunctionEnd +%26 = OpFunction %2 None %27 +%25 = OpFunctionParameter %5 +%24 = OpLabel +OpBranch %28 +%28 = OpLabel +OpReturn +OpFunctionEnd +%31 = OpFunction %2 None %32 +%30 = OpFunctionParameter %6 +%29 = OpLabel +OpBranch %33 +%33 = OpLabel +OpReturn +OpFunctionEnd +%36 = OpFunction %2 None %37 +%35 = OpFunctionParameter %7 +%34 = OpLabel +OpBranch %38 +%38 = OpLabel +OpReturn +OpFunctionEnd +%41 = OpFunction %2 None %42 +%40 = OpFunctionParameter %8 +%39 = OpLabel +OpBranch %43 +%43 = OpLabel +OpReturn +OpFunctionEnd +%46 = OpFunction %2 None %47 +%45 = OpFunctionParameter %9 +%44 = OpLabel +OpBranch %48 +%48 = OpLabel +OpReturn +OpFunctionEnd +%51 = OpFunction %2 None %52 +%50 = OpFunctionParameter %10 +%49 = OpLabel +OpBranch %53 +%53 = OpLabel +OpReturn +OpFunctionEnd +%56 = OpFunction %2 None %57 +%55 = OpFunctionParameter %12 +%54 = OpLabel +OpBranch %58 +%58 = OpLabel +OpReturn +OpFunctionEnd +%61 = OpFunction %2 None %62 +%60 = OpFunctionParameter %13 +%59 = OpLabel +OpBranch %63 +%63 = OpLabel +OpReturn +OpFunctionEnd +%67 = OpFunction %2 None %68 +%65 = OpFunctionParameter %3 +%66 = OpFunctionParameter %4 +%64 = OpLabel +OpBranch %69 +%69 = OpLabel +OpReturn +OpFunctionEnd +%71 = OpFunction %2 None %72 +%70 = OpLabel +OpBranch %83 +%83 = OpLabel +%84 = OpFunctionCall %2 %16 %73 +%85 = OpFunctionCall %2 %16 %73 +%86 = OpFunctionCall %2 %21 %74 +%87 = OpFunctionCall %2 %26 %75 +%88 = OpFunctionCall %2 %31 %76 +%89 = OpFunctionCall %2 %31 %76 +%90 = OpFunctionCall %2 %36 %77 +%91 = OpFunctionCall %2 %41 %78 +%92 = OpFunctionCall %2 %46 %79 +%93 = OpFunctionCall %2 %46 %79 +%94 = OpFunctionCall %2 %51 %80 +%95 = OpFunctionCall %2 %51 %80 +%96 = OpFunctionCall %2 %56 %81 +%97 = OpFunctionCall %2 %61 %82 +%98 = OpFunctionCall %2 %67 %73 %74 +%99 = OpFunctionCall %2 %67 %73 %74 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/abstract-types-function-calls.wgsl b/naga/tests/out/wgsl/abstract-types-function-calls.wgsl new file mode 100644 index 0000000000..419125bc6b --- /dev/null +++ b/naga/tests/out/wgsl/abstract-types-function-calls.wgsl @@ -0,0 +1,64 @@ +fn func_f(a: f32) { + return; +} + +fn func_i(a_1: i32) { + return; +} + +fn func_u(a_2: u32) { + return; +} + +fn func_vf(a_3: vec2) { + return; +} + +fn func_vi(a_4: vec2) { + return; +} + +fn func_vu(a_5: vec2) { + return; +} + +fn func_mf(a_6: mat2x2) { + return; +} + +fn func_af(a_7: array) { + return; +} + +fn func_ai(a_8: array) { + return; +} + +fn func_au(a_9: array) { + return; +} + +fn func_f_i(a_10: f32, b: i32) { + return; +} + +fn main() { + func_f(0f); + func_f(0f); + func_i(0i); + func_u(0u); + func_vf(vec2(0f)); + func_vf(vec2(0f)); + func_vi(vec2(0i)); + func_vu(vec2(0u)); + func_mf(mat2x2(vec2(0f), vec2(0f))); + func_mf(mat2x2(vec2(0f), vec2(0f))); + func_af(array(0f, 0f)); + func_af(array(0f, 0f)); + func_ai(array(0i, 0i)); + func_au(array(0u, 0u)); + func_f_i(0f, 0i); + func_f_i(0f, 0i); + return; +} + diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 7a463946ef..a5116d7501 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -893,6 +893,10 @@ fn convert_wgsl() { "abstract-types-const", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, ), + ( + "abstract-types-function-calls", + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, + ), ( "abstract-types-var", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,