[wgsl-in] Don't allow redefinition of module scope identifiers

src/front/wgsl/mod.rs

@@ -171,6 +171,10 @@ pub enum Error<'a> {
     NotPointer(Span),
     NotReference(&'static str, Span),
     ReservedKeyword(Span),
+    Redefinition {
+        previous: Span,
+        current: Span,
+    },
     Other,
 }
 
@@ -188,7 +192,7 @@ impl<'a> Error<'a> {
                                 Token::Number { value, .. } => {
                                     format!("number ({})", value)
                                 }
-                                Token::String(s) => format!("string literal ('{}')", s.to_string()),
+                                Token::String(s) => format!("string literal ('{}')", s),
                                 Token::Word(s) => s.to_string(),
                                 Token::Operation(c) => format!("operation ('{}')", c),
                                 Token::LogicalOperation(c) => format!("logical operation ('{}')", c),
@@ -439,10 +443,17 @@ impl<'a> Error<'a> {
                 notes: vec![],
             },
             Error::ReservedKeyword(ref name_span) => ParseError {
-                message: format!("Name `{}` is a reserved keyword", &source[name_span.clone()]),
+                message: format!("name `{}` is a reserved keyword", &source[name_span.clone()]),
                 labels: vec![(name_span.clone(), format!("definition of `{}`", &source[name_span.clone()]).into())],
                 notes: vec![],
             },
+            Error::Redefinition { ref previous, ref current } => ParseError {
+                message: format!("redefinition of `{}`", &source[current.clone()]),
+                labels: vec![(current.clone(), format!("redefinition of `{}`", &source[current.clone()]).into()),
+                             (previous.clone(), format!("previous definition of `{}`", &source[previous.clone()]).into())
+                ],
+                notes: vec![],
+            },
             Error::Other => ParseError {
                 message: "other error".to_string(),
                 labels: vec![],
@@ -1204,6 +1215,7 @@ impl std::error::Error for ParseError {
 
 pub struct Parser {
     scopes: Vec<(Scope, usize)>,
+    module_scope_identifiers: FastHashMap<String, Span>,
     lookup_type: FastHashMap<String, Handle<crate::Type>>,
     layouter: Layouter,
 }
@@ -1212,6 +1224,7 @@ impl Parser {
     pub fn new() -> Self {
         Parser {
             scopes: Vec::new(),
+            module_scope_identifiers: FastHashMap::default(),
             lookup_type: FastHashMap::default(),
             layouter: Default::default(),
         }
@@ -3837,6 +3850,15 @@ impl Parser {
         if crate::keywords::wgsl::RESERVED.contains(&fun_name) {
             return Err(Error::ReservedKeyword(span));
         }
+        if let Some(entry) = self
+            .module_scope_identifiers
+            .insert(String::from(fun_name), span.clone())
+        {
+            return Err(Error::Redefinition {
+                previous: entry,
+                current: span,
+            });
+        }
         // populate initial expressions
         let mut expressions = Arena::new();
         for (&name, expression) in lookup_global_expression.iter() {
@@ -4083,6 +4105,15 @@ impl Parser {
                 if crate::keywords::wgsl::RESERVED.contains(&name) {
                     return Err(Error::ReservedKeyword(name_span));
                 }
+                if let Some(entry) = self
+                    .module_scope_identifiers
+                    .insert(String::from(name), name_span.clone())
+                {
+                    return Err(Error::Redefinition {
+                        previous: entry,
+                        current: name_span,
+                    });
+                }
                 let given_ty = if lexer.skip(Token::Separator(':')) {
                     let (ty, _access) = self.parse_type_decl(
                         lexer,
@@ -4131,6 +4162,15 @@ impl Parser {
                 if crate::keywords::wgsl::RESERVED.contains(&pvar.name) {
                     return Err(Error::ReservedKeyword(pvar.name_span));
                 }
+                if let Some(entry) = self
+                    .module_scope_identifiers
+                    .insert(String::from(pvar.name), pvar.name_span.clone())
+                {
+                    return Err(Error::Redefinition {
+                        previous: entry,
+                        current: pvar.name_span,
+                    });
+                }
                 let var_handle = module.global_variables.append(
                     crate::GlobalVariable {
                         name: Some(pvar.name.to_owned()),

src/front/wgsl/tests.rs

@@ -408,7 +408,7 @@ fn parse_array_length() {
         [[group(0), binding(1)]]
         var<storage> bar: array<u32>;
 
-        fn foo() {
+        fn baz() {
             var x: u32 = arrayLength(foo.data);
             var y: u32 = arrayLength(bar);
         }

tests/in/globals.wgsl

@@ -1,12 +1,16 @@
-// Global variable & constant declarations
-
-let Foo: bool = true;
-
-var<workgroup> wg : array<f32, 10u>;
-var<workgroup> at: atomic<u32>;
-
-[[stage(compute), workgroup_size(1)]]
-fn main() {
-	wg[3] = 1.0;
-	atomicStore(&at, 2u);
-}
+// Global variable & constant declarations
+
+let Foo: bool = true;
+
+var<workgroup> wg : array<f32, 10u>;
+var<workgroup> at: atomic<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+	wg[3] = 1.0;
+	atomicStore(&at, 2u);
+
+    // Valid, Foo and at is in function scope
+    var Foo: f32 = 1.0;
+    var at: bool = true;
+}

tests/out/glsl/globals.main.Compute.glsl

@@ -7,12 +7,14 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
 shared float wg[10];
 
-shared uint at;
+shared uint at_1;
 
 
 void main() {
+    float Foo = 1.0;
+    bool at = true;
     wg[3] = 1.0;
-    at = 2u;
+    at_1 = 2u;
     return;
 }
 

tests/out/hlsl/globals.hlsl

@@ -1,12 +1,15 @@
-static const bool Foo = true;
+static const bool Foo_1 = true;
 
 groupshared float wg[10];
-groupshared uint at;
+groupshared uint at_1;
 
 [numthreads(1, 1, 1)]
 void main()
 {
+    float Foo = 1.0;
+    bool at = true;
+
     wg[3] = 1.0;
-    at = 2u;
+    at_1 = 2u;
     return;
 }

tests/out/msl/globals.msl

@@ -2,16 +2,18 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
-constexpr constant bool Foo = true;
+constexpr constant bool Foo_1 = true;
 struct type_2 {
     float inner[10u];
 };
 
 kernel void main_(
   threadgroup type_2& wg
-, threadgroup metal::atomic_uint& at
+, threadgroup metal::atomic_uint& at_1
 ) {
+    float Foo = 1.0;
+    bool at = true;
     wg.inner[3] = 1.0;
-    metal::atomic_store_explicit(&at, 2u, metal::memory_order_relaxed);
+    metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);
     return;
 }

tests/out/spv/globals.spvasm

@@ -1,13 +1,13 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 26
+; Bound: 31
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %18 "main"
-OpExecutionMode %18 LocalSize 1 1 1
-OpDecorate %12 ArrayStride 4
+OpEntryPoint GLCompute %23 "main"
+OpExecutionMode %23 LocalSize 1 1 1
+OpDecorate %13 ArrayStride 4
 %2 = OpTypeVoid
 %4 = OpTypeBool
 %3 = OpConstantTrue  %4
@@ -18,22 +18,27 @@ OpDecorate %12 ArrayStride 4
 %10 = OpTypeFloat 32
 %9 = OpConstant  %10  1.0
 %11 = OpConstant  %6  2
-%12 = OpTypeArray %10 %5
-%14 = OpTypePointer Workgroup %12
-%13 = OpVariable  %14  Workgroup
-%16 = OpTypePointer Workgroup %6
-%15 = OpVariable  %16  Workgroup
-%19 = OpTypeFunction %2
-%21 = OpTypePointer Workgroup %10
-%22 = OpConstant  %6  3
-%24 = OpConstant  %8  2
-%25 = OpConstant  %6  256
-%18 = OpFunction  %2  None %19
-%17 = OpLabel
-OpBranch %20
-%20 = OpLabel
-%23 = OpAccessChain  %21  %13 %22
-OpStore %23 %9
-OpAtomicStore %15 %24 %25 %11
+%12 = OpConstantTrue  %4
+%13 = OpTypeArray %10 %5
+%15 = OpTypePointer Workgroup %13
+%14 = OpVariable  %15  Workgroup
+%17 = OpTypePointer Workgroup %6
+%16 = OpVariable  %17  Workgroup
+%19 = OpTypePointer Function %10
+%21 = OpTypePointer Function %4
+%24 = OpTypeFunction %2
+%26 = OpTypePointer Workgroup %10
+%27 = OpConstant  %6  3
+%29 = OpConstant  %8  2
+%30 = OpConstant  %6  256
+%23 = OpFunction  %2  None %24
+%22 = OpLabel
+%18 = OpVariable  %19  Function %9
+%20 = OpVariable  %21  Function %12
+OpBranch %25
+%25 = OpLabel
+%28 = OpAccessChain  %26  %14 %27
+OpStore %28 %9
+OpAtomicStore %16 %29 %30 %11
 OpReturn
 OpFunctionEnd

tests/out/wgsl/globals.wgsl

@@ -1,11 +1,14 @@
-let Foo: bool = true;
+let Foo_1: bool = true;
 
 var<workgroup> wg: array<f32,10u>;
-var<workgroup> at: atomic<u32>;
+var<workgroup> at_1: atomic<u32>;
 
 [[stage(compute), workgroup_size(1, 1, 1)]]
 fn main() {
+    var Foo: f32 = 1.0;
+    var at: bool = true;
+
     wg[3] = 1.0;
-    atomicStore((&at), 2u);
+    atomicStore((&at_1), 2u);
     return;
 }