From 8fc60fc07a40e0aecf518d7b879d9e61a3b2304b Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Tue, 12 Dec 2023 14:05:26 -0800 Subject: [PATCH] [naga wgsl] Let unary operators accept and produce abstract types. Fixes #4445. Fixes #4492. Fixes #4435. --- CHANGELOG.md | 2 +- naga/src/back/wgsl/writer.rs | 11 +- naga/src/front/wgsl/lower/mod.rs | 2 +- naga/src/front/wgsl/parse/lexer.rs | 103 ++++------ naga/src/front/wgsl/parse/number.rs | 59 ++---- naga/src/proc/constant_evaluator.rs | 5 +- naga/tests/in/abstract-types-operators.wgsl | 24 +++ .../out/msl/abstract-types-operators.msl | 20 ++ .../out/spv/abstract-types-operators.spvasm | 190 ++++++++++-------- .../out/wgsl/abstract-types-operators.wgsl | 16 ++ naga/tests/wgsl_errors.rs | 15 ++ 11 files changed, 259 insertions(+), 188 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index acb6c87a1b..909ff8ad40 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -68,7 +68,7 @@ This feature allowed you to call `global_id` on any wgpu opaque handle to get a #### Naga -- Naga's WGSL front end now allows binary operators to produce values with abstract types, rather than concretizing thir operands. By @jimblandy in [#4850](https://github.com/gfx-rs/wgpu/pull/4850). +- Naga's WGSL front end now allows operators to produce values with abstract types, rather than concretizing thir operands. By @jimblandy in [#4850](https://github.com/gfx-rs/wgpu/pull/4850) and [#4870](https://github.com/gfx-rs/wgpu/pull/4870). - 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). diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index 25a6979515..c737934f5e 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -1090,7 +1090,16 @@ impl Writer { Expression::Literal(literal) => match literal { crate::Literal::F32(value) => write!(self.out, "{}f", value)?, crate::Literal::U32(value) => write!(self.out, "{}u", value)?, - crate::Literal::I32(value) => write!(self.out, "{}i", value)?, + crate::Literal::I32(value) => { + // `-2147483648i` is not valid WGSL. The most negative `i32` + // value can only be expressed in WGSL using AbstractInt and + // a unary negation operator. + if value == i32::MIN { + write!(self.out, "i32(-2147483648)")?; + } else { + write!(self.out, "{}i", value)?; + } + } crate::Literal::Bool(value) => write!(self.out, "{}", value)?, crate::Literal::F64(value) => write!(self.out, "{:?}lf", value)?, crate::Literal::I64(_) => { diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index a4a49da061..c38c57267a 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1574,7 +1574,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { return Ok(Typed::Plain(handle)); } ast::Expression::Unary { op, expr } => { - let expr = self.expression(expr, ctx)?; + let expr = self.expression_for_abstract(expr, ctx)?; Typed::Plain(crate::Expression::Unary { op, expr }) } ast::Expression::AddrOf(expr) => { diff --git a/naga/src/front/wgsl/parse/lexer.rs b/naga/src/front/wgsl/parse/lexer.rs index 9fa9416f11..d03a448561 100644 --- a/naga/src/front/wgsl/parse/lexer.rs +++ b/naga/src/front/wgsl/parse/lexer.rs @@ -117,7 +117,6 @@ fn consume_token(input: &str, generic: bool) -> (Token<'_>, &str) { let og_chars = chars.as_str(); match chars.next() { Some('>') => (Token::Arrow, chars.as_str()), - Some('0'..='9' | '.') => consume_number(input), Some('-') => (Token::DecrementOperation, chars.as_str()), Some('=') => (Token::AssignmentOperation(cur), chars.as_str()), _ => (Token::Operation(cur), og_chars), @@ -496,44 +495,60 @@ fn test_numbers() { // MIN / MAX // - // min / max decimal signed integer + // min / max decimal integer sub_test( - "-2147483648i 2147483647i -2147483649i 2147483648i", + "0i 2147483647i 2147483648i", &[ - Token::Number(Ok(Number::I32(i32::MIN))), + Token::Number(Ok(Number::I32(0))), Token::Number(Ok(Number::I32(i32::MAX))), Token::Number(Err(NumberError::NotRepresentable)), - Token::Number(Err(NumberError::NotRepresentable)), ], ); // min / max decimal unsigned integer sub_test( - "0u 4294967295u -1u 4294967296u", + "0u 4294967295u 4294967296u", &[ Token::Number(Ok(Number::U32(u32::MIN))), Token::Number(Ok(Number::U32(u32::MAX))), Token::Number(Err(NumberError::NotRepresentable)), - Token::Number(Err(NumberError::NotRepresentable)), ], ); // min / max hexadecimal signed integer sub_test( - "-0x80000000i 0x7FFFFFFFi -0x80000001i 0x80000000i", + "0x0i 0x7FFFFFFFi 0x80000000i", &[ - Token::Number(Ok(Number::I32(i32::MIN))), + Token::Number(Ok(Number::I32(0))), Token::Number(Ok(Number::I32(i32::MAX))), Token::Number(Err(NumberError::NotRepresentable)), - Token::Number(Err(NumberError::NotRepresentable)), ], ); // min / max hexadecimal unsigned integer sub_test( - "0x0u 0xFFFFFFFFu -0x1u 0x100000000u", + "0x0u 0xFFFFFFFFu 0x100000000u", &[ Token::Number(Ok(Number::U32(u32::MIN))), Token::Number(Ok(Number::U32(u32::MAX))), Token::Number(Err(NumberError::NotRepresentable)), + ], + ); + + // min/max decimal abstract int + sub_test( + "0 9223372036854775807 9223372036854775808", + &[ + Token::Number(Ok(Number::AbstractInt(0))), + Token::Number(Ok(Number::AbstractInt(i64::MAX))), + Token::Number(Err(NumberError::NotRepresentable)), + ], + ); + + // min/max hexadecimal abstract int + sub_test( + "0 0x7fffffffffffffff 0x8000000000000000", + &[ + Token::Number(Ok(Number::AbstractInt(0))), + Token::Number(Ok(Number::AbstractInt(i64::MAX))), Token::Number(Err(NumberError::NotRepresentable)), ], ); @@ -548,77 +563,43 @@ fn test_numbers() { const LARGEST_F32_LESS_THAN_ONE: f32 = 0.99999994; /// ≈ 1 + 2^−23 const SMALLEST_F32_LARGER_THAN_ONE: f32 = 1.0000001; - /// ≈ -(2^127 * (2 − 2^−23)) - const SMALLEST_NORMAL_F32: f32 = f32::MIN; /// ≈ 2^127 * (2 − 2^−23) const LARGEST_NORMAL_F32: f32 = f32::MAX; // decimal floating point sub_test( - "1e-45f 1.1754942e-38f 1.17549435e-38f 0.99999994f 1.0000001f -3.40282347e+38f 3.40282347e+38f", + "1e-45f 1.1754942e-38f 1.17549435e-38f 0.99999994f 1.0000001f 3.40282347e+38f", &[ - Token::Number(Ok(Number::F32( - SMALLEST_POSITIVE_SUBNORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - LARGEST_SUBNORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - SMALLEST_POSITIVE_NORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - LARGEST_F32_LESS_THAN_ONE, - ))), - Token::Number(Ok(Number::F32( - SMALLEST_F32_LARGER_THAN_ONE, - ))), - Token::Number(Ok(Number::F32( - SMALLEST_NORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - LARGEST_NORMAL_F32, - ))), + Token::Number(Ok(Number::F32(SMALLEST_POSITIVE_SUBNORMAL_F32))), + Token::Number(Ok(Number::F32(LARGEST_SUBNORMAL_F32))), + Token::Number(Ok(Number::F32(SMALLEST_POSITIVE_NORMAL_F32))), + Token::Number(Ok(Number::F32(LARGEST_F32_LESS_THAN_ONE))), + Token::Number(Ok(Number::F32(SMALLEST_F32_LARGER_THAN_ONE))), + Token::Number(Ok(Number::F32(LARGEST_NORMAL_F32))), ], ); sub_test( - "-3.40282367e+38f 3.40282367e+38f", + "3.40282367e+38f", &[ - Token::Number(Err(NumberError::NotRepresentable)), // ≈ -2^128 Token::Number(Err(NumberError::NotRepresentable)), // ≈ 2^128 ], ); // hexadecimal floating point sub_test( - "0x1p-149f 0x7FFFFFp-149f 0x1p-126f 0xFFFFFFp-24f 0x800001p-23f -0xFFFFFFp+104f 0xFFFFFFp+104f", + "0x1p-149f 0x7FFFFFp-149f 0x1p-126f 0xFFFFFFp-24f 0x800001p-23f 0xFFFFFFp+104f", &[ - Token::Number(Ok(Number::F32( - SMALLEST_POSITIVE_SUBNORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - LARGEST_SUBNORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - SMALLEST_POSITIVE_NORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - LARGEST_F32_LESS_THAN_ONE, - ))), - Token::Number(Ok(Number::F32( - SMALLEST_F32_LARGER_THAN_ONE, - ))), - Token::Number(Ok(Number::F32( - SMALLEST_NORMAL_F32, - ))), - Token::Number(Ok(Number::F32( - LARGEST_NORMAL_F32, - ))), + Token::Number(Ok(Number::F32(SMALLEST_POSITIVE_SUBNORMAL_F32))), + Token::Number(Ok(Number::F32(LARGEST_SUBNORMAL_F32))), + Token::Number(Ok(Number::F32(SMALLEST_POSITIVE_NORMAL_F32))), + Token::Number(Ok(Number::F32(LARGEST_F32_LESS_THAN_ONE))), + Token::Number(Ok(Number::F32(SMALLEST_F32_LARGER_THAN_ONE))), + Token::Number(Ok(Number::F32(LARGEST_NORMAL_F32))), ], ); sub_test( - "-0x1p128f 0x1p128f 0x1.000001p0f", + "0x1p128f 0x1.000001p0f", &[ - Token::Number(Err(NumberError::NotRepresentable)), // = -2^128 Token::Number(Err(NumberError::NotRepresentable)), // = 2^128 Token::Number(Err(NumberError::NotRepresentable)), ], diff --git a/naga/src/front/wgsl/parse/number.rs b/naga/src/front/wgsl/parse/number.rs index fde5e3cee6..7b09ac59bb 100644 --- a/naga/src/front/wgsl/parse/number.rs +++ b/naga/src/front/wgsl/parse/number.rs @@ -1,5 +1,3 @@ -use std::borrow::Cow; - use crate::front::wgsl::error::NumberError; use crate::front::wgsl::parse::lexer::Token; @@ -20,8 +18,6 @@ pub enum Number { F64(f64), } -// TODO: when implementing Creation-Time Expressions, remove the ability to match the minus sign - pub(in crate::front::wgsl) fn consume_number(input: &str) -> (Token<'_>, &str) { let (result, rest) = parse(input); (Token::Number(result), rest) @@ -64,7 +60,9 @@ enum FloatKind { // | / 0[xX][0-9a-fA-F]+ [pP][+-]?[0-9]+ [fh]? / // You could visualize the regex below via https://debuggex.com to get a rough idea what `parse` is doing -// -?(?:0[xX](?:([0-9a-fA-F]+\.[0-9a-fA-F]*|[0-9a-fA-F]*\.[0-9a-fA-F]+)(?:([pP][+-]?[0-9]+)([fh]?))?|([0-9a-fA-F]+)([pP][+-]?[0-9]+)([fh]?)|([0-9a-fA-F]+)([iu]?))|((?:[0-9]+[eE][+-]?[0-9]+|(?:[0-9]+\.[0-9]*|[0-9]*\.[0-9]+)(?:[eE][+-]?[0-9]+)?))([fh]?)|((?:[0-9]|[1-9][0-9]+))([iufh]?)) +// (?:0[xX](?:([0-9a-fA-F]+\.[0-9a-fA-F]*|[0-9a-fA-F]*\.[0-9a-fA-F]+)(?:([pP][+-]?[0-9]+)([fh]?))?|([0-9a-fA-F]+)([pP][+-]?[0-9]+)([fh]?)|([0-9a-fA-F]+)([iu]?))|((?:[0-9]+[eE][+-]?[0-9]+|(?:[0-9]+\.[0-9]*|[0-9]*\.[0-9]+)(?:[eE][+-]?[0-9]+)?))([fh]?)|((?:[0-9]|[1-9][0-9]+))([iufh]?)) + +// Leading signs are handled as unary operators. fn parse(input: &str) -> (Result, &str) { /// returns `true` and consumes `X` bytes from the given byte buffer @@ -152,8 +150,6 @@ fn parse(input: &str) -> (Result, &str) { let general_extract = ExtractSubStr::start(input, bytes); - let is_negative = consume!(bytes, b'-'); - if consume!(bytes, b'0', b'x' | b'X') { let digits_extract = ExtractSubStr::start(input, bytes); @@ -216,10 +212,7 @@ fn parse(input: &str) -> (Result, &str) { } else { let kind = consume_map!(bytes, [b'i' => IntKind::I32, b'u' => IntKind::U32]); - ( - parse_hex_int(is_negative, digits, kind), - rest_to_str!(bytes), - ) + (parse_hex_int(digits, kind), rest_to_str!(bytes)) } } } else { @@ -272,7 +265,7 @@ fn parse(input: &str) -> (Result, &str) { return (Err(NumberError::Invalid), rest_to_str!(bytes)); } - let digits_with_sign = general_extract.end(bytes); + let digits = general_extract.end(bytes); let kind = consume_map!(bytes, [ b'i' => Kind::Int(IntKind::I32), @@ -282,17 +275,14 @@ fn parse(input: &str) -> (Result, &str) { b'l', b'f' => Kind::Float(FloatKind::F64), ]); - ( - parse_dec(is_negative, digits_with_sign, kind), - rest_to_str!(bytes), - ) + (parse_dec(digits, kind), rest_to_str!(bytes)) } } } } fn parse_hex_float_missing_exponent( - // format: -?0[xX] ( [0-9a-fA-F]+\.[0-9a-fA-F]* | [0-9a-fA-F]*\.[0-9a-fA-F]+ ) + // format: 0[xX] ( [0-9a-fA-F]+\.[0-9a-fA-F]* | [0-9a-fA-F]*\.[0-9a-fA-F]+ ) significand: &str, kind: Option, ) -> Result { @@ -301,7 +291,7 @@ fn parse_hex_float_missing_exponent( } fn parse_hex_float_missing_period( - // format: -?0[xX] [0-9a-fA-F]+ + // format: 0[xX] [0-9a-fA-F]+ significand: &str, // format: [pP][+-]?[0-9]+ exponent: &str, @@ -312,29 +302,22 @@ fn parse_hex_float_missing_period( } fn parse_hex_int( - is_negative: bool, // format: [0-9a-fA-F]+ digits: &str, kind: Option, ) -> Result { - let digits_with_sign = if is_negative { - Cow::Owned(format!("-{digits}")) - } else { - Cow::Borrowed(digits) - }; - parse_int(&digits_with_sign, kind, 16, is_negative) + parse_int(digits, kind, 16) } fn parse_dec( - is_negative: bool, - // format: -? ( [0-9] | [1-9][0-9]+ ) - digits_with_sign: &str, + // format: ( [0-9] | [1-9][0-9]+ ) + digits: &str, kind: Option, ) -> Result { match kind { - None => parse_int(digits_with_sign, None, 10, is_negative), - Some(Kind::Int(kind)) => parse_int(digits_with_sign, Some(kind), 10, is_negative), - Some(Kind::Float(kind)) => parse_dec_float(digits_with_sign, Some(kind)), + None => parse_int(digits, None, 10), + Some(Kind::Int(kind)) => parse_int(digits, Some(kind), 10), + Some(Kind::Float(kind)) => parse_dec_float(digits, Some(kind)), } } @@ -363,7 +346,7 @@ fn parse_dec( // Therefore we only check for overflow manually for decimal floating point literals -// input format: -?0[xX] ( [0-9a-fA-F]+\.[0-9a-fA-F]* | [0-9a-fA-F]*\.[0-9a-fA-F]+ ) [pP][+-]?[0-9]+ +// input format: 0[xX] ( [0-9a-fA-F]+\.[0-9a-fA-F]* | [0-9a-fA-F]*\.[0-9a-fA-F]+ ) [pP][+-]?[0-9]+ fn parse_hex_float(input: &str, kind: Option) -> Result { match kind { None => match hexf_parse::parse_hexf64(input, false) { @@ -385,8 +368,8 @@ fn parse_hex_float(input: &str, kind: Option) -> Result) -> Result { match kind { None => { @@ -411,12 +394,7 @@ fn parse_dec_float(input: &str, kind: Option) -> Result, - radix: u32, - is_negative: bool, -) -> Result { +fn parse_int(input: &str, kind: Option, radix: u32) -> Result { fn map_err(e: core::num::ParseIntError) -> NumberError { match *e.kind() { core::num::IntErrorKind::PosOverflow | core::num::IntErrorKind::NegOverflow => { @@ -434,7 +412,6 @@ fn parse_int( Ok(num) => Ok(Number::I32(num)), Err(e) => Err(map_err(e)), }, - Some(IntKind::U32) if is_negative => Err(NumberError::NotRepresentable), Some(IntKind::U32) => match u32::from_str_radix(input, radix) { Ok(num) => Ok(Number::U32(num)), Err(e) => Err(map_err(e)), diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 81bcb35e58..82efeece4e 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1191,8 +1191,10 @@ impl<'a> ConstantEvaluator<'a> { let expr = match self.expressions[expr] { Expression::Literal(value) => Expression::Literal(match op { UnaryOperator::Negate => match value { - Literal::I32(v) => Literal::I32(-v), + Literal::I32(v) => Literal::I32(v.wrapping_neg()), Literal::F32(v) => Literal::F32(-v), + Literal::AbstractInt(v) => Literal::AbstractInt(v.wrapping_neg()), + Literal::AbstractFloat(v) => Literal::AbstractFloat(-v), _ => return Err(ConstantEvaluatorError::InvalidUnaryOpArg), }, UnaryOperator::LogicalNot => match value { @@ -1202,6 +1204,7 @@ impl<'a> ConstantEvaluator<'a> { UnaryOperator::BitwiseNot => match value { Literal::I32(v) => Literal::I32(!v), Literal::U32(v) => Literal::U32(!v), + Literal::AbstractInt(v) => Literal::AbstractInt(!v), _ => return Err(ConstantEvaluatorError::InvalidUnaryOpArg), }, }), diff --git a/naga/tests/in/abstract-types-operators.wgsl b/naga/tests/in/abstract-types-operators.wgsl index 2c06ef43e4..25f9236426 100644 --- a/naga/tests/in/abstract-types-operators.wgsl +++ b/naga/tests/in/abstract-types-operators.wgsl @@ -18,6 +18,12 @@ const plus_uai_u: u32 = 1 + 2u; const plus_u_uai: u32 = 1u + 2; const plus_u_u_u: u32 = 1u + 2u; +const bitflip_u_u: u32 = ~0xffffffffu; +const bitflip_uai: u32 = ~0xffffffff & (0x100000000 - 1); + +const least_i32: i32 = -2147483648; +const least_f32: f32 = -3.40282347e+38; + fn runtime_values() { var f: f32 = 42; var i: i32 = 43; @@ -43,3 +49,21 @@ fn runtime_values() { var plus_u_uai: u32 = u + 2; var plus_u_u_u: u32 = u + u; } + +fn wgpu_4445() { + // This ok: + let a = (3.0*2.0-(1.0)) * 1.0; + let b = (3.0*2.0+1.0) * 1.0; + // This fails: + let c = (3.0*2.0-1.0) * 1.0; +} + +const wgpu_4492 = i32(-0x80000000); +const wgpu_4492_2 = -2147483648; + +var a: array; + +fn wgpu_4435() { + let x = 1; + let y = a[x-1]; +} diff --git a/naga/tests/out/msl/abstract-types-operators.msl b/naga/tests/out/msl/abstract-types-operators.msl index b018573152..f273b06610 100644 --- a/naga/tests/out/msl/abstract-types-operators.msl +++ b/naga/tests/out/msl/abstract-types-operators.msl @@ -4,6 +4,9 @@ using metal::uint; +struct type_3 { + uint inner[64]; +}; constant float plus_fafaf_1 = 3.0; constant float plus_fafai_1 = 3.0; constant float plus_faf_f_1 = 3.0; @@ -21,6 +24,12 @@ constant uint plus_uaiai_1 = 3u; constant uint plus_uai_u_1 = 3u; constant uint plus_u_uai_1 = 3u; constant uint plus_u_u_u_1 = 3u; +constant uint bitflip_u_u = 0u; +constant uint bitflip_uai = 0u; +constant int least_i32_ = -2147483648; +constant float least_f32_ = -340282350000000000000000000000000000000.0; +constant int wgpu_4492_ = -2147483648; +constant int wgpu_4492_2_ = -2147483648; void runtime_values( ) { @@ -71,3 +80,14 @@ void runtime_values( plus_u_u_u = _e52 + _e53; return; } + +void wgpu_4445_( +) { + return; +} + +void wgpu_4435_( + threadgroup type_3& a +) { + uint y = a.inner[1 - 1]; +} diff --git a/naga/tests/out/spv/abstract-types-operators.spvasm b/naga/tests/out/spv/abstract-types-operators.spvasm index ea9f858295..d1ecc56397 100644 --- a/naga/tests/out/spv/abstract-types-operators.spvasm +++ b/naga/tests/out/spv/abstract-types-operators.spvasm @@ -1,101 +1,127 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 81 +; Bound: 100 OpCapability Shader OpCapability Linkage %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 +OpDecorate %6 ArrayStride 4 %2 = OpTypeVoid %3 = OpTypeFloat 32 %4 = OpTypeInt 32 1 %5 = OpTypeInt 32 0 -%6 = OpConstant %3 3.0 -%7 = OpConstant %4 3 -%8 = OpConstant %5 3 -%11 = OpTypeFunction %2 -%12 = OpConstant %3 42.0 -%13 = OpConstant %4 43 -%14 = OpConstant %5 44 -%15 = OpConstant %3 1.0 -%16 = OpConstant %3 2.0 -%17 = OpConstant %4 1 -%18 = OpConstant %4 2 -%19 = OpConstant %5 1 -%20 = OpConstant %5 2 -%22 = OpTypePointer Function %3 -%24 = OpTypePointer Function %4 -%26 = OpTypePointer Function %5 -%30 = OpConstantNull %3 -%34 = OpConstantNull %3 -%36 = OpConstantNull %3 -%38 = OpConstantNull %3 -%40 = OpConstantNull %3 -%43 = OpConstantNull %4 -%45 = OpConstantNull %4 -%47 = OpConstantNull %4 -%50 = OpConstantNull %5 -%52 = OpConstantNull %5 -%54 = OpConstantNull %5 -%10 = OpFunction %2 None %11 -%9 = OpLabel -%53 = OpVariable %26 Function %54 -%48 = OpVariable %26 Function %8 -%42 = OpVariable %24 Function %43 -%37 = OpVariable %22 Function %38 -%32 = OpVariable %22 Function %6 -%28 = OpVariable %22 Function %6 -%23 = OpVariable %24 Function %13 -%51 = OpVariable %26 Function %52 -%46 = OpVariable %24 Function %47 -%41 = OpVariable %24 Function %7 -%35 = OpVariable %22 Function %36 -%31 = OpVariable %22 Function %6 -%27 = OpVariable %22 Function %6 -%21 = OpVariable %22 Function %12 -%49 = OpVariable %26 Function %50 -%44 = OpVariable %24 Function %45 -%39 = OpVariable %22 Function %40 -%33 = OpVariable %22 Function %34 -%29 = OpVariable %22 Function %30 -%25 = OpVariable %26 Function %14 -OpBranch %55 -%55 = OpLabel -%56 = OpLoad %3 %21 -%57 = OpFAdd %3 %15 %56 -OpStore %29 %57 -%58 = OpLoad %3 %21 -%59 = OpFAdd %3 %15 %58 -OpStore %33 %59 -%60 = OpLoad %3 %21 -%61 = OpFAdd %3 %60 %16 -OpStore %35 %61 -%62 = OpLoad %3 %21 -%63 = OpFAdd %3 %62 %16 -OpStore %37 %63 -%64 = OpLoad %3 %21 -%65 = OpLoad %3 %21 -%66 = OpFAdd %3 %64 %65 -OpStore %39 %66 -%67 = OpLoad %4 %23 -%68 = OpIAdd %4 %17 %67 +%7 = OpConstant %5 64 +%6 = OpTypeArray %5 %7 +%8 = OpConstant %3 3.0 +%9 = OpConstant %4 3 +%10 = OpConstant %5 3 +%11 = OpConstant %5 0 +%12 = OpConstant %4 -2147483648 +%13 = OpConstant %3 -3.4028235e38 +%15 = OpTypePointer Workgroup %6 +%14 = OpVariable %15 Workgroup +%18 = OpTypeFunction %2 +%19 = OpConstant %3 42.0 +%20 = OpConstant %4 43 +%21 = OpConstant %5 44 +%22 = OpConstant %3 1.0 +%23 = OpConstant %3 2.0 +%24 = OpConstant %4 1 +%25 = OpConstant %4 2 +%26 = OpConstant %5 1 +%27 = OpConstant %5 2 +%29 = OpTypePointer Function %3 +%31 = OpTypePointer Function %4 +%33 = OpTypePointer Function %5 +%37 = OpConstantNull %3 +%41 = OpConstantNull %3 +%43 = OpConstantNull %3 +%45 = OpConstantNull %3 +%47 = OpConstantNull %3 +%50 = OpConstantNull %4 +%52 = OpConstantNull %4 +%54 = OpConstantNull %4 +%57 = OpConstantNull %5 +%59 = OpConstantNull %5 +%61 = OpConstantNull %5 +%90 = OpConstant %3 5.0 +%91 = OpConstant %3 7.0 +%97 = OpTypePointer Workgroup %5 +%17 = OpFunction %2 None %18 +%16 = OpLabel +%60 = OpVariable %33 Function %61 +%55 = OpVariable %33 Function %10 +%49 = OpVariable %31 Function %50 +%44 = OpVariable %29 Function %45 +%39 = OpVariable %29 Function %8 +%35 = OpVariable %29 Function %8 +%30 = OpVariable %31 Function %20 +%58 = OpVariable %33 Function %59 +%53 = OpVariable %31 Function %54 +%48 = OpVariable %31 Function %9 +%42 = OpVariable %29 Function %43 +%38 = OpVariable %29 Function %8 +%34 = OpVariable %29 Function %8 +%28 = OpVariable %29 Function %19 +%56 = OpVariable %33 Function %57 +%51 = OpVariable %31 Function %52 +%46 = OpVariable %29 Function %47 +%40 = OpVariable %29 Function %41 +%36 = OpVariable %29 Function %37 +%32 = OpVariable %33 Function %21 +OpBranch %62 +%62 = OpLabel +%63 = OpLoad %3 %28 +%64 = OpFAdd %3 %22 %63 +OpStore %36 %64 +%65 = OpLoad %3 %28 +%66 = OpFAdd %3 %22 %65 +OpStore %40 %66 +%67 = OpLoad %3 %28 +%68 = OpFAdd %3 %67 %23 OpStore %42 %68 -%69 = OpLoad %4 %23 -%70 = OpIAdd %4 %69 %18 +%69 = OpLoad %3 %28 +%70 = OpFAdd %3 %69 %23 OpStore %44 %70 -%71 = OpLoad %4 %23 -%72 = OpLoad %4 %23 -%73 = OpIAdd %4 %71 %72 +%71 = OpLoad %3 %28 +%72 = OpLoad %3 %28 +%73 = OpFAdd %3 %71 %72 OpStore %46 %73 -%74 = OpLoad %5 %25 -%75 = OpIAdd %5 %19 %74 +%74 = OpLoad %4 %30 +%75 = OpIAdd %4 %24 %74 OpStore %49 %75 -%76 = OpLoad %5 %25 -%77 = OpIAdd %5 %76 %20 +%76 = OpLoad %4 %30 +%77 = OpIAdd %4 %76 %25 OpStore %51 %77 -%78 = OpLoad %5 %25 -%79 = OpLoad %5 %25 -%80 = OpIAdd %5 %78 %79 +%78 = OpLoad %4 %30 +%79 = OpLoad %4 %30 +%80 = OpIAdd %4 %78 %79 OpStore %53 %80 +%81 = OpLoad %5 %32 +%82 = OpIAdd %5 %26 %81 +OpStore %56 %82 +%83 = OpLoad %5 %32 +%84 = OpIAdd %5 %83 %27 +OpStore %58 %84 +%85 = OpLoad %5 %32 +%86 = OpLoad %5 %32 +%87 = OpIAdd %5 %85 %86 +OpStore %60 %87 +OpReturn +OpFunctionEnd +%89 = OpFunction %2 None %18 +%88 = OpLabel +OpBranch %92 +%92 = OpLabel +OpReturn +OpFunctionEnd +%94 = OpFunction %2 None %18 +%93 = OpLabel +OpBranch %95 +%95 = OpLabel +%96 = OpISub %4 %24 %24 +%98 = OpAccessChain %97 %14 %96 +%99 = OpLoad %5 %98 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/abstract-types-operators.wgsl b/naga/tests/out/wgsl/abstract-types-operators.wgsl index a0278712ad..de2b073074 100644 --- a/naga/tests/out/wgsl/abstract-types-operators.wgsl +++ b/naga/tests/out/wgsl/abstract-types-operators.wgsl @@ -15,6 +15,14 @@ const plus_uaiai_1: u32 = 3u; const plus_uai_u_1: u32 = 3u; const plus_u_uai_1: u32 = 3u; const plus_u_u_u_1: u32 = 3u; +const bitflip_u_u: u32 = 0u; +const bitflip_uai: u32 = 0u; +const least_i32_: i32 = i32(-2147483648); +const least_f32_: f32 = -340282350000000000000000000000000000000f; +const wgpu_4492_: i32 = i32(-2147483648); +const wgpu_4492_2_: i32 = i32(-2147483648); + +var a: array; fn runtime_values() { var f: f32 = 42f; @@ -66,3 +74,11 @@ fn runtime_values() { return; } +fn wgpu_4445_() { + return; +} + +fn wgpu_4435_() { + let y = a[(1i - 1i)]; +} + diff --git a/naga/tests/wgsl_errors.rs b/naga/tests/wgsl_errors.rs index 7dddf633a4..2f62491b3f 100644 --- a/naga/tests/wgsl_errors.rs +++ b/naga/tests/wgsl_errors.rs @@ -19,6 +19,21 @@ fn check(input: &str, snapshot: &str) { } } +#[test] +fn very_negative_integers() { + // wgpu#4492 + check( + "const i32min = -0x80000000i;", + r###"error: numeric literal not representable by target type: `0x80000000i` + ┌─ wgsl:1:17 + │ +1 │ const i32min = -0x80000000i; + │ ^^^^^^^^^^^ numeric literal not representable by target type + +"###, + ); +} + #[test] fn reserved_identifier_prefix() { check(