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

CHANGELOG.md

@@ -47,6 +47,9 @@ Bottom level categories:
 
 #### 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(_) => {
-                        return Err(Error::Custom("unsupported f64 literal".to_string()));
-                    }
+                    crate::Literal::F64(value) => write!(self.out, "{:?}lf", value)?,
                     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/src/proc/constant_evaluator.rs

@@ -998,6 +998,14 @@ impl<'a> ConstantEvaluator<'a> {
                             return Err(ConstantEvaluatorError::InvalidCastArg)
                         }
                     }),
+                    Sc::F64 => Literal::F64(match literal {
+                        Literal::I32(v) => v as f64,
+                        Literal::U32(v) => v as f64,
+                        Literal::F32(v) => v as f64,
+                        Literal::Bool(v) => v as u32 as f64,
+                        Literal::F64(v) => v,
+                        Literal::I64(_) => return Err(ConstantEvaluatorError::InvalidCastArg),
+                    }),
                     Sc::BOOL => Literal::Bool(match literal {
                         Literal::I32(v) => v != 0,
                         Literal::U32(v) => v != 0,

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 + 5.0LF);
+    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 + 5.0L);
+    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,45 @@
+; SPIR-V
+; Version: 1.0
+; Generator: rspirv
+; Bound: 30
+OpCapability Shader
+OpCapability Float64
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %25 "main"
+OpExecutionMode %25 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 = OpConstant  %3  5.0
+%16 = OpTypePointer Function %3
+%17 = OpConstantNull  %3
+%26 = OpTypeFunction %2
+%27 = OpConstant  %3  6.0
+%10 = OpFunction  %3  None %11
+%9 = OpFunctionParameter  %3
+%8 = OpLabel
+%15 = OpVariable  %16  Function %17
+OpBranch %18
+%18 = OpLabel
+%19 = OpFAdd  %3  %12 %13
+%20 = OpFAdd  %3  %19 %14
+OpStore %15 %20
+%21 = OpFAdd  %3  %9 %19
+%22 = OpFAdd  %3  %21 %5
+%23 = OpFAdd  %3  %22 %14
+OpReturnValue %23
+OpFunctionEnd
+%25 = OpFunction  %2  None %26
+%24 = OpLabel
+OpBranch %28
+%28 = OpLabel
+%29 = OpFunctionCall  %3  %10 %27
+OpReturn
+OpFunctionEnd

naga/tests/out/wgsl/double-math-functions.frag.wgsl

@@ -1,6 +1,6 @@
 fn main_1() {
-    var a: vec4<f64>;
-    var b: vec4<f64>;
+    var a: vec4<f64> = vec4(1.0lf);
+    var b: vec4<f64> = vec4(2.0lf);
     var m: mat4x4<f64>;
     var i: i32 = 5;
     var ceilOut: vec4<f64>;
@@ -29,8 +29,6 @@ fn main_1() {
     var smoothStepVector: vec4<f64>;
     var smoothStepMixed: vec4<f64>;
 
-    a = vec4(f64(1.0));
-    b = vec4(f64(2.0));
     let _e8 = a;
     let _e9 = b;
     let _e10 = a;
@@ -95,8 +93,8 @@ fn main_1() {
     let _e152 = i;
     ldexpOut = ldexp(_e150.x, _e152);
     smoothStepScalar = f64(smoothstep(0.0, 1.0, 0.5));
-    smoothStepVector = smoothstep(vec4(f64(0.0)), vec4(f64(1.0)), vec4(f64(0.5)));
-    smoothStepMixed = smoothstep(vec4(f64(0.0)), vec4(f64(1.0)), vec4(f64(0.5)));
+    smoothStepVector = smoothstep(vec4(0.0lf), vec4(1.0lf), vec4(0.5lf));
+    smoothStepMixed = smoothstep(vec4(0.0lf), vec4(1.0lf), vec4(0.5lf));
     return;
 }
 

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 + 5.0lf);
+    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() {