From b11f094287e8044fe295aaf34ba3e6853345209f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=A3o=20Capucho?= Date: Mon, 16 Aug 2021 21:33:56 +0100 Subject: [PATCH] [wgsl-out] Write pointers types and loads --- src/back/wgsl/writer.rs | 143 +++-- src/front/wgsl/mod.rs | 901 ++++++++++++++------------- tests/out/wgsl/246-collatz-comp.wgsl | 2 +- tests/out/wgsl/access.wgsl | 2 +- tests/out/wgsl/boids.wgsl | 2 +- tests/out/wgsl/collatz.wgsl | 2 +- 6 files changed, 583 insertions(+), 469 deletions(-) diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index ec2683fcab..829f93380b 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -6,11 +6,6 @@ use crate::{ }; use std::fmt::Write; -// This is a hack: we need to pass a pointer to an atomic, -// but generally the backend isn't putting "&" in front of every pointer. -// Some more general handling of pointers is needed to be implemented here. -const ATOMIC_REFERENCE: &str = "&"; - /// Shorthand result used internally by the backend type BackendResult = Result<(), Error>; @@ -537,9 +532,16 @@ impl Writer { )?; } TypeInner::Pointer { base, class } => { + let (storage, maybe_access) = storage_class_str(class); + if let Some(class) = storage { + write!(self.out, "ptr<{}, ", class)?; + if let Some(access) = maybe_access { + write!(self.out, ", {}", access)?; + } + } self.write_type(module, base)?; - if let Some(storage_class) = storage_class_str(class) { - write!(self.out, "<{}>", storage_class)?; + if storage.is_some() { + write!(self.out, ">")?; } } _ => { @@ -673,15 +675,22 @@ impl Writer { } write!(self.out, "{}", INDENT.repeat(indent))?; - let is_atomic = match *func_ctx.info[pointer].ty.inner_with(&module.types) { - crate::TypeInner::Pointer { base, .. } => match module.types[base].inner { - crate::TypeInner::Atomic { .. } => true, - _ => false, - }, - _ => false, + let (is_ptr, is_atomic) = match *func_ctx.info[pointer].ty.inner_with(&module.types) + { + crate::TypeInner::Pointer { base, .. } => ( + func_ctx.expressions[pointer].should_deref(), + match module.types[base].inner { + crate::TypeInner::Atomic { .. } => true, + _ => false, + }, + ), + _ => (false, false), }; if is_atomic { - write!(self.out, "atomicStore({}", ATOMIC_REFERENCE)?; + write!(self.out, "atomicStore(")?; + if !is_ptr { + write!(self.out, "&")?; + } self.write_expr(module, pointer, func_ctx)?; write!(self.out, ", ")?; self.write_expr(module, value, func_ctx)?; @@ -727,8 +736,13 @@ impl Writer { self.start_named_expr(module, result, func_ctx, &res_name)?; self.named_expressions.insert(result, res_name); + let is_ptr = func_ctx.expressions[pointer].should_deref(); + let fun_str = fun.to_wgsl(); - write!(self.out, "atomic{}({}", fun_str, ATOMIC_REFERENCE)?; + write!(self.out, "atomic{}(", fun_str)?; + if !is_ptr { + write!(self.out, "&")?; + } self.write_expr(module, pointer, func_ctx)?; if let crate::AtomicFunction::Exchange { compare: Some(cmp) } = *fun { write!(self.out, ", ")?; @@ -973,10 +987,22 @@ impl Writer { } // TODO: copy-paste from glsl-out Expression::AccessIndex { base, index } => { - self.write_expr(module, base, func_ctx)?; - let base_ty_res = &func_ctx.info[base].ty; let mut resolved = base_ty_res.inner_with(&module.types); + + let deref = match *resolved { + TypeInner::Pointer { .. } => func_ctx.expressions[base].should_deref(), + _ => false, + }; + + if deref { + write!(self.out, "(*")?; + } + self.write_expr(module, base, func_ctx)?; + if deref { + write!(self.out, ")")?; + } + let base_ty_handle = match *resolved { TypeInner::Pointer { base, class: _ } => { resolved = &module.types[base].inner; @@ -1169,15 +1195,26 @@ impl Writer { write!(self.out, ")")?; } Expression::Load { pointer } => { - let is_atomic = match *func_ctx.info[pointer].ty.inner_with(&module.types) { - crate::TypeInner::Pointer { base, .. } => match module.types[base].inner { - crate::TypeInner::Atomic { .. } => true, - _ => false, - }, - _ => false, - }; + let (is_pointer, is_atomic) = + match *func_ctx.info[pointer].ty.inner_with(&module.types) { + crate::TypeInner::Pointer { base, .. } => ( + func_ctx.expressions[pointer].should_deref(), + match module.types[base].inner { + crate::TypeInner::Atomic { .. } => true, + _ => false, + }, + ), + _ => (false, false), + }; if is_atomic { - write!(self.out, "atomicLoad({}", ATOMIC_REFERENCE)?; + write!(self.out, "atomicLoad(")?; + if !is_pointer { + // Write an indirection in case the underlying + // expression isn't a pointer but a reference + write!(self.out, "&")?; + } + } else if is_pointer { + write!(self.out, "*")?; } self.write_expr(module, pointer, func_ctx)?; if is_atomic { @@ -1377,8 +1414,13 @@ impl Writer { // First write global name and storage class if supported write!(self.out, "var")?; - if let Some(storage_class) = storage_class_str(global.class) { - write!(self.out, "<{}>", storage_class)?; + let (storage, maybe_access) = storage_class_str(global.class); + if let Some(class) = storage { + write!(self.out, "<{}", class)?; + if let Some(access) = maybe_access { + write!(self.out, ", {}", access)?; + } + write!(self.out, ">")?; } write!(self.out, " {}: ", name)?; @@ -1525,6 +1567,21 @@ impl Writer { } } +impl crate::Expression { + /// Wether an expression should be dereferenced, this is false when the + /// expression returns a reference instead of a pointer + fn should_deref(&self) -> bool { + match *self { + // Variables in the typifier have pointer types but in wgsl they + // have reference types and shouldn't be dereferenced + crate::Expression::LocalVariable(_) | crate::Expression::GlobalVariable(_) + // Access chains might have pointer types but wgsl considers them as references + | crate::Expression::AccessIndex {..} | crate::Expression::Access {..} => false, + _ => true, + } + } +} + fn builtin_str(built_in: crate::BuiltIn) -> Option<&'static str> { use crate::BuiltIn as Bi; @@ -1629,21 +1686,29 @@ fn sampling_str(sampling: crate::Sampling) -> &'static str { } } -fn storage_class_str(storage_class: crate::StorageClass) -> Option<&'static str> { +fn storage_class_str( + storage_class: crate::StorageClass, +) -> (Option<&'static str>, Option<&'static str>) { use crate::StorageClass as Sc; - match storage_class { - Sc::Private => Some("private"), - Sc::Uniform => Some("uniform"), - Sc::Storage { access } => Some(if access.contains(crate::StorageAccess::STORE) { - "storage,read_write" - } else { - "storage" + ( + Some(match storage_class { + Sc::Private => "private", + Sc::Uniform => "uniform", + Sc::Storage { access } => { + if access.contains(crate::StorageAccess::STORE) { + return (Some("storage"), Some("read_write")); + } else { + "storage" + } + } + Sc::PushConstant => "push_constant", + Sc::WorkGroup => "workgroup", + Sc::Handle => return (None, None), + Sc::Function => "function", }), - Sc::PushConstant => Some("push_constant"), - Sc::WorkGroup => Some("workgroup"), - Sc::Function | Sc::Handle => None, - } + None, + ) } fn map_binding_to_attribute( diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index 2f2b5e5c6a..4bbc70231d 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -2103,6 +2103,17 @@ impl Parser { ctx.expressions.append(expr, NagaSpan::ByteRange(span)), ) } + Token::Operation('*') => { + let _ = lexer.next(); + let expr = crate::Expression::Load { + pointer: self.parse_singular_expression(lexer, ctx.reborrow())?, + }; + let span = self.peek_scope(lexer); + ( + false, + ctx.expressions.append(expr, NagaSpan::ByteRange(span)), + ) + } Token::Operation('&') => { let _ = lexer.next(); let handle = self.parse_primary_expression(lexer, ctx.reborrow())?; @@ -2854,7 +2865,7 @@ impl Parser { is_uniform_control_flow: bool, ) -> Result<(), Error<'a>> { self.push_scope(Scope::Statement, lexer); - let (word, word_span) = match lexer.peek() { + match lexer.peek() { (Token::Separator(';'), _) => { let _ = lexer.next(); self.pop_scope(lexer); @@ -2880,83 +2891,36 @@ impl Parser { ); return Ok(()); } - (Token::Word(word), span) => (word, span), - other => return Err(Error::Unexpected(other, ExpectedToken::Statement)), - }; - let mut emitter = super::Emitter::default(); - let statement = match word { - "let" => { - let _ = lexer.next(); - emitter.start(context.expressions); - let (name, name_span) = lexer.next_ident_with_span()?; - let given_ty = if lexer.skip(Token::Separator(':')) { - let (ty, _access) = - self.parse_type_decl(lexer, None, context.types, context.constants)?; - Some(ty) - } else { - None - }; - lexer.expect(Token::Operation('='))?; - let expr_id = self - .parse_general_expression(lexer, context.as_expression(block, &mut emitter))?; - lexer.expect(Token::Separator(';'))?; - if let Some(ty) = given_ty { - // prepare the typifier, but work around mutable borrowing... - let _ = context - .as_expression(block, &mut emitter) - .resolve_type(expr_id)?; - let expr_inner = context.typifier.get(expr_id, context.types); - let given_inner = &context.types[ty].inner; - if given_inner != expr_inner { - log::error!( - "Given type {:?} doesn't match expected {:?}", - given_inner, - expr_inner - ); - return Err(Error::InitializationTypeMismatch(name_span, ty)); - } - } - block.extend(emitter.finish(context.expressions)); - context.lookup_ident.insert(name, expr_id); - context - .named_expressions - .insert(expr_id, String::from(name)); - None - } - "var" => { - let _ = lexer.next(); - enum Init { - Empty, - Constant(Handle), - Variable(Handle), - } - - let (name, name_span) = lexer.next_ident_with_span()?; - let given_ty = if lexer.skip(Token::Separator(':')) { - let (ty, _access) = - self.parse_type_decl(lexer, None, context.types, context.constants)?; - Some(ty) - } else { - None - }; - - let (init, ty) = if lexer.skip(Token::Operation('=')) { - emitter.start(context.expressions); - let value = self.parse_general_expression( - lexer, - context.as_expression(block, &mut emitter), - )?; - block.extend(emitter.finish(context.expressions)); - - // prepare the typifier, but work around mutable borrowing... - let _ = context - .as_expression(block, &mut emitter) - .resolve_type(value)?; - - //TODO: share more of this code with `let` arm - let ty = match given_ty { - Some(ty) => { - let expr_inner = context.typifier.get(value, context.types); + (Token::Word(word), word_span) => { + let mut emitter = super::Emitter::default(); + let statement = match word { + "let" => { + let _ = lexer.next(); + emitter.start(context.expressions); + let (name, name_span) = lexer.next_ident_with_span()?; + let given_ty = if lexer.skip(Token::Separator(':')) { + let (ty, _access) = self.parse_type_decl( + lexer, + None, + context.types, + context.constants, + )?; + Some(ty) + } else { + None + }; + lexer.expect(Token::Operation('='))?; + let expr_id = self.parse_general_expression( + lexer, + context.as_expression(block, &mut emitter), + )?; + lexer.expect(Token::Separator(';'))?; + if let Some(ty) = given_ty { + // prepare the typifier, but work around mutable borrowing... + let _ = context + .as_expression(block, &mut emitter) + .resolve_type(expr_id)?; + let expr_inner = context.typifier.get(expr_id, context.types); let given_inner = &context.types[ty].inner; if given_inner != expr_inner { log::error!( @@ -2966,378 +2930,463 @@ impl Parser { ); return Err(Error::InitializationTypeMismatch(name_span, ty)); } - ty - } - None => { - // register the type, if needed - match context.typifier[value].clone() { - TypeResolution::Handle(ty) => ty, - TypeResolution::Value(inner) => context.types.fetch_or_append( - crate::Type { name: None, inner }, - Default::default(), - ), - } - } - }; - - let init = match context.expressions[value] { - crate::Expression::Constant(handle) if is_uniform_control_flow => { - Init::Constant(handle) - } - _ => Init::Variable(value), - }; - (init, ty) - } else { - match given_ty { - Some(ty) => (Init::Empty, ty), - None => { - log::error!("Variable '{}' without an initializer needs a type", name); - return Err(Error::MissingType(name_span)); } + block.extend(emitter.finish(context.expressions)); + context.lookup_ident.insert(name, expr_id); + context + .named_expressions + .insert(expr_id, String::from(name)); + None } - }; + "var" => { + let _ = lexer.next(); + enum Init { + Empty, + Constant(Handle), + Variable(Handle), + } - lexer.expect(Token::Separator(';'))?; - let var_id = context.variables.append( - crate::LocalVariable { - name: Some(name.to_owned()), - ty, - init: match init { - Init::Constant(value) => Some(value), - _ => None, - }, - }, - NagaSpan::ByteRange(name_span), - ); + let (name, name_span) = lexer.next_ident_with_span()?; + let given_ty = if lexer.skip(Token::Separator(':')) { + let (ty, _access) = self.parse_type_decl( + lexer, + None, + context.types, + context.constants, + )?; + Some(ty) + } else { + None + }; - // Doesn't make sense to assign a span to cached lookup - let expr_id = context - .expressions - .append(crate::Expression::LocalVariable(var_id), Default::default()); - context.lookup_ident.insert(name, expr_id); + let (init, ty) = if lexer.skip(Token::Operation('=')) { + emitter.start(context.expressions); + let value = self.parse_general_expression( + lexer, + context.as_expression(block, &mut emitter), + )?; + block.extend(emitter.finish(context.expressions)); - if let Init::Variable(value) = init { - Some(crate::Statement::Store { - pointer: expr_id, - value, - }) - } else { - None - } - } - "return" => { - let _ = lexer.next(); - let value = if lexer.peek().0 != Token::Separator(';') { - emitter.start(context.expressions); - let handle = self.parse_general_expression( - lexer, - context.as_expression(block, &mut emitter), - )?; - block.extend(emitter.finish(context.expressions)); - Some(handle) - } else { - None - }; - lexer.expect(Token::Separator(';'))?; - Some(crate::Statement::Return { value }) - } - "if" => { - let _ = lexer.next(); - lexer.expect(Token::Paren('('))?; - emitter.start(context.expressions); - let condition = self - .parse_general_expression(lexer, context.as_expression(block, &mut emitter))?; - block.extend(emitter.finish(context.expressions)); - lexer.expect(Token::Paren(')'))?; + // prepare the typifier, but work around mutable borrowing... + let _ = context + .as_expression(block, &mut emitter) + .resolve_type(value)?; - let accept = self.parse_block(lexer, context.reborrow(), false)?; - let mut elsif_stack = Vec::new(); - let mut elseif_span_start = lexer.current_byte_offset(); - while lexer.skip(Token::Word("elseif")) { - let mut sub_emitter = super::Emitter::default(); - - lexer.expect(Token::Paren('('))?; - sub_emitter.start(context.expressions); - let other_condition = self.parse_general_expression( - lexer, - context.as_expression(block, &mut sub_emitter), - )?; - let other_emit = sub_emitter.finish(context.expressions); - lexer.expect(Token::Paren(')'))?; - let other_block = self.parse_block(lexer, context.reborrow(), false)?; - elsif_stack.push((elseif_span_start, other_condition, other_emit, other_block)); - elseif_span_start = lexer.current_byte_offset(); - } - let mut reject = if lexer.skip(Token::Word("else")) { - self.parse_block(lexer, context.reborrow(), false)? - } else { - crate::Block::new() - }; - let span_end = lexer.current_byte_offset(); - // reverse-fold the else-if blocks - //Note: we may consider uplifting this to the IR - for (other_span_start, other_cond, other_emit, other_block) in - elsif_stack.into_iter().rev() - { - let sub_stmt = crate::Statement::If { - condition: other_cond, - accept: other_block, - reject, - }; - reject = crate::Block::new(); - reject.extend(other_emit); - reject.push(sub_stmt, NagaSpan::ByteRange(other_span_start..span_end)) - } - - Some(crate::Statement::If { - condition, - accept, - reject, - }) - } - "switch" => { - let _ = lexer.next(); - emitter.start(context.expressions); - lexer.expect(Token::Paren('('))?; - let selector = self - .parse_general_expression(lexer, context.as_expression(block, &mut emitter))?; - lexer.expect(Token::Paren(')'))?; - block.extend(emitter.finish(context.expressions)); - lexer.expect(Token::Paren('{'))?; - let mut cases = Vec::new(); - let mut default = crate::Block::new(); - - loop { - // cases + default - match lexer.next() { - (Token::Word("case"), _) => { - // parse a list of values - let value = loop { - let value = lexer.next_sint_literal()?; - if lexer.skip(Token::Separator(',')) { - if lexer.skip(Token::Separator(':')) { - break value; + //TODO: share more of this code with `let` arm + let ty = match given_ty { + Some(ty) => { + let expr_inner = context.typifier.get(value, context.types); + let given_inner = &context.types[ty].inner; + if given_inner != expr_inner { + log::error!( + "Given type {:?} doesn't match expected {:?}", + given_inner, + expr_inner + ); + return Err(Error::InitializationTypeMismatch( + name_span, ty, + )); + } + ty + } + None => { + // register the type, if needed + match context.typifier[value].clone() { + TypeResolution::Handle(ty) => ty, + TypeResolution::Value(inner) => { + context.types.fetch_or_append( + crate::Type { name: None, inner }, + Default::default(), + ) + } } - } else { - lexer.expect(Token::Separator(':'))?; - break value; } - cases.push(crate::SwitchCase { - value, - body: crate::Block::new(), - fall_through: true, - }); }; - let mut body = crate::Block::new(); - lexer.expect(Token::Paren('{'))?; - let fall_through = loop { - // default statements - if lexer.skip(Token::Word("fallthrough")) { - lexer.expect(Token::Separator(';'))?; - lexer.expect(Token::Paren('}'))?; - break true; + let init = match context.expressions[value] { + crate::Expression::Constant(handle) if is_uniform_control_flow => { + Init::Constant(handle) } - if lexer.skip(Token::Paren('}')) { - break false; - } - self.parse_statement(lexer, context.reborrow(), &mut body, false)?; + _ => Init::Variable(value), }; + (init, ty) + } else { + match given_ty { + Some(ty) => (Init::Empty, ty), + None => { + log::error!( + "Variable '{}' without an initializer needs a type", + name + ); + return Err(Error::MissingType(name_span)); + } + } + }; - cases.push(crate::SwitchCase { + lexer.expect(Token::Separator(';'))?; + let var_id = context.variables.append( + crate::LocalVariable { + name: Some(name.to_owned()), + ty, + init: match init { + Init::Constant(value) => Some(value), + _ => None, + }, + }, + NagaSpan::ByteRange(name_span), + ); + + // Doesn't make sense to assign a span to cached lookup + let expr_id = context + .expressions + .append(crate::Expression::LocalVariable(var_id), Default::default()); + context.lookup_ident.insert(name, expr_id); + + if let Init::Variable(value) = init { + Some(crate::Statement::Store { + pointer: expr_id, value, - body, - fall_through, - }); - } - (Token::Word("default"), _) => { - lexer.expect(Token::Separator(':'))?; - default = self.parse_block(lexer, context.reborrow(), false)?; - } - (Token::Paren('}'), _) => break, - other => return Err(Error::Unexpected(other, ExpectedToken::SwitchItem)), - } - } - - Some(crate::Statement::Switch { - selector, - cases, - default, - }) - } - "loop" => { - let _ = lexer.next(); - let mut body = crate::Block::new(); - let mut continuing = crate::Block::new(); - lexer.expect(Token::Paren('{'))?; - - loop { - if lexer.skip(Token::Word("continuing")) { - continuing = self.parse_block(lexer, context.reborrow(), false)?; - lexer.expect(Token::Paren('}'))?; - break; - } - if lexer.skip(Token::Paren('}')) { - break; - } - self.parse_statement(lexer, context.reborrow(), &mut body, false)?; - } - - Some(crate::Statement::Loop { body, continuing }) - } - "for" => { - let _ = lexer.next(); - lexer.expect(Token::Paren('('))?; - if !lexer.skip(Token::Separator(';')) { - let num_statements = block.len(); - let (_, span) = lexer.capture_span(|lexer| { - self.parse_statement( - lexer, - context.reborrow(), - block, - is_uniform_control_flow, - ) - })?; - - if block.len() != num_statements { - match *block.last().unwrap() { - crate::Statement::Store { .. } | crate::Statement::Call { .. } => {} - _ => return Err(Error::InvalidForInitializer(span)), + }) + } else { + None } } - }; - - let mut body = crate::Block::new(); - if !lexer.skip(Token::Separator(';')) { - let (condition, span) = lexer.capture_span(|lexer| { + "return" => { + let _ = lexer.next(); + let value = if lexer.peek().0 != Token::Separator(';') { + emitter.start(context.expressions); + let handle = self.parse_general_expression( + lexer, + context.as_expression(block, &mut emitter), + )?; + block.extend(emitter.finish(context.expressions)); + Some(handle) + } else { + None + }; + lexer.expect(Token::Separator(';'))?; + Some(crate::Statement::Return { value }) + } + "if" => { + let _ = lexer.next(); + lexer.expect(Token::Paren('('))?; emitter.start(context.expressions); let condition = self.parse_general_expression( lexer, - context.as_expression(&mut body, &mut emitter), + context.as_expression(block, &mut emitter), + )?; + block.extend(emitter.finish(context.expressions)); + lexer.expect(Token::Paren(')'))?; + + let accept = self.parse_block(lexer, context.reborrow(), false)?; + let mut elsif_stack = Vec::new(); + let mut elseif_span_start = lexer.current_byte_offset(); + while lexer.skip(Token::Word("elseif")) { + let mut sub_emitter = super::Emitter::default(); + + lexer.expect(Token::Paren('('))?; + sub_emitter.start(context.expressions); + let other_condition = self.parse_general_expression( + lexer, + context.as_expression(block, &mut sub_emitter), + )?; + let other_emit = sub_emitter.finish(context.expressions); + lexer.expect(Token::Paren(')'))?; + let other_block = self.parse_block(lexer, context.reborrow(), false)?; + elsif_stack.push(( + elseif_span_start, + other_condition, + other_emit, + other_block, + )); + elseif_span_start = lexer.current_byte_offset(); + } + let mut reject = if lexer.skip(Token::Word("else")) { + self.parse_block(lexer, context.reborrow(), false)? + } else { + crate::Block::new() + }; + let span_end = lexer.current_byte_offset(); + // reverse-fold the else-if blocks + //Note: we may consider uplifting this to the IR + for (other_span_start, other_cond, other_emit, other_block) in + elsif_stack.into_iter().rev() + { + let sub_stmt = crate::Statement::If { + condition: other_cond, + accept: other_block, + reject, + }; + reject = crate::Block::new(); + reject.extend(other_emit); + reject.push(sub_stmt, NagaSpan::ByteRange(other_span_start..span_end)) + } + + Some(crate::Statement::If { + condition, + accept, + reject, + }) + } + "switch" => { + let _ = lexer.next(); + emitter.start(context.expressions); + lexer.expect(Token::Paren('('))?; + let selector = self.parse_general_expression( + lexer, + context.as_expression(block, &mut emitter), + )?; + lexer.expect(Token::Paren(')'))?; + block.extend(emitter.finish(context.expressions)); + lexer.expect(Token::Paren('{'))?; + let mut cases = Vec::new(); + let mut default = crate::Block::new(); + + loop { + // cases + default + match lexer.next() { + (Token::Word("case"), _) => { + // parse a list of values + let value = loop { + let value = lexer.next_sint_literal()?; + if lexer.skip(Token::Separator(',')) { + if lexer.skip(Token::Separator(':')) { + break value; + } + } else { + lexer.expect(Token::Separator(':'))?; + break value; + } + cases.push(crate::SwitchCase { + value, + body: crate::Block::new(), + fall_through: true, + }); + }; + + let mut body = crate::Block::new(); + lexer.expect(Token::Paren('{'))?; + let fall_through = loop { + // default statements + if lexer.skip(Token::Word("fallthrough")) { + lexer.expect(Token::Separator(';'))?; + lexer.expect(Token::Paren('}'))?; + break true; + } + if lexer.skip(Token::Paren('}')) { + break false; + } + self.parse_statement( + lexer, + context.reborrow(), + &mut body, + false, + )?; + }; + + cases.push(crate::SwitchCase { + value, + body, + fall_through, + }); + } + (Token::Word("default"), _) => { + lexer.expect(Token::Separator(':'))?; + default = self.parse_block(lexer, context.reborrow(), false)?; + } + (Token::Paren('}'), _) => break, + other => { + return Err(Error::Unexpected(other, ExpectedToken::SwitchItem)) + } + } + } + + Some(crate::Statement::Switch { + selector, + cases, + default, + }) + } + "loop" => { + let _ = lexer.next(); + let mut body = crate::Block::new(); + let mut continuing = crate::Block::new(); + lexer.expect(Token::Paren('{'))?; + + loop { + if lexer.skip(Token::Word("continuing")) { + continuing = self.parse_block(lexer, context.reborrow(), false)?; + lexer.expect(Token::Paren('}'))?; + break; + } + if lexer.skip(Token::Paren('}')) { + break; + } + self.parse_statement(lexer, context.reborrow(), &mut body, false)?; + } + + Some(crate::Statement::Loop { body, continuing }) + } + "for" => { + let _ = lexer.next(); + lexer.expect(Token::Paren('('))?; + if !lexer.skip(Token::Separator(';')) { + let num_statements = block.len(); + let (_, span) = lexer.capture_span(|lexer| { + self.parse_statement( + lexer, + context.reborrow(), + block, + is_uniform_control_flow, + ) + })?; + + if block.len() != num_statements { + match *block.last().unwrap() { + crate::Statement::Store { .. } + | crate::Statement::Call { .. } => {} + _ => return Err(Error::InvalidForInitializer(span)), + } + } + }; + + let mut body = crate::Block::new(); + if !lexer.skip(Token::Separator(';')) { + let (condition, span) = lexer.capture_span(|lexer| { + emitter.start(context.expressions); + let condition = self.parse_general_expression( + lexer, + context.as_expression(&mut body, &mut emitter), + )?; + lexer.expect(Token::Separator(';'))?; + body.extend(emitter.finish(context.expressions)); + Ok(condition) + })?; + let mut reject = crate::Block::new(); + reject.push(crate::Statement::Break, NagaSpan::Unknown); + body.push( + crate::Statement::If { + condition, + accept: crate::Block::new(), + reject, + }, + NagaSpan::ByteRange(span), + ); + }; + + let mut continuing = crate::Block::new(); + if let (Token::Word(ident), ident_span) = lexer.peek() { + // manually parse the next statement here instead of calling parse_statement + // because the statement is not terminated with a semicolon + self.parse_statement_restricted( + lexer, + ident, + ident_span, + context.as_expression(&mut continuing, &mut emitter), + )?; + } + lexer.expect(Token::Paren(')'))?; + lexer.expect(Token::Paren('{'))?; + + while !lexer.skip(Token::Paren('}')) { + self.parse_statement(lexer, context.reborrow(), &mut body, false)?; + } + + Some(crate::Statement::Loop { body, continuing }) + } + "break" => { + let _ = lexer.next(); + Some(crate::Statement::Break) + } + "continue" => { + let _ = lexer.next(); + Some(crate::Statement::Continue) + } + "discard" => { + let _ = lexer.next(); + Some(crate::Statement::Kill) + } + "storageBarrier" => { + let _ = lexer.next(); + lexer.expect(Token::Paren('('))?; + lexer.expect(Token::Paren(')'))?; + Some(crate::Statement::Barrier(crate::Barrier::STORAGE)) + } + "workgroupBarrier" => { + let _ = lexer.next(); + lexer.expect(Token::Paren('('))?; + lexer.expect(Token::Paren(')'))?; + Some(crate::Statement::Barrier(crate::Barrier::WORK_GROUP)) + } + "atomicStore" => { + let _ = lexer.next(); + emitter.start(context.expressions); + lexer.open_arguments()?; + let mut expression_ctx = context.as_expression(block, &mut emitter); + let pointer = + self.parse_atomic_pointer(lexer, expression_ctx.reborrow())?; + lexer.expect(Token::Separator(','))?; + let value = self.parse_general_expression(lexer, expression_ctx)?; + lexer.close_arguments()?; + block.extend(emitter.finish(context.expressions)); + Some(crate::Statement::Store { pointer, value }) + } + "textureStore" => { + let _ = lexer.next(); + emitter.start(context.expressions); + lexer.open_arguments()?; + let (image_name, image_span) = lexer.next_ident_with_span()?; + let image = context + .lookup_ident + .lookup(image_name, image_span.clone())?; + lexer.expect(Token::Separator(','))?; + let mut expr_context = context.as_expression(block, &mut emitter); + let arrayed = match *expr_context.resolve_type(image)? { + crate::TypeInner::Image { arrayed, .. } => arrayed, + _ => return Err(Error::BadTexture(image_span)), + }; + let coordinate = self.parse_general_expression(lexer, expr_context)?; + let array_index = if arrayed { + lexer.expect(Token::Separator(','))?; + Some(self.parse_general_expression( + lexer, + context.as_expression(block, &mut emitter), + )?) + } else { + None + }; + lexer.expect(Token::Separator(','))?; + let value = self.parse_general_expression( + lexer, + context.as_expression(block, &mut emitter), + )?; + lexer.close_arguments()?; + block.extend(emitter.finish(context.expressions)); + Some(crate::Statement::ImageStore { + image, + coordinate, + array_index, + value, + }) + } + // assignment or a function call + ident => { + self.parse_statement_restricted( + lexer, + ident, + word_span, + context.as_expression(block, &mut emitter), )?; lexer.expect(Token::Separator(';'))?; - body.extend(emitter.finish(context.expressions)); - Ok(condition) - })?; - let mut reject = crate::Block::new(); - reject.push(crate::Statement::Break, NagaSpan::Unknown); - body.push( - crate::Statement::If { - condition, - accept: crate::Block::new(), - reject, - }, - NagaSpan::ByteRange(span), - ); + None + } }; - - let mut continuing = crate::Block::new(); - if let (Token::Word(ident), ident_span) = lexer.peek() { - // manually parse the next statement here instead of calling parse_statement - // because the statement is not terminated with a semicolon - self.parse_statement_restricted( - lexer, - ident, - ident_span, - context.as_expression(&mut continuing, &mut emitter), - )?; + let span = self.pop_scope(lexer); + if let Some(statement) = statement { + block.push(statement, NagaSpan::ByteRange(span)); } - lexer.expect(Token::Paren(')'))?; - lexer.expect(Token::Paren('{'))?; - - while !lexer.skip(Token::Paren('}')) { - self.parse_statement(lexer, context.reborrow(), &mut body, false)?; - } - - Some(crate::Statement::Loop { body, continuing }) } - "break" => { - let _ = lexer.next(); - Some(crate::Statement::Break) - } - "continue" => { - let _ = lexer.next(); - Some(crate::Statement::Continue) - } - "discard" => { - let _ = lexer.next(); - Some(crate::Statement::Kill) - } - "storageBarrier" => { - let _ = lexer.next(); - lexer.expect(Token::Paren('('))?; - lexer.expect(Token::Paren(')'))?; - Some(crate::Statement::Barrier(crate::Barrier::STORAGE)) - } - "workgroupBarrier" => { - let _ = lexer.next(); - lexer.expect(Token::Paren('('))?; - lexer.expect(Token::Paren(')'))?; - Some(crate::Statement::Barrier(crate::Barrier::WORK_GROUP)) - } - "atomicStore" => { - let _ = lexer.next(); - emitter.start(context.expressions); - lexer.open_arguments()?; - let mut expression_ctx = context.as_expression(block, &mut emitter); - let pointer = self.parse_atomic_pointer(lexer, expression_ctx.reborrow())?; - lexer.expect(Token::Separator(','))?; - let value = self.parse_general_expression(lexer, expression_ctx)?; - lexer.close_arguments()?; - block.extend(emitter.finish(context.expressions)); - Some(crate::Statement::Store { pointer, value }) - } - "textureStore" => { - let _ = lexer.next(); - emitter.start(context.expressions); - lexer.open_arguments()?; - let (image_name, image_span) = lexer.next_ident_with_span()?; - let image = context - .lookup_ident - .lookup(image_name, image_span.clone())?; - lexer.expect(Token::Separator(','))?; - let mut expr_context = context.as_expression(block, &mut emitter); - let arrayed = match *expr_context.resolve_type(image)? { - crate::TypeInner::Image { arrayed, .. } => arrayed, - _ => return Err(Error::BadTexture(image_span)), - }; - let coordinate = self.parse_general_expression(lexer, expr_context)?; - let array_index = if arrayed { - lexer.expect(Token::Separator(','))?; - Some(self.parse_general_expression( - lexer, - context.as_expression(block, &mut emitter), - )?) - } else { - None - }; - lexer.expect(Token::Separator(','))?; - let value = self - .parse_general_expression(lexer, context.as_expression(block, &mut emitter))?; - lexer.close_arguments()?; - block.extend(emitter.finish(context.expressions)); - Some(crate::Statement::ImageStore { - image, - coordinate, - array_index, - value, - }) - } - // assignment or a function call - ident => { - self.parse_statement_restricted( - lexer, - ident, - word_span, - context.as_expression(block, &mut emitter), - )?; - lexer.expect(Token::Separator(';'))?; - None - } - }; - let span = self.pop_scope(lexer); - if let Some(statement) = statement { - block.push(statement, NagaSpan::ByteRange(span)); + other => return Err(Error::Unexpected(other, ExpectedToken::Statement)), } Ok(()) } diff --git a/tests/out/wgsl/246-collatz-comp.wgsl b/tests/out/wgsl/246-collatz-comp.wgsl index 21fde3265c..a0ee765350 100644 --- a/tests/out/wgsl/246-collatz-comp.wgsl +++ b/tests/out/wgsl/246-collatz-comp.wgsl @@ -4,7 +4,7 @@ struct PrimeIndices { }; [[group(0), binding(0)]] -var global: PrimeIndices; +var global: PrimeIndices; var gl_GlobalInvocationID: vec3; fn collatz_iterations(n: u32) -> u32 { diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 7706cda4fb..d9a00535a1 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -7,7 +7,7 @@ struct Bar { }; [[group(0), binding(0)]] -var bar: Bar; +var bar: Bar; [[stage(vertex)]] fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { diff --git a/tests/out/wgsl/boids.wgsl b/tests/out/wgsl/boids.wgsl index 4546f1721c..283169b58f 100644 --- a/tests/out/wgsl/boids.wgsl +++ b/tests/out/wgsl/boids.wgsl @@ -26,7 +26,7 @@ var params: SimParams; [[group(0), binding(1)]] var particlesSrc: Particles; [[group(0), binding(2)]] -var particlesDst: Particles; +var particlesDst: Particles; [[stage(compute), workgroup_size(64, 1, 1)]] fn main([[builtin(global_invocation_id)]] global_invocation_id: vec3) { diff --git a/tests/out/wgsl/collatz.wgsl b/tests/out/wgsl/collatz.wgsl index 7d7a8dedc6..e3bbb2cabd 100644 --- a/tests/out/wgsl/collatz.wgsl +++ b/tests/out/wgsl/collatz.wgsl @@ -4,7 +4,7 @@ struct PrimeIndices { }; [[group(0), binding(0)]] -var v_indices: PrimeIndices; +var v_indices: PrimeIndices; fn collatz_iterations(n_base: u32) -> u32 { var n: u32;