[naga wgsl-in] Attempt automatic conversion for arguments to user defined function calls

CHANGELOG.md

@@ -115,6 +115,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148]
 - Expose Ray Query flags as constants in WGSL. Implement candidate intersections. By @kvark in [#5429](https://github.com/gfx-rs/wgpu/pull/5429)
 - Allow for override-expressions in `workgroup_size`. By @KentSlaney in [#6635](https://github.com/gfx-rs/wgpu/pull/6635).
 - Add support for OpAtomicCompareExchange in SPIR-V frontend. By @schell in [#6590](https://github.com/gfx-rs/wgpu/pull/6590).
+- Implement type inference for abstract arguments to user-defined functions. By @jamienicol in [#6577](https://github.com/gfx-rs/wgpu/pull/6577).
 
 #### General
 

naga/src/front/wgsl/lower/mod.rs

@@ -2184,7 +2184,25 @@ 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
+                        let Some(&crate::FunctionArgument {
+                            ty: parameter_ty, ..
+                        }) = ctx.module.functions[function].arguments.get(i)
+                        else {
+                            // Wrong number of arguments... just concretize the type here
+                            // and let the validator report the error.
+                            return self.expression(arg, ctx);
+                        };
+
+                        let expr = self.expression_for_abstract(arg, ctx)?;
+                        ctx.try_automatic_conversions(
+                            expr,
+                            &crate::proc::TypeResolution::Handle(parameter_ty),
+                            ctx.ast_expressions.get_span(arg),
+                        )
+                    })
                     .collect::<Result<Vec<_>, _>>()?;
 
                 let has_result = ctx.module.functions[function].result.is_some();

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<private, f32>) -> f32 { return *a; }
+        "fn foo(a: ptr<function, f32>) -> f32 { return *a; }
     fn bar() {
         var x: f32 = 1.0;
         let px = &x;

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<f32>) {}
+fn func_vi(a: vec2<i32>) {}
+fn func_vu(a: vec2<u32>) {}
+
+fn func_mf(a: mat2x2<f32>) {}
+
+fn func_af(a: array<f32, 2>) {}
+fn func_ai(a: array<i32, 2>) {}
+fn func_au(a: array<u32, 2>) {}
+
+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);
+}

naga/tests/out/hlsl/abstract-types-function-calls.ron

@@ -0,0 +1,8 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+    ],
+)

naga/tests/out/msl/abstract-types-function-calls.msl

@@ -0,0 +1,103 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+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;
+}

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