From 8b6fa54564eb6bad5221dda7edc733c363f8effa Mon Sep 17 00:00:00 2001 From: Jamie Nicol Date: Thu, 30 Jan 2025 16:28:21 +0000 Subject: [PATCH] [naga wgsl-in] Allow abstract literals to be used as return values --- naga/src/front/wgsl/lower/mod.rs | 24 ++++++++- naga/tests/in/abstract-types-return.wgsl | 19 +++++++ naga/tests/out/msl/abstract-types-return.msl | 31 +++++++++++ .../out/spv/abstract-types-return.spvasm | 51 +++++++++++++++++++ .../tests/out/wgsl/abstract-types-return.wgsl | 20 ++++++++ naga/tests/snapshots.rs | 4 ++ 6 files changed, 148 insertions(+), 1 deletion(-) create mode 100644 naga/tests/in/abstract-types-return.wgsl create mode 100644 naga/tests/out/msl/abstract-types-return.msl create mode 100644 naga/tests/out/spv/abstract-types-return.spvasm create mode 100644 naga/tests/out/wgsl/abstract-types-return.wgsl diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index d25eb362c1..b8d638a99e 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1677,7 +1677,29 @@ 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) { + let mut ctx = ctx.as_expression(block, &mut emitter); + let expr_ty = resolve_inner!(ctx, expr); + if expr_ty.scalar().is_some_and(|s| s.is_abstract()) { + ctx + .try_automatic_conversions( + expr, + &crate::proc::TypeResolution::Handle(result_ty), + Span::default(), + ) + } else { + Ok(expr) + } + } else { + Ok(expr) + } + }) .transpose()?; block.extend(emitter.finish(&ctx.function.expressions)); diff --git a/naga/tests/in/abstract-types-return.wgsl b/naga/tests/in/abstract-types-return.wgsl new file mode 100644 index 0000000000..76c8222649 --- /dev/null +++ b/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 { + return vec2(1); +} diff --git a/naga/tests/out/msl/abstract-types-return.msl b/naga/tests/out/msl/abstract-types-return.msl new file mode 100644 index 0000000000..a926cd56a4 --- /dev/null +++ b/naga/tests/out/msl/abstract-types-return.msl @@ -0,0 +1,31 @@ +// language: metal1.0 +#include +#include + +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); +} diff --git a/naga/tests/out/spv/abstract-types-return.spvasm b/naga/tests/out/spv/abstract-types-return.spvasm new file mode 100644 index 0000000000..d2255080d9 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/naga/tests/out/wgsl/abstract-types-return.wgsl b/naga/tests/out/wgsl/abstract-types-return.wgsl new file mode 100644 index 0000000000..1a9634dc92 --- /dev/null +++ b/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 { + return vec2(1f); +} + diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 829019d99b..7e303f2302 100644 --- a/naga/tests/snapshots.rs +++ b/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,