[naga wgsl-in] Allow abstract literals to be used as return values

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

@@ -1677,7 +1677,23 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                 emitter.start(&ctx.function.expressions);
 
                 let value = value
-                    .map(|expr| self.expression(expr, &mut ctx.as_expression(block, &mut emitter)))
+                    .map(|expr| {
+                        let expr = self.expression_for_abstract(
+                            expr,
+                            &mut ctx.as_expression(block, &mut emitter),
+                        )?;
+
+                        if let Some(result_ty) = ctx.function.result.as_ref().map(|r| r.ty) {
+                            ctx.as_expression(block, &mut emitter)
+                                .try_automatic_conversions(
+                                    expr,
+                                    &crate::proc::TypeResolution::Handle(result_ty),
+                                    Span::default(),
+                                )
+                        } else {
+                            Ok(expr)
+                        }
+                    })
                     .transpose()?;
                 block.extend(emitter.finish(&ctx.function.expressions));
 

naga/tests/in/abstract-types-return.wgsl

@@ -0,0 +1,19 @@
+fn return_i32_ai() -> i32 {
+    return 1;
+}
+
+fn return_u32_ai() -> u32 {
+    return 1;
+}
+
+fn return_f32_ai() -> f32 {
+    return 1;
+}
+
+fn return_f32_af() -> f32 {
+    return 1.0;
+}
+
+fn return_vec2f32_ai() -> vec2<f32> {
+    return vec2(1);
+}

naga/tests/out/msl/abstract-types-return.msl

@@ -0,0 +1,31 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+int return_i32_ai(
+) {
+    return 1;
+}
+
+uint return_u32_ai(
+) {
+    return 1u;
+}
+
+float return_f32_ai(
+) {
+    return 1.0;
+}
+
+float return_f32_af(
+) {
+    return 1.0;
+}
+
+metal::float2 return_vec2f32_ai(
+) {
+    return metal::float2(1.0);
+}

naga/tests/out/spv/abstract-types-return.spvasm

@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 30
+OpCapability Shader
+OpCapability Linkage
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+%2 = OpTypeVoid
+%3 = OpTypeInt 32 1
+%4 = OpTypeInt 32 0
+%5 = OpTypeFloat 32
+%6 = OpTypeVector %5 2
+%9 = OpTypeFunction %3
+%10 = OpConstant  %3  1
+%14 = OpTypeFunction %4
+%15 = OpConstant  %4  1
+%19 = OpTypeFunction %5
+%20 = OpConstant  %5  1.0
+%27 = OpTypeFunction %6
+%28 = OpConstantComposite  %6  %20 %20
+%8 = OpFunction  %3  None %9
+%7 = OpLabel
+OpBranch %11
+%11 = OpLabel
+OpReturnValue %10
+OpFunctionEnd
+%13 = OpFunction  %4  None %14
+%12 = OpLabel
+OpBranch %16
+%16 = OpLabel
+OpReturnValue %15
+OpFunctionEnd
+%18 = OpFunction  %5  None %19
+%17 = OpLabel
+OpBranch %21
+%21 = OpLabel
+OpReturnValue %20
+OpFunctionEnd
+%23 = OpFunction  %5  None %19
+%22 = OpLabel
+OpBranch %24
+%24 = OpLabel
+OpReturnValue %20
+OpFunctionEnd
+%26 = OpFunction  %6  None %27
+%25 = OpLabel
+OpBranch %29
+%29 = OpLabel
+OpReturnValue %28
+OpFunctionEnd

naga/tests/out/wgsl/abstract-types-return.wgsl

@@ -0,0 +1,20 @@
+fn return_i32_ai() -> i32 {
+    return 1i;
+}
+
+fn return_u32_ai() -> u32 {
+    return 1u;
+}
+
+fn return_f32_ai() -> f32 {
+    return 1f;
+}
+
+fn return_f32_af() -> f32 {
+    return 1f;
+}
+
+fn return_vec2f32_ai() -> vec2<f32> {
+    return vec2(1f);
+}
+

naga/tests/snapshots.rs

@@ -921,6 +921,10 @@ fn convert_wgsl() {
             "abstract-types-operators",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
         ),
+        (
+            "abstract-types-return",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
+        ),
         (
             "int64",
             Targets::SPIRV | Targets::HLSL | Targets::WGSL | Targets::METAL,