[naga wgsl] Experimental 64-bit floating-point literals.

In the WGSL front and back ends, support an `lf` suffix on
floating-point literals to yield 64-bit integer literals.

CHANGELOG.md

@@ -51,6 +51,9 @@ This adds a way to allow a Vulkan driver which is non-compliant per VK_KHR_drive
 
 #### OpenGL
 - `@builtin(instance_index)` now properly reflects the range provided in the draw call instead of always counting from 0. By @cwfitzgerald in [#4722](https://github.com/gfx-rs/wgpu/pull/4722).
+#### Naga
+
+- Naga's WGSL front and back ends now have experimental support for 64-bit floating-point literals: `1.0lf` denotes an `f64` value. There has been experimental support for an `f64` type for a while, but until now there was no syntax for writing literals with that type. As before, Naga module validation rejects `f64` values unless `naga::valid::Capabilities::FLOAT64` is requested. By @jimblandy in [#4747](https://github.com/gfx-rs/wgpu/pull/4747).
 
 ### Changes
 

naga/src/back/wgsl/writer.rs

@@ -1095,9 +1095,7 @@ impl<W: Write> Writer<W> {
                     crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
                     crate::Literal::I32(value) => write!(self.out, "{}", value)?,
                     crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
-                    crate::Literal::F64(_) => {
+                    crate::Literal::F64(value) => write!(self.out, "{:?}lf", value)?,
-                        return Err(Error::Custom("unsupported f64 literal".to_string()));
-                    }
                     crate::Literal::I64(_) => {
                         return Err(Error::Custom("unsupported i64 literal".to_string()));
                     }

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

@@ -1472,6 +1472,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                     ast::Literal::Number(Number::F32(f)) => crate::Literal::F32(f),
                     ast::Literal::Number(Number::I32(i)) => crate::Literal::I32(i),
                     ast::Literal::Number(Number::U32(u)) => crate::Literal::U32(u),
+                    ast::Literal::Number(Number::F64(f)) => crate::Literal::F64(f),
                     ast::Literal::Number(_) => {
                         unreachable!("got abstract numeric type when not expected");
                     }

naga/src/front/wgsl/parse/lexer.rs

@@ -625,6 +625,22 @@ fn test_numbers() {
     );
 }
 
+#[test]
+fn double_floats() {
+    sub_test(
+        "0x1.2p4lf 0x1p8lf 0.0625lf 625e-4lf 10lf 10l",
+        &[
+            Token::Number(Ok(Number::F64(18.0))),
+            Token::Number(Ok(Number::F64(256.0))),
+            Token::Number(Ok(Number::F64(0.0625))),
+            Token::Number(Ok(Number::F64(0.0625))),
+            Token::Number(Ok(Number::F64(10.0))),
+            Token::Number(Ok(Number::I32(10))),
+            Token::Word("l"),
+        ],
+    )
+}
+
 #[test]
 fn test_tokens() {
     sub_test("id123_OK", &[Token::Word("id123_OK")]);
@@ -679,7 +695,7 @@ fn test_tokens() {
     // Type suffixes are only allowed on hex float literals
     // if you provided an exponent.
     sub_test(
-        "0x1.2f 0x1.2f 0x1.2h 0x1.2H",
+        "0x1.2f 0x1.2f 0x1.2h 0x1.2H 0x1.2lf",
         &[
             // The 'f' suffixes are taken as a hex digit:
             // the fractional part is 0x2f / 256.
@@ -689,6 +705,8 @@ fn test_tokens() {
             Token::Word("h"),
             Token::Number(Ok(Number::F32(1.125))),
             Token::Word("H"),
+            Token::Number(Ok(Number::F32(1.125))),
+            Token::Word("lf"),
         ],
     )
 }

naga/src/front/wgsl/parse/number.rs

@@ -16,6 +16,8 @@ pub enum Number {
     U32(u32),
     /// Concrete f32
     F32(f32),
+    /// Concrete f64
+    F64(f64),
 }
 
 impl Number {
@@ -61,9 +63,11 @@ enum IntKind {
     U32,
 }
 
+#[derive(Debug)]
 enum FloatKind {
-    F32,
     F16,
+    F32,
+    F64,
 }
 
 // The following regexes (from the WGSL spec) will be matched:
@@ -104,9 +108,9 @@ fn parse(input: &str) -> (Result<Number, NumberError>, &str) {
     /// if one of the given patterns are found at the start of the buffer
     /// returning the corresponding expr for the matched pattern
     macro_rules! consume_map {
-        ($bytes:ident, [$($pattern:pat_param => $to:expr),*]) => {
+        ($bytes:ident, [$( $($pattern:pat_param),* => $to:expr),* $(,)?]) => {
             match $bytes {
-                $( &[$pattern, ref rest @ ..] => { $bytes = rest; Some($to) }, )*
+                $( &[ $($pattern),*, ref rest @ ..] => { $bytes = rest; Some($to) }, )*
                 _ => None,
             }
         };
@@ -136,6 +140,16 @@ fn parse(input: &str) -> (Result<Number, NumberError>, &str) {
         }};
     }
 
+    macro_rules! consume_float_suffix {
+        ($bytes:ident) => {
+            consume_map!($bytes, [
+                b'h' => FloatKind::F16,
+                b'f' => FloatKind::F32,
+                b'l', b'f' => FloatKind::F64,
+            ])
+        };
+    }
+
     /// maps the given `&[u8]` (tail of the initial `input: &str`) to a `&str`
     macro_rules! rest_to_str {
         ($bytes:ident) => {
@@ -190,7 +204,7 @@ fn parse(input: &str) -> (Result<Number, NumberError>, &str) {
 
                 let number = general_extract.end(bytes);
 
-                let kind = consume_map!(bytes, [b'f' => FloatKind::F32, b'h' => FloatKind::F16]);
+                let kind = consume_float_suffix!(bytes);
 
                 (parse_hex_float(number, kind), rest_to_str!(bytes))
             } else {
@@ -219,7 +233,7 @@ fn parse(input: &str) -> (Result<Number, NumberError>, &str) {
 
                 let exponent = exp_extract.end(bytes);
 
-                let kind = consume_map!(bytes, [b'f' => FloatKind::F32, b'h' => FloatKind::F16]);
+                let kind = consume_float_suffix!(bytes);
 
                 (
                     parse_hex_float_missing_period(significand, exponent, kind),
@@ -257,7 +271,7 @@ fn parse(input: &str) -> (Result<Number, NumberError>, &str) {
 
             let number = general_extract.end(bytes);
 
-            let kind = consume_map!(bytes, [b'f' => FloatKind::F32, b'h' => FloatKind::F16]);
+            let kind = consume_float_suffix!(bytes);
 
             (parse_dec_float(number, kind), rest_to_str!(bytes))
         } else {
@@ -275,7 +289,7 @@ fn parse(input: &str) -> (Result<Number, NumberError>, &str) {
 
                 let number = general_extract.end(bytes);
 
-                let kind = consume_map!(bytes, [b'f' => FloatKind::F32, b'h' => FloatKind::F16]);
+                let kind = consume_float_suffix!(bytes);
 
                 (parse_dec_float(number, kind), rest_to_str!(bytes))
             } else {
@@ -289,8 +303,9 @@ fn parse(input: &str) -> (Result<Number, NumberError>, &str) {
                 let kind = consume_map!(bytes, [
                     b'i' => Kind::Int(IntKind::I32),
                     b'u' => Kind::Int(IntKind::U32),
+                    b'h' => Kind::Float(FloatKind::F16),
                     b'f' => Kind::Float(FloatKind::F32),
-                    b'h' => Kind::Float(FloatKind::F16)
+                    b'l', b'f' => Kind::Float(FloatKind::F64),
                 ]);
 
                 (
@@ -382,12 +397,17 @@ fn parse_hex_float(input: &str, kind: Option<FloatKind>) -> Result<Number, Numbe
             // can only be ParseHexfErrorKind::Inexact but we can't check since it's private
             _ => Err(NumberError::NotRepresentable),
         },
+        Some(FloatKind::F16) => Err(NumberError::UnimplementedF16),
         Some(FloatKind::F32) => match hexf_parse::parse_hexf32(input, false) {
             Ok(num) => Ok(Number::F32(num)),
             // can only be ParseHexfErrorKind::Inexact but we can't check since it's private
             _ => Err(NumberError::NotRepresentable),
         },
-        Some(FloatKind::F16) => Err(NumberError::UnimplementedF16),
+        Some(FloatKind::F64) => match hexf_parse::parse_hexf64(input, false) {
+            Ok(num) => Ok(Number::F64(num)),
+            // can only be ParseHexfErrorKind::Inexact but we can't check since it's private
+            _ => Err(NumberError::NotRepresentable),
+        },
     }
 }
 
@@ -407,6 +427,12 @@ fn parse_dec_float(input: &str, kind: Option<FloatKind>) -> Result<Number, Numbe
                 .then_some(Number::F32(num))
                 .ok_or(NumberError::NotRepresentable)
         }
+        Some(FloatKind::F64) => {
+            let num = input.parse::<f64>().unwrap(); // will never fail
+            num.is_finite()
+                .then_some(Number::F64(num))
+                .ok_or(NumberError::NotRepresentable)
+        }
         Some(FloatKind::F16) => Err(NumberError::UnimplementedF16),
     }
 }

naga/tests/in/f64.param.ron

@@ -0,0 +1,12 @@
+(
+	god_mode: true,
+	spv: (
+		version: (1, 0),
+	),
+	glsl: (
+		version: Desktop(420),
+		writer_flags: (""),
+		binding_map: { },
+		zero_initialize_workgroup_memory: true,
+	),
+)

naga/tests/in/f64.wgsl

@@ -0,0 +1,13 @@
+var<private> v: f64 = 1lf;
+const k: f64 = 2.0lf;
+
+fn f(x: f64) -> f64 {
+   let y: f64 = 3e1lf + 4.0e2lf;
+   var z = y + f64(5);
+   return x + y + k + 5.0lf;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+   f(6.0lf);
+}

naga/tests/out/glsl/f64.main.Compute.glsl

@@ -0,0 +1,19 @@
+#version 420 core
+#extension GL_ARB_compute_shader : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+const double k = 2.0LF;
+
+
+double f(double x) {
+    double z = 0.0;
+    double y = (30.0LF + 400.0LF);
+    z = (y + double(5));
+    return (((x + y) + k) + 5.0LF);
+}
+
+void main() {
+    double _e1 = f(6.0LF);
+    return;
+}
+

naga/tests/out/hlsl/f64.hlsl

@@ -0,0 +1,19 @@
+static const double k = 2.0L;
+
+static double v = 1.0L;
+
+double f(double x)
+{
+    double z = (double)0;
+
+    double y = (30.0L + 400.0L);
+    z = (y + double(5));
+    return (((x + y) + k) + 5.0L);
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    const double _e1 = f(6.0L);
+    return;
+}

naga/tests/out/hlsl/f64.ron

@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)

naga/tests/out/spv/f64.spvasm

@@ -0,0 +1,48 @@
+; SPIR-V
+; Version: 1.0
+; Generator: rspirv
+; Bound: 33
+OpCapability Shader
+OpCapability Float64
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %28 "main"
+OpExecutionMode %28 LocalSize 1 1 1
+%2 = OpTypeVoid
+%3 = OpTypeFloat 64
+%4 = OpConstant  %3  1.0
+%5 = OpConstant  %3  2.0
+%7 = OpTypePointer Private %3
+%6 = OpVariable  %7  Private %4
+%11 = OpTypeFunction %3 %3
+%12 = OpConstant  %3  30.0
+%13 = OpConstant  %3  400.0
+%14 = OpTypeInt 32 1
+%15 = OpConstant  %14  5
+%16 = OpConstant  %3  5.0
+%18 = OpTypePointer Function %3
+%19 = OpConstantNull  %3
+%29 = OpTypeFunction %2
+%30 = OpConstant  %3  6.0
+%10 = OpFunction  %3  None %11
+%9 = OpFunctionParameter  %3
+%8 = OpLabel
+%17 = OpVariable  %18  Function %19
+OpBranch %20
+%20 = OpLabel
+%21 = OpFAdd  %3  %12 %13
+%22 = OpConvertSToF  %3  %15
+%23 = OpFAdd  %3  %21 %22
+OpStore %17 %23
+%24 = OpFAdd  %3  %9 %21
+%25 = OpFAdd  %3  %24 %5
+%26 = OpFAdd  %3  %25 %16
+OpReturnValue %26
+OpFunctionEnd
+%28 = OpFunction  %2  None %29
+%27 = OpLabel
+OpBranch %31
+%31 = OpLabel
+%32 = OpFunctionCall  %3  %10 %30
+OpReturn
+OpFunctionEnd

naga/tests/out/wgsl/f64.wgsl

@@ -0,0 +1,17 @@
+const k: f64 = 2.0lf;
+
+var<private> v: f64 = 1.0lf;
+
+fn f(x: f64) -> f64 {
+    var z: f64;
+
+    let y = (30.0lf + 400.0lf);
+    z = (y + f64(5));
+    return (((x + y) + k) + 5.0lf);
+}
+
+@compute @workgroup_size(1, 1, 1) 
+fn main() {
+    let _e1 = f(6.0lf);
+    return;
+}

naga/tests/snapshots.rs

@@ -779,6 +779,10 @@ fn convert_wgsl() {
             "struct-layout",
             Targets::WGSL | Targets::GLSL | Targets::SPIRV | Targets::HLSL | Targets::METAL,
         ),
+        (
+            "f64",
+            Targets::SPIRV | Targets::GLSL | Targets::HLSL | Targets::WGSL,
+        ),
     ];
 
     for &(name, targets) in inputs.iter() {