diff --git a/CHANGELOG.md b/CHANGELOG.md index d0d5f82cb..c0da41d01 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -98,6 +98,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148] - Fix an issue where `naga` CLI would incorrectly skip the first positional argument when `--stdin-file-path` was specified. By @ErichDonGubler in [#6480](https://github.com/gfx-rs/wgpu/pull/6480). - Fix textureNumLevels in the GLSL backend. By @magcius in [#6483](https://github.com/gfx-rs/wgpu/pull/6483). +- Support 64-bit hex literals and unary operations in constants [#6616](https://github.com/gfx-rs/wgpu/pull/6616). - Implement `quantizeToF16()` for WGSL frontend, and WGSL, SPIR-V, HLSL, MSL, and GLSL backends. By @jamienicol in [#6519](https://github.com/gfx-rs/wgpu/pull/6519). - Add support for GLSL `usampler*` and `isampler*`. By @DavidPeicho in [#6513](https://github.com/gfx-rs/wgpu/pull/6513). @@ -128,7 +129,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148] - Handle query set creation failure as an internal error that loses the `Device`, rather than panicking. By @ErichDonGubler in [#6505](https://github.com/gfx-rs/wgpu/pull/6505). - Ensure that `Features::TIMESTAMP_QUERY` is set when using timestamp writes in render and compute passes. By @ErichDonGubler in [#6497](https://github.com/gfx-rs/wgpu/pull/6497). -- Check for device mismatches when beginning render and compute passes. By @ErichDonGubler in [#6497](https://github.com/gfx-rs/wgpu/pull/6497). +- Check for device mismatches when beginning render and compute passes. By @ErichDonGubler in [#6497](https://github.com/gfx-rs/wgpu/pull/6497). - Lower `QUERY_SET_MAX_QUERIES` (and enforced limits) from 8192 to 4096 to match WebGPU spec. By @ErichDonGubler in [#6525](https://github.com/gfx-rs/wgpu/pull/6525). - Allow non-filterable float on texture bindings never used with samplers when using a derived bind group layout. By @ErichDonGubler in [#6531](https://github.com/gfx-rs/wgpu/pull/6531/). - Replace potentially unsound usage of `PreHashedMap` with `FastHashMap`. By @jamienicol in [#6541](https://github.com/gfx-rs/wgpu/pull/6541). diff --git a/naga/src/front/wgsl/parse/number.rs b/naga/src/front/wgsl/parse/number.rs index ceb2cb336..72795de6b 100644 --- a/naga/src/front/wgsl/parse/number.rs +++ b/naga/src/front/wgsl/parse/number.rs @@ -216,7 +216,12 @@ fn parse(input: &str) -> (Result, &str) { rest_to_str!(bytes), ) } else { - let kind = consume_map!(bytes, [b'i' => IntKind::I32, b'u' => IntKind::U32]); + let kind = consume_map!(bytes, [ + b'i' => IntKind::I32, + b'u' => IntKind::U32, + b'l', b'i' => IntKind::I64, + b'l', b'u' => IntKind::U64, + ]); (parse_hex_int(digits, kind), rest_to_str!(bytes)) } diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index c12983c0c..585174384 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1742,6 +1742,7 @@ impl<'a> ConstantEvaluator<'a> { Expression::Literal(value) => Expression::Literal(match op { UnaryOperator::Negate => match value { Literal::I32(v) => Literal::I32(v.wrapping_neg()), + Literal::I64(v) => Literal::I64(v.wrapping_neg()), Literal::F32(v) => Literal::F32(-v), Literal::AbstractInt(v) => Literal::AbstractInt(v.wrapping_neg()), Literal::AbstractFloat(v) => Literal::AbstractFloat(-v), @@ -1753,7 +1754,9 @@ impl<'a> ConstantEvaluator<'a> { }, UnaryOperator::BitwiseNot => match value { Literal::I32(v) => Literal::I32(!v), + Literal::I64(v) => Literal::I64(!v), Literal::U32(v) => Literal::U32(!v), + Literal::U64(v) => Literal::U64(!v), Literal::AbstractInt(v) => Literal::AbstractInt(!v), _ => return Err(ConstantEvaluatorError::InvalidUnaryOpArg), }, diff --git a/naga/tests/in/int64.wgsl b/naga/tests/in/int64.wgsl index 3da5fefc1..966eb0b55 100644 --- a/naga/tests/in/int64.wgsl +++ b/naga/tests/in/int64.wgsl @@ -41,7 +41,7 @@ var output_arrays: StorageCompatible; fn int64_function(x: i64) -> i64 { var val: i64 = i64(constant_variable); // A number too big for i32 - val += 31li - 1002003004005006li; + val += 31li - 1002003004005006li + -0x7fffffffffffffffli; // Constructing an i64 from an AbstractInt val += val + i64(5); // Constructing a i64 from other types and other types from u64. @@ -89,8 +89,8 @@ fn int64_function(x: i64) -> i64 { fn uint64_function(x: u64) -> u64 { var val: u64 = u64(constant_variable); - // A number too big for u32 - val += 31lu + 1002003004005006lu; + // Numbers too big for u32 (u64::MAX) + val += 31lu + 18446744073709551615lu - 0xfffffffffffffffflu; // Constructing a u64 from an AbstractInt val += val + u64(5); // Constructing a u64 from other types and other types from u64. diff --git a/naga/tests/out/hlsl/int64.hlsl b/naga/tests/out/hlsl/int64.hlsl index ccce279ba..26441380b 100644 --- a/naga/tests/out/hlsl/int64.hlsl +++ b/naga/tests/out/hlsl/int64.hlsl @@ -63,81 +63,81 @@ int64_t int64_function(int64_t x) { int64_t val = 20L; - int64_t _e6 = val; - val = (_e6 + (31L - 1002003004005006L)); int64_t _e8 = val; - int64_t _e11 = val; - val = (_e11 + (_e8 + 5L)); - uint _e15 = input_uniform.val_u32_; - int64_t _e16 = val; - int64_t _e20 = val; - val = (_e20 + int64_t((_e15 + uint(_e16)))); - int _e24 = input_uniform.val_i32_; - int64_t _e25 = val; - int64_t _e29 = val; - val = (_e29 + int64_t((_e24 + int(_e25)))); - float _e33 = input_uniform.val_f32_; - int64_t _e34 = val; - int64_t _e38 = val; - val = (_e38 + int64_t((_e33 + float(_e34)))); - int64_t _e42 = input_uniform.val_i64_; - int64_t _e45 = val; - val = (_e45 + (_e42).xxx.z); - uint64_t _e49 = input_uniform.val_u64_; - int64_t _e51 = val; - val = (_e51 + _e49); - uint64_t2 _e55 = input_uniform.val_u64_2_; - int64_t _e58 = val; - val = (_e58 + _e55.y); - uint64_t3 _e62 = input_uniform.val_u64_3_; - int64_t _e65 = val; - val = (_e65 + _e62.z); - uint64_t4 _e69 = input_uniform.val_u64_4_; - int64_t _e72 = val; - val = (_e72 + _e69.w); - int64_t _e78 = input_uniform.val_i64_; - int64_t _e81 = input_storage.Load(128); - output.Store(128, (_e78 + _e81)); - int64_t2 _e87 = input_uniform.val_i64_2_; - int64_t2 _e90 = input_storage.Load(144); - output.Store(144, (_e87 + _e90)); - int64_t3 _e96 = input_uniform.val_i64_3_; - int64_t3 _e99 = input_storage.Load(160); - output.Store(160, (_e96 + _e99)); - int64_t4 _e105 = input_uniform.val_i64_4_; - int64_t4 _e108 = input_storage.Load(192); - output.Store(192, (_e105 + _e108)); - int64_t _e114[2] = Constructarray2_int64_t_(input_arrays.Load(16+0), input_arrays.Load(16+8)); + val = (_e8 + ((31L - 1002003004005006L) + -9223372036854775807L)); + int64_t _e10 = val; + int64_t _e13 = val; + val = (_e13 + (_e10 + 5L)); + uint _e17 = input_uniform.val_u32_; + int64_t _e18 = val; + int64_t _e22 = val; + val = (_e22 + int64_t((_e17 + uint(_e18)))); + int _e26 = input_uniform.val_i32_; + int64_t _e27 = val; + int64_t _e31 = val; + val = (_e31 + int64_t((_e26 + int(_e27)))); + float _e35 = input_uniform.val_f32_; + int64_t _e36 = val; + int64_t _e40 = val; + val = (_e40 + int64_t((_e35 + float(_e36)))); + int64_t _e44 = input_uniform.val_i64_; + int64_t _e47 = val; + val = (_e47 + (_e44).xxx.z); + uint64_t _e51 = input_uniform.val_u64_; + int64_t _e53 = val; + val = (_e53 + _e51); + uint64_t2 _e57 = input_uniform.val_u64_2_; + int64_t _e60 = val; + val = (_e60 + _e57.y); + uint64_t3 _e64 = input_uniform.val_u64_3_; + int64_t _e67 = val; + val = (_e67 + _e64.z); + uint64_t4 _e71 = input_uniform.val_u64_4_; + int64_t _e74 = val; + val = (_e74 + _e71.w); + int64_t _e80 = input_uniform.val_i64_; + int64_t _e83 = input_storage.Load(128); + output.Store(128, (_e80 + _e83)); + int64_t2 _e89 = input_uniform.val_i64_2_; + int64_t2 _e92 = input_storage.Load(144); + output.Store(144, (_e89 + _e92)); + int64_t3 _e98 = input_uniform.val_i64_3_; + int64_t3 _e101 = input_storage.Load(160); + output.Store(160, (_e98 + _e101)); + int64_t4 _e107 = input_uniform.val_i64_4_; + int64_t4 _e110 = input_storage.Load(192); + output.Store(192, (_e107 + _e110)); + int64_t _e116[2] = Constructarray2_int64_t_(input_arrays.Load(16+0), input_arrays.Load(16+8)); { - int64_t _value2[2] = _e114; + int64_t _value2[2] = _e116; output_arrays.Store(16+0, _value2[0]); output_arrays.Store(16+8, _value2[1]); } - int64_t _e115 = val; int64_t _e117 = val; - val = (_e117 + abs(_e115)); int64_t _e119 = val; - int64_t _e120 = val; + val = (_e119 + abs(_e117)); int64_t _e121 = val; + int64_t _e122 = val; int64_t _e123 = val; - val = (_e123 + clamp(_e119, _e120, _e121)); int64_t _e125 = val; + val = (_e125 + clamp(_e121, _e122, _e123)); int64_t _e127 = val; - int64_t _e130 = val; - val = (_e130 + dot((_e125).xx, (_e127).xx)); + int64_t _e129 = val; int64_t _e132 = val; - int64_t _e133 = val; + val = (_e132 + dot((_e127).xx, (_e129).xx)); + int64_t _e134 = val; int64_t _e135 = val; - val = (_e135 + max(_e132, _e133)); int64_t _e137 = val; - int64_t _e138 = val; + val = (_e137 + max(_e134, _e135)); + int64_t _e139 = val; int64_t _e140 = val; - val = (_e140 + min(_e137, _e138)); int64_t _e142 = val; + val = (_e142 + min(_e139, _e140)); int64_t _e144 = val; - val = (_e144 + sign(_e142)); int64_t _e146 = val; - return _e146; + val = (_e146 + sign(_e144)); + int64_t _e148 = val; + return _e148; } typedef uint64_t ret_Constructarray2_uint64_t_[2]; @@ -150,78 +150,78 @@ uint64_t uint64_function(uint64_t x_1) { uint64_t val_1 = 20uL; - uint64_t _e6 = val_1; - val_1 = (_e6 + (31uL + 1002003004005006uL)); uint64_t _e8 = val_1; - uint64_t _e11 = val_1; - val_1 = (_e11 + (_e8 + 5uL)); - uint _e15 = input_uniform.val_u32_; - uint64_t _e16 = val_1; - uint64_t _e20 = val_1; - val_1 = (_e20 + uint64_t((_e15 + uint(_e16)))); - int _e24 = input_uniform.val_i32_; - uint64_t _e25 = val_1; - uint64_t _e29 = val_1; - val_1 = (_e29 + uint64_t((_e24 + int(_e25)))); - float _e33 = input_uniform.val_f32_; - uint64_t _e34 = val_1; - uint64_t _e38 = val_1; - val_1 = (_e38 + uint64_t((_e33 + float(_e34)))); - uint64_t _e42 = input_uniform.val_u64_; - uint64_t _e45 = val_1; - val_1 = (_e45 + (_e42).xxx.z); - int64_t _e49 = input_uniform.val_i64_; - uint64_t _e51 = val_1; - val_1 = (_e51 + _e49); - int64_t2 _e55 = input_uniform.val_i64_2_; - uint64_t _e58 = val_1; - val_1 = (_e58 + _e55.y); - int64_t3 _e62 = input_uniform.val_i64_3_; - uint64_t _e65 = val_1; - val_1 = (_e65 + _e62.z); - int64_t4 _e69 = input_uniform.val_i64_4_; - uint64_t _e72 = val_1; - val_1 = (_e72 + _e69.w); - uint64_t _e78 = input_uniform.val_u64_; - uint64_t _e81 = input_storage.Load(16); - output.Store(16, (_e78 + _e81)); - uint64_t2 _e87 = input_uniform.val_u64_2_; - uint64_t2 _e90 = input_storage.Load(32); - output.Store(32, (_e87 + _e90)); - uint64_t3 _e96 = input_uniform.val_u64_3_; - uint64_t3 _e99 = input_storage.Load(64); - output.Store(64, (_e96 + _e99)); - uint64_t4 _e105 = input_uniform.val_u64_4_; - uint64_t4 _e108 = input_storage.Load(96); - output.Store(96, (_e105 + _e108)); - uint64_t _e114[2] = Constructarray2_uint64_t_(input_arrays.Load(0+0), input_arrays.Load(0+8)); + val_1 = (_e8 + ((31uL + 18446744073709551615uL) - 18446744073709551615uL)); + uint64_t _e10 = val_1; + uint64_t _e13 = val_1; + val_1 = (_e13 + (_e10 + 5uL)); + uint _e17 = input_uniform.val_u32_; + uint64_t _e18 = val_1; + uint64_t _e22 = val_1; + val_1 = (_e22 + uint64_t((_e17 + uint(_e18)))); + int _e26 = input_uniform.val_i32_; + uint64_t _e27 = val_1; + uint64_t _e31 = val_1; + val_1 = (_e31 + uint64_t((_e26 + int(_e27)))); + float _e35 = input_uniform.val_f32_; + uint64_t _e36 = val_1; + uint64_t _e40 = val_1; + val_1 = (_e40 + uint64_t((_e35 + float(_e36)))); + uint64_t _e44 = input_uniform.val_u64_; + uint64_t _e47 = val_1; + val_1 = (_e47 + (_e44).xxx.z); + int64_t _e51 = input_uniform.val_i64_; + uint64_t _e53 = val_1; + val_1 = (_e53 + _e51); + int64_t2 _e57 = input_uniform.val_i64_2_; + uint64_t _e60 = val_1; + val_1 = (_e60 + _e57.y); + int64_t3 _e64 = input_uniform.val_i64_3_; + uint64_t _e67 = val_1; + val_1 = (_e67 + _e64.z); + int64_t4 _e71 = input_uniform.val_i64_4_; + uint64_t _e74 = val_1; + val_1 = (_e74 + _e71.w); + uint64_t _e80 = input_uniform.val_u64_; + uint64_t _e83 = input_storage.Load(16); + output.Store(16, (_e80 + _e83)); + uint64_t2 _e89 = input_uniform.val_u64_2_; + uint64_t2 _e92 = input_storage.Load(32); + output.Store(32, (_e89 + _e92)); + uint64_t3 _e98 = input_uniform.val_u64_3_; + uint64_t3 _e101 = input_storage.Load(64); + output.Store(64, (_e98 + _e101)); + uint64_t4 _e107 = input_uniform.val_u64_4_; + uint64_t4 _e110 = input_storage.Load(96); + output.Store(96, (_e107 + _e110)); + uint64_t _e116[2] = Constructarray2_uint64_t_(input_arrays.Load(0+0), input_arrays.Load(0+8)); { - uint64_t _value2[2] = _e114; + uint64_t _value2[2] = _e116; output_arrays.Store(0+0, _value2[0]); output_arrays.Store(0+8, _value2[1]); } - uint64_t _e115 = val_1; uint64_t _e117 = val_1; - val_1 = (_e117 + abs(_e115)); uint64_t _e119 = val_1; - uint64_t _e120 = val_1; + val_1 = (_e119 + abs(_e117)); uint64_t _e121 = val_1; + uint64_t _e122 = val_1; uint64_t _e123 = val_1; - val_1 = (_e123 + clamp(_e119, _e120, _e121)); uint64_t _e125 = val_1; + val_1 = (_e125 + clamp(_e121, _e122, _e123)); uint64_t _e127 = val_1; - uint64_t _e130 = val_1; - val_1 = (_e130 + dot((_e125).xx, (_e127).xx)); + uint64_t _e129 = val_1; uint64_t _e132 = val_1; - uint64_t _e133 = val_1; + val_1 = (_e132 + dot((_e127).xx, (_e129).xx)); + uint64_t _e134 = val_1; uint64_t _e135 = val_1; - val_1 = (_e135 + max(_e132, _e133)); uint64_t _e137 = val_1; - uint64_t _e138 = val_1; + val_1 = (_e137 + max(_e134, _e135)); + uint64_t _e139 = val_1; uint64_t _e140 = val_1; - val_1 = (_e140 + min(_e137, _e138)); uint64_t _e142 = val_1; - return _e142; + val_1 = (_e142 + min(_e139, _e140)); + uint64_t _e144 = val_1; + return _e144; } [numthreads(1, 1, 1)] diff --git a/naga/tests/out/msl/int64.msl b/naga/tests/out/msl/int64.msl index 2ef03d9ae..79a3c0688 100644 --- a/naga/tests/out/msl/int64.msl +++ b/naga/tests/out/msl/int64.msl @@ -43,79 +43,79 @@ long int64_function( device StorageCompatible& output_arrays ) { long val = 20L; - long _e6 = val; - val = _e6 + (31L - 1002003004005006L); long _e8 = val; - long _e11 = val; - val = _e11 + (_e8 + 5L); - uint _e15 = input_uniform.val_u32_; - long _e16 = val; - long _e20 = val; - val = _e20 + static_cast(_e15 + static_cast(_e16)); - int _e24 = input_uniform.val_i32_; - long _e25 = val; - long _e29 = val; - val = _e29 + static_cast(_e24 + static_cast(_e25)); - float _e33 = input_uniform.val_f32_; - long _e34 = val; - long _e38 = val; - val = _e38 + static_cast(_e33 + static_cast(_e34)); - long _e42 = input_uniform.val_i64_; - long _e45 = val; - val = _e45 + metal::long3(_e42).z; - ulong _e49 = input_uniform.val_u64_; - long _e51 = val; - val = _e51 + as_type(_e49); - metal::ulong2 _e55 = input_uniform.val_u64_2_; - long _e58 = val; - val = _e58 + as_type(_e55).y; - metal::ulong3 _e62 = input_uniform.val_u64_3_; - long _e65 = val; - val = _e65 + as_type(_e62).z; - metal::ulong4 _e69 = input_uniform.val_u64_4_; - long _e72 = val; - val = _e72 + as_type(_e69).w; - long _e78 = input_uniform.val_i64_; - long _e81 = input_storage.val_i64_; - output.val_i64_ = _e78 + _e81; - metal::long2 _e87 = input_uniform.val_i64_2_; - metal::long2 _e90 = input_storage.val_i64_2_; - output.val_i64_2_ = _e87 + _e90; - metal::long3 _e96 = input_uniform.val_i64_3_; - metal::long3 _e99 = input_storage.val_i64_3_; - output.val_i64_3_ = _e96 + _e99; - metal::long4 _e105 = input_uniform.val_i64_4_; - metal::long4 _e108 = input_storage.val_i64_4_; - output.val_i64_4_ = _e105 + _e108; - type_12 _e114 = input_arrays.val_i64_array_2_; - output_arrays.val_i64_array_2_ = _e114; - long _e115 = val; + val = _e8 + ((31L - 1002003004005006L) + -9223372036854775807L); + long _e10 = val; + long _e13 = val; + val = _e13 + (_e10 + 5L); + uint _e17 = input_uniform.val_u32_; + long _e18 = val; + long _e22 = val; + val = _e22 + static_cast(_e17 + static_cast(_e18)); + int _e26 = input_uniform.val_i32_; + long _e27 = val; + long _e31 = val; + val = _e31 + static_cast(_e26 + static_cast(_e27)); + float _e35 = input_uniform.val_f32_; + long _e36 = val; + long _e40 = val; + val = _e40 + static_cast(_e35 + static_cast(_e36)); + long _e44 = input_uniform.val_i64_; + long _e47 = val; + val = _e47 + metal::long3(_e44).z; + ulong _e51 = input_uniform.val_u64_; + long _e53 = val; + val = _e53 + as_type(_e51); + metal::ulong2 _e57 = input_uniform.val_u64_2_; + long _e60 = val; + val = _e60 + as_type(_e57).y; + metal::ulong3 _e64 = input_uniform.val_u64_3_; + long _e67 = val; + val = _e67 + as_type(_e64).z; + metal::ulong4 _e71 = input_uniform.val_u64_4_; + long _e74 = val; + val = _e74 + as_type(_e71).w; + long _e80 = input_uniform.val_i64_; + long _e83 = input_storage.val_i64_; + output.val_i64_ = _e80 + _e83; + metal::long2 _e89 = input_uniform.val_i64_2_; + metal::long2 _e92 = input_storage.val_i64_2_; + output.val_i64_2_ = _e89 + _e92; + metal::long3 _e98 = input_uniform.val_i64_3_; + metal::long3 _e101 = input_storage.val_i64_3_; + output.val_i64_3_ = _e98 + _e101; + metal::long4 _e107 = input_uniform.val_i64_4_; + metal::long4 _e110 = input_storage.val_i64_4_; + output.val_i64_4_ = _e107 + _e110; + type_12 _e116 = input_arrays.val_i64_array_2_; + output_arrays.val_i64_array_2_ = _e116; long _e117 = val; - val = _e117 + metal::abs(_e115); long _e119 = val; - long _e120 = val; + val = _e119 + metal::abs(_e117); long _e121 = val; + long _e122 = val; long _e123 = val; - val = _e123 + metal::clamp(_e119, _e120, _e121); long _e125 = val; - metal::long2 _e126 = metal::long2(_e125); + val = _e125 + metal::clamp(_e121, _e122, _e123); long _e127 = val; metal::long2 _e128 = metal::long2(_e127); - long _e130 = val; - val = _e130 + ( + _e126.x * _e128.x + _e126.y * _e128.y); + long _e129 = val; + metal::long2 _e130 = metal::long2(_e129); long _e132 = val; - long _e133 = val; + val = _e132 + ( + _e128.x * _e130.x + _e128.y * _e130.y); + long _e134 = val; long _e135 = val; - val = _e135 + metal::max(_e132, _e133); long _e137 = val; - long _e138 = val; + val = _e137 + metal::max(_e134, _e135); + long _e139 = val; long _e140 = val; - val = _e140 + metal::min(_e137, _e138); long _e142 = val; + val = _e142 + metal::min(_e139, _e140); long _e144 = val; - val = _e144 + metal::select(metal::select(-1, 1, (_e142 > 0)), 0, (_e142 == 0)); long _e146 = val; - return _e146; + val = _e146 + metal::select(metal::select(-1, 1, (_e144 > 0)), 0, (_e144 == 0)); + long _e148 = val; + return _e148; } ulong uint64_function( @@ -127,76 +127,76 @@ ulong uint64_function( device StorageCompatible& output_arrays ) { ulong val_1 = 20uL; - ulong _e6 = val_1; - val_1 = _e6 + (31uL + 1002003004005006uL); ulong _e8 = val_1; - ulong _e11 = val_1; - val_1 = _e11 + (_e8 + 5uL); - uint _e15 = input_uniform.val_u32_; - ulong _e16 = val_1; - ulong _e20 = val_1; - val_1 = _e20 + static_cast(_e15 + static_cast(_e16)); - int _e24 = input_uniform.val_i32_; - ulong _e25 = val_1; - ulong _e29 = val_1; - val_1 = _e29 + static_cast(_e24 + static_cast(_e25)); - float _e33 = input_uniform.val_f32_; - ulong _e34 = val_1; - ulong _e38 = val_1; - val_1 = _e38 + static_cast(_e33 + static_cast(_e34)); - ulong _e42 = input_uniform.val_u64_; - ulong _e45 = val_1; - val_1 = _e45 + metal::ulong3(_e42).z; - long _e49 = input_uniform.val_i64_; - ulong _e51 = val_1; - val_1 = _e51 + as_type(_e49); - metal::long2 _e55 = input_uniform.val_i64_2_; - ulong _e58 = val_1; - val_1 = _e58 + as_type(_e55).y; - metal::long3 _e62 = input_uniform.val_i64_3_; - ulong _e65 = val_1; - val_1 = _e65 + as_type(_e62).z; - metal::long4 _e69 = input_uniform.val_i64_4_; - ulong _e72 = val_1; - val_1 = _e72 + as_type(_e69).w; - ulong _e78 = input_uniform.val_u64_; - ulong _e81 = input_storage.val_u64_; - output.val_u64_ = _e78 + _e81; - metal::ulong2 _e87 = input_uniform.val_u64_2_; - metal::ulong2 _e90 = input_storage.val_u64_2_; - output.val_u64_2_ = _e87 + _e90; - metal::ulong3 _e96 = input_uniform.val_u64_3_; - metal::ulong3 _e99 = input_storage.val_u64_3_; - output.val_u64_3_ = _e96 + _e99; - metal::ulong4 _e105 = input_uniform.val_u64_4_; - metal::ulong4 _e108 = input_storage.val_u64_4_; - output.val_u64_4_ = _e105 + _e108; - type_11 _e114 = input_arrays.val_u64_array_2_; - output_arrays.val_u64_array_2_ = _e114; - ulong _e115 = val_1; + val_1 = _e8 + ((31uL + 18446744073709551615uL) - 18446744073709551615uL); + ulong _e10 = val_1; + ulong _e13 = val_1; + val_1 = _e13 + (_e10 + 5uL); + uint _e17 = input_uniform.val_u32_; + ulong _e18 = val_1; + ulong _e22 = val_1; + val_1 = _e22 + static_cast(_e17 + static_cast(_e18)); + int _e26 = input_uniform.val_i32_; + ulong _e27 = val_1; + ulong _e31 = val_1; + val_1 = _e31 + static_cast(_e26 + static_cast(_e27)); + float _e35 = input_uniform.val_f32_; + ulong _e36 = val_1; + ulong _e40 = val_1; + val_1 = _e40 + static_cast(_e35 + static_cast(_e36)); + ulong _e44 = input_uniform.val_u64_; + ulong _e47 = val_1; + val_1 = _e47 + metal::ulong3(_e44).z; + long _e51 = input_uniform.val_i64_; + ulong _e53 = val_1; + val_1 = _e53 + as_type(_e51); + metal::long2 _e57 = input_uniform.val_i64_2_; + ulong _e60 = val_1; + val_1 = _e60 + as_type(_e57).y; + metal::long3 _e64 = input_uniform.val_i64_3_; + ulong _e67 = val_1; + val_1 = _e67 + as_type(_e64).z; + metal::long4 _e71 = input_uniform.val_i64_4_; + ulong _e74 = val_1; + val_1 = _e74 + as_type(_e71).w; + ulong _e80 = input_uniform.val_u64_; + ulong _e83 = input_storage.val_u64_; + output.val_u64_ = _e80 + _e83; + metal::ulong2 _e89 = input_uniform.val_u64_2_; + metal::ulong2 _e92 = input_storage.val_u64_2_; + output.val_u64_2_ = _e89 + _e92; + metal::ulong3 _e98 = input_uniform.val_u64_3_; + metal::ulong3 _e101 = input_storage.val_u64_3_; + output.val_u64_3_ = _e98 + _e101; + metal::ulong4 _e107 = input_uniform.val_u64_4_; + metal::ulong4 _e110 = input_storage.val_u64_4_; + output.val_u64_4_ = _e107 + _e110; + type_11 _e116 = input_arrays.val_u64_array_2_; + output_arrays.val_u64_array_2_ = _e116; ulong _e117 = val_1; - val_1 = _e117 + metal::abs(_e115); ulong _e119 = val_1; - ulong _e120 = val_1; + val_1 = _e119 + metal::abs(_e117); ulong _e121 = val_1; + ulong _e122 = val_1; ulong _e123 = val_1; - val_1 = _e123 + metal::clamp(_e119, _e120, _e121); ulong _e125 = val_1; - metal::ulong2 _e126 = metal::ulong2(_e125); + val_1 = _e125 + metal::clamp(_e121, _e122, _e123); ulong _e127 = val_1; metal::ulong2 _e128 = metal::ulong2(_e127); - ulong _e130 = val_1; - val_1 = _e130 + ( + _e126.x * _e128.x + _e126.y * _e128.y); + ulong _e129 = val_1; + metal::ulong2 _e130 = metal::ulong2(_e129); ulong _e132 = val_1; - ulong _e133 = val_1; + val_1 = _e132 + ( + _e128.x * _e130.x + _e128.y * _e130.y); + ulong _e134 = val_1; ulong _e135 = val_1; - val_1 = _e135 + metal::max(_e132, _e133); ulong _e137 = val_1; - ulong _e138 = val_1; + val_1 = _e137 + metal::max(_e134, _e135); + ulong _e139 = val_1; ulong _e140 = val_1; - val_1 = _e140 + metal::min(_e137, _e138); ulong _e142 = val_1; - return _e142; + val_1 = _e142 + metal::min(_e139, _e140); + ulong _e144 = val_1; + return _e144; } kernel void main_( diff --git a/naga/tests/out/spv/int64.spvasm b/naga/tests/out/spv/int64.spvasm index a60a14d75..fdd2bcab1 100644 --- a/naga/tests/out/spv/int64.spvasm +++ b/naga/tests/out/spv/int64.spvasm @@ -1,14 +1,14 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 372 +; Bound: 375 OpCapability Shader OpCapability Int64 OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %356 "main" -OpExecutionMode %356 LocalSize 1 1 1 +OpEntryPoint GLCompute %359 "main" +OpExecutionMode %359 LocalSize 1 1 1 OpMemberDecorate %14 0 Offset 0 OpMemberDecorate %14 1 Offset 4 OpMemberDecorate %14 2 Offset 8 @@ -91,380 +91,383 @@ OpMemberDecorate %36 0 Offset 0 %51 = OpConstant %3 20 %52 = OpConstant %3 31 %53 = OpConstant %3 1002003004005006 -%54 = OpConstant %3 5 -%56 = OpTypePointer Function %3 -%65 = OpTypePointer Uniform %5 -%74 = OpTypePointer Uniform %6 -%75 = OpConstant %5 1 -%84 = OpTypePointer Uniform %7 -%93 = OpTypePointer Uniform %3 -%94 = OpConstant %5 7 -%101 = OpTypePointer Uniform %4 -%102 = OpConstant %5 3 -%108 = OpTypePointer Uniform %8 -%109 = OpConstant %5 4 -%116 = OpTypePointer Uniform %9 -%117 = OpConstant %5 5 -%124 = OpTypePointer Uniform %10 -%125 = OpConstant %5 6 -%132 = OpTypePointer StorageBuffer %3 -%139 = OpTypePointer StorageBuffer %11 -%140 = OpTypePointer Uniform %11 -%141 = OpConstant %5 8 -%148 = OpTypePointer StorageBuffer %12 -%149 = OpTypePointer Uniform %12 -%150 = OpConstant %5 9 -%157 = OpTypePointer StorageBuffer %13 -%158 = OpTypePointer Uniform %13 -%159 = OpConstant %5 10 -%166 = OpTypePointer StorageBuffer %17 -%186 = OpConstantNull %3 -%214 = OpTypeFunction %4 %4 -%220 = OpConstant %4 31 -%221 = OpConstant %4 1002003004005006 -%222 = OpConstant %4 5 -%224 = OpTypePointer Function %4 -%286 = OpTypePointer StorageBuffer %4 -%293 = OpTypePointer StorageBuffer %8 -%300 = OpTypePointer StorageBuffer %9 -%307 = OpTypePointer StorageBuffer %10 -%314 = OpTypePointer StorageBuffer %15 -%334 = OpConstantNull %4 -%357 = OpTypeFunction %2 -%363 = OpConstant %4 67 -%364 = OpConstant %3 60 -%370 = OpConstant %5 11 +%54 = OpConstant %3 -9223372036854775807 +%55 = OpConstant %3 5 +%57 = OpTypePointer Function %3 +%67 = OpTypePointer Uniform %5 +%76 = OpTypePointer Uniform %6 +%77 = OpConstant %5 1 +%86 = OpTypePointer Uniform %7 +%95 = OpTypePointer Uniform %3 +%96 = OpConstant %5 7 +%103 = OpTypePointer Uniform %4 +%104 = OpConstant %5 3 +%110 = OpTypePointer Uniform %8 +%111 = OpConstant %5 4 +%118 = OpTypePointer Uniform %9 +%119 = OpConstant %5 5 +%126 = OpTypePointer Uniform %10 +%127 = OpConstant %5 6 +%134 = OpTypePointer StorageBuffer %3 +%141 = OpTypePointer StorageBuffer %11 +%142 = OpTypePointer Uniform %11 +%143 = OpConstant %5 8 +%150 = OpTypePointer StorageBuffer %12 +%151 = OpTypePointer Uniform %12 +%152 = OpConstant %5 9 +%159 = OpTypePointer StorageBuffer %13 +%160 = OpTypePointer Uniform %13 +%161 = OpConstant %5 10 +%168 = OpTypePointer StorageBuffer %17 +%188 = OpConstantNull %3 +%216 = OpTypeFunction %4 %4 +%222 = OpConstant %4 31 +%223 = OpConstant %4 18446744073709551615 +%224 = OpConstant %4 5 +%226 = OpTypePointer Function %4 +%289 = OpTypePointer StorageBuffer %4 +%296 = OpTypePointer StorageBuffer %8 +%303 = OpTypePointer StorageBuffer %9 +%310 = OpTypePointer StorageBuffer %10 +%317 = OpTypePointer StorageBuffer %15 +%337 = OpConstantNull %4 +%360 = OpTypeFunction %2 +%366 = OpConstant %4 67 +%367 = OpConstant %3 60 +%373 = OpConstant %5 11 %40 = OpFunction %3 None %41 %39 = OpFunctionParameter %3 %38 = OpLabel -%55 = OpVariable %56 Function %51 +%56 = OpVariable %57 Function %51 %44 = OpAccessChain %42 %23 %43 %46 = OpAccessChain %45 %26 %43 %48 = OpAccessChain %47 %29 %43 %49 = OpAccessChain %45 %32 %43 %50 = OpAccessChain %47 %35 %43 -OpBranch %57 -%57 = OpLabel -%58 = OpISub %3 %52 %53 -%59 = OpLoad %3 %55 -%60 = OpIAdd %3 %59 %58 -OpStore %55 %60 -%61 = OpLoad %3 %55 -%62 = OpIAdd %3 %61 %54 -%63 = OpLoad %3 %55 -%64 = OpIAdd %3 %63 %62 -OpStore %55 %64 -%66 = OpAccessChain %65 %44 %43 -%67 = OpLoad %5 %66 -%68 = OpLoad %3 %55 -%69 = OpUConvert %5 %68 -%70 = OpIAdd %5 %67 %69 -%71 = OpSConvert %3 %70 -%72 = OpLoad %3 %55 -%73 = OpIAdd %3 %72 %71 -OpStore %55 %73 -%76 = OpAccessChain %74 %44 %75 -%77 = OpLoad %6 %76 -%78 = OpLoad %3 %55 -%79 = OpSConvert %6 %78 -%80 = OpIAdd %6 %77 %79 -%81 = OpSConvert %3 %80 -%82 = OpLoad %3 %55 -%83 = OpIAdd %3 %82 %81 -OpStore %55 %83 -%85 = OpAccessChain %84 %44 %16 -%86 = OpLoad %7 %85 -%87 = OpLoad %3 %55 -%88 = OpConvertSToF %7 %87 -%89 = OpFAdd %7 %86 %88 -%90 = OpConvertFToS %3 %89 -%91 = OpLoad %3 %55 -%92 = OpIAdd %3 %91 %90 -OpStore %55 %92 -%95 = OpAccessChain %93 %44 %94 -%96 = OpLoad %3 %95 -%97 = OpCompositeConstruct %12 %96 %96 %96 -%98 = OpCompositeExtract %3 %97 2 -%99 = OpLoad %3 %55 -%100 = OpIAdd %3 %99 %98 -OpStore %55 %100 -%103 = OpAccessChain %101 %44 %102 -%104 = OpLoad %4 %103 -%105 = OpBitcast %3 %104 -%106 = OpLoad %3 %55 -%107 = OpIAdd %3 %106 %105 -OpStore %55 %107 -%110 = OpAccessChain %108 %44 %109 -%111 = OpLoad %8 %110 -%112 = OpBitcast %11 %111 -%113 = OpCompositeExtract %3 %112 1 -%114 = OpLoad %3 %55 -%115 = OpIAdd %3 %114 %113 -OpStore %55 %115 -%118 = OpAccessChain %116 %44 %117 -%119 = OpLoad %9 %118 -%120 = OpBitcast %12 %119 -%121 = OpCompositeExtract %3 %120 2 -%122 = OpLoad %3 %55 -%123 = OpIAdd %3 %122 %121 -OpStore %55 %123 -%126 = OpAccessChain %124 %44 %125 -%127 = OpLoad %10 %126 -%128 = OpBitcast %13 %127 -%129 = OpCompositeExtract %3 %128 3 -%130 = OpLoad %3 %55 -%131 = OpIAdd %3 %130 %129 -OpStore %55 %131 -%133 = OpAccessChain %93 %44 %94 -%134 = OpLoad %3 %133 -%135 = OpAccessChain %132 %46 %94 +OpBranch %58 +%58 = OpLabel +%59 = OpISub %3 %52 %53 +%60 = OpIAdd %3 %59 %54 +%61 = OpLoad %3 %56 +%62 = OpIAdd %3 %61 %60 +OpStore %56 %62 +%63 = OpLoad %3 %56 +%64 = OpIAdd %3 %63 %55 +%65 = OpLoad %3 %56 +%66 = OpIAdd %3 %65 %64 +OpStore %56 %66 +%68 = OpAccessChain %67 %44 %43 +%69 = OpLoad %5 %68 +%70 = OpLoad %3 %56 +%71 = OpUConvert %5 %70 +%72 = OpIAdd %5 %69 %71 +%73 = OpSConvert %3 %72 +%74 = OpLoad %3 %56 +%75 = OpIAdd %3 %74 %73 +OpStore %56 %75 +%78 = OpAccessChain %76 %44 %77 +%79 = OpLoad %6 %78 +%80 = OpLoad %3 %56 +%81 = OpSConvert %6 %80 +%82 = OpIAdd %6 %79 %81 +%83 = OpSConvert %3 %82 +%84 = OpLoad %3 %56 +%85 = OpIAdd %3 %84 %83 +OpStore %56 %85 +%87 = OpAccessChain %86 %44 %16 +%88 = OpLoad %7 %87 +%89 = OpLoad %3 %56 +%90 = OpConvertSToF %7 %89 +%91 = OpFAdd %7 %88 %90 +%92 = OpConvertFToS %3 %91 +%93 = OpLoad %3 %56 +%94 = OpIAdd %3 %93 %92 +OpStore %56 %94 +%97 = OpAccessChain %95 %44 %96 +%98 = OpLoad %3 %97 +%99 = OpCompositeConstruct %12 %98 %98 %98 +%100 = OpCompositeExtract %3 %99 2 +%101 = OpLoad %3 %56 +%102 = OpIAdd %3 %101 %100 +OpStore %56 %102 +%105 = OpAccessChain %103 %44 %104 +%106 = OpLoad %4 %105 +%107 = OpBitcast %3 %106 +%108 = OpLoad %3 %56 +%109 = OpIAdd %3 %108 %107 +OpStore %56 %109 +%112 = OpAccessChain %110 %44 %111 +%113 = OpLoad %8 %112 +%114 = OpBitcast %11 %113 +%115 = OpCompositeExtract %3 %114 1 +%116 = OpLoad %3 %56 +%117 = OpIAdd %3 %116 %115 +OpStore %56 %117 +%120 = OpAccessChain %118 %44 %119 +%121 = OpLoad %9 %120 +%122 = OpBitcast %12 %121 +%123 = OpCompositeExtract %3 %122 2 +%124 = OpLoad %3 %56 +%125 = OpIAdd %3 %124 %123 +OpStore %56 %125 +%128 = OpAccessChain %126 %44 %127 +%129 = OpLoad %10 %128 +%130 = OpBitcast %13 %129 +%131 = OpCompositeExtract %3 %130 3 +%132 = OpLoad %3 %56 +%133 = OpIAdd %3 %132 %131 +OpStore %56 %133 +%135 = OpAccessChain %95 %44 %96 %136 = OpLoad %3 %135 -%137 = OpIAdd %3 %134 %136 -%138 = OpAccessChain %132 %49 %94 -OpStore %138 %137 -%142 = OpAccessChain %140 %44 %141 -%143 = OpLoad %11 %142 -%144 = OpAccessChain %139 %46 %141 +%137 = OpAccessChain %134 %46 %96 +%138 = OpLoad %3 %137 +%139 = OpIAdd %3 %136 %138 +%140 = OpAccessChain %134 %49 %96 +OpStore %140 %139 +%144 = OpAccessChain %142 %44 %143 %145 = OpLoad %11 %144 -%146 = OpIAdd %11 %143 %145 -%147 = OpAccessChain %139 %49 %141 -OpStore %147 %146 -%151 = OpAccessChain %149 %44 %150 -%152 = OpLoad %12 %151 -%153 = OpAccessChain %148 %46 %150 +%146 = OpAccessChain %141 %46 %143 +%147 = OpLoad %11 %146 +%148 = OpIAdd %11 %145 %147 +%149 = OpAccessChain %141 %49 %143 +OpStore %149 %148 +%153 = OpAccessChain %151 %44 %152 %154 = OpLoad %12 %153 -%155 = OpIAdd %12 %152 %154 -%156 = OpAccessChain %148 %49 %150 -OpStore %156 %155 -%160 = OpAccessChain %158 %44 %159 -%161 = OpLoad %13 %160 -%162 = OpAccessChain %157 %46 %159 +%155 = OpAccessChain %150 %46 %152 +%156 = OpLoad %12 %155 +%157 = OpIAdd %12 %154 %156 +%158 = OpAccessChain %150 %49 %152 +OpStore %158 %157 +%162 = OpAccessChain %160 %44 %161 %163 = OpLoad %13 %162 -%164 = OpIAdd %13 %161 %163 -%165 = OpAccessChain %157 %49 %159 -OpStore %165 %164 -%167 = OpAccessChain %166 %48 %75 -%168 = OpLoad %17 %167 -%169 = OpAccessChain %166 %50 %75 -OpStore %169 %168 -%170 = OpLoad %3 %55 -%171 = OpExtInst %3 %1 SAbs %170 -%172 = OpLoad %3 %55 -%173 = OpIAdd %3 %172 %171 -OpStore %55 %173 -%174 = OpLoad %3 %55 -%175 = OpLoad %3 %55 -%176 = OpLoad %3 %55 -%178 = OpExtInst %3 %1 SMax %174 %175 -%177 = OpExtInst %3 %1 SMin %178 %176 -%179 = OpLoad %3 %55 -%180 = OpIAdd %3 %179 %177 -OpStore %55 %180 -%181 = OpLoad %3 %55 -%182 = OpCompositeConstruct %11 %181 %181 -%183 = OpLoad %3 %55 +%164 = OpAccessChain %159 %46 %161 +%165 = OpLoad %13 %164 +%166 = OpIAdd %13 %163 %165 +%167 = OpAccessChain %159 %49 %161 +OpStore %167 %166 +%169 = OpAccessChain %168 %48 %77 +%170 = OpLoad %17 %169 +%171 = OpAccessChain %168 %50 %77 +OpStore %171 %170 +%172 = OpLoad %3 %56 +%173 = OpExtInst %3 %1 SAbs %172 +%174 = OpLoad %3 %56 +%175 = OpIAdd %3 %174 %173 +OpStore %56 %175 +%176 = OpLoad %3 %56 +%177 = OpLoad %3 %56 +%178 = OpLoad %3 %56 +%180 = OpExtInst %3 %1 SMax %176 %177 +%179 = OpExtInst %3 %1 SMin %180 %178 +%181 = OpLoad %3 %56 +%182 = OpIAdd %3 %181 %179 +OpStore %56 %182 +%183 = OpLoad %3 %56 %184 = OpCompositeConstruct %11 %183 %183 -%187 = OpCompositeExtract %3 %182 0 -%188 = OpCompositeExtract %3 %184 0 -%189 = OpIMul %3 %187 %188 -%190 = OpIAdd %3 %186 %189 -%191 = OpCompositeExtract %3 %182 1 -%192 = OpCompositeExtract %3 %184 1 -%193 = OpIMul %3 %191 %192 -%185 = OpIAdd %3 %190 %193 -%194 = OpLoad %3 %55 -%195 = OpIAdd %3 %194 %185 -OpStore %55 %195 -%196 = OpLoad %3 %55 -%197 = OpLoad %3 %55 -%198 = OpExtInst %3 %1 SMax %196 %197 -%199 = OpLoad %3 %55 -%200 = OpIAdd %3 %199 %198 -OpStore %55 %200 -%201 = OpLoad %3 %55 -%202 = OpLoad %3 %55 -%203 = OpExtInst %3 %1 SMin %201 %202 -%204 = OpLoad %3 %55 -%205 = OpIAdd %3 %204 %203 -OpStore %55 %205 -%206 = OpLoad %3 %55 -%207 = OpExtInst %3 %1 SSign %206 -%208 = OpLoad %3 %55 -%209 = OpIAdd %3 %208 %207 -OpStore %55 %209 -%210 = OpLoad %3 %55 -OpReturnValue %210 +%185 = OpLoad %3 %56 +%186 = OpCompositeConstruct %11 %185 %185 +%189 = OpCompositeExtract %3 %184 0 +%190 = OpCompositeExtract %3 %186 0 +%191 = OpIMul %3 %189 %190 +%192 = OpIAdd %3 %188 %191 +%193 = OpCompositeExtract %3 %184 1 +%194 = OpCompositeExtract %3 %186 1 +%195 = OpIMul %3 %193 %194 +%187 = OpIAdd %3 %192 %195 +%196 = OpLoad %3 %56 +%197 = OpIAdd %3 %196 %187 +OpStore %56 %197 +%198 = OpLoad %3 %56 +%199 = OpLoad %3 %56 +%200 = OpExtInst %3 %1 SMax %198 %199 +%201 = OpLoad %3 %56 +%202 = OpIAdd %3 %201 %200 +OpStore %56 %202 +%203 = OpLoad %3 %56 +%204 = OpLoad %3 %56 +%205 = OpExtInst %3 %1 SMin %203 %204 +%206 = OpLoad %3 %56 +%207 = OpIAdd %3 %206 %205 +OpStore %56 %207 +%208 = OpLoad %3 %56 +%209 = OpExtInst %3 %1 SSign %208 +%210 = OpLoad %3 %56 +%211 = OpIAdd %3 %210 %209 +OpStore %56 %211 +%212 = OpLoad %3 %56 +OpReturnValue %212 OpFunctionEnd -%213 = OpFunction %4 None %214 -%212 = OpFunctionParameter %4 -%211 = OpLabel -%223 = OpVariable %224 Function %20 -%215 = OpAccessChain %42 %23 %43 -%216 = OpAccessChain %45 %26 %43 -%217 = OpAccessChain %47 %29 %43 -%218 = OpAccessChain %45 %32 %43 -%219 = OpAccessChain %47 %35 %43 -OpBranch %225 -%225 = OpLabel -%226 = OpIAdd %4 %220 %221 -%227 = OpLoad %4 %223 -%228 = OpIAdd %4 %227 %226 -OpStore %223 %228 -%229 = OpLoad %4 %223 -%230 = OpIAdd %4 %229 %222 -%231 = OpLoad %4 %223 -%232 = OpIAdd %4 %231 %230 -OpStore %223 %232 -%233 = OpAccessChain %65 %215 %43 -%234 = OpLoad %5 %233 -%235 = OpLoad %4 %223 -%236 = OpUConvert %5 %235 -%237 = OpIAdd %5 %234 %236 -%238 = OpUConvert %4 %237 -%239 = OpLoad %4 %223 -%240 = OpIAdd %4 %239 %238 -OpStore %223 %240 -%241 = OpAccessChain %74 %215 %75 -%242 = OpLoad %6 %241 -%243 = OpLoad %4 %223 -%244 = OpSConvert %6 %243 -%245 = OpIAdd %6 %242 %244 -%246 = OpUConvert %4 %245 -%247 = OpLoad %4 %223 -%248 = OpIAdd %4 %247 %246 -OpStore %223 %248 -%249 = OpAccessChain %84 %215 %16 -%250 = OpLoad %7 %249 -%251 = OpLoad %4 %223 -%252 = OpConvertUToF %7 %251 -%253 = OpFAdd %7 %250 %252 -%254 = OpConvertFToU %4 %253 -%255 = OpLoad %4 %223 -%256 = OpIAdd %4 %255 %254 -OpStore %223 %256 -%257 = OpAccessChain %101 %215 %102 -%258 = OpLoad %4 %257 -%259 = OpCompositeConstruct %9 %258 %258 %258 -%260 = OpCompositeExtract %4 %259 2 -%261 = OpLoad %4 %223 -%262 = OpIAdd %4 %261 %260 -OpStore %223 %262 -%263 = OpAccessChain %93 %215 %94 -%264 = OpLoad %3 %263 -%265 = OpBitcast %4 %264 -%266 = OpLoad %4 %223 -%267 = OpIAdd %4 %266 %265 -OpStore %223 %267 -%268 = OpAccessChain %140 %215 %141 -%269 = OpLoad %11 %268 -%270 = OpBitcast %8 %269 -%271 = OpCompositeExtract %4 %270 1 -%272 = OpLoad %4 %223 -%273 = OpIAdd %4 %272 %271 -OpStore %223 %273 -%274 = OpAccessChain %149 %215 %150 -%275 = OpLoad %12 %274 -%276 = OpBitcast %9 %275 -%277 = OpCompositeExtract %4 %276 2 -%278 = OpLoad %4 %223 -%279 = OpIAdd %4 %278 %277 -OpStore %223 %279 -%280 = OpAccessChain %158 %215 %159 -%281 = OpLoad %13 %280 -%282 = OpBitcast %10 %281 -%283 = OpCompositeExtract %4 %282 3 -%284 = OpLoad %4 %223 -%285 = OpIAdd %4 %284 %283 -OpStore %223 %285 -%287 = OpAccessChain %101 %215 %102 -%288 = OpLoad %4 %287 -%289 = OpAccessChain %286 %216 %102 -%290 = OpLoad %4 %289 -%291 = OpIAdd %4 %288 %290 -%292 = OpAccessChain %286 %218 %102 -OpStore %292 %291 -%294 = OpAccessChain %108 %215 %109 -%295 = OpLoad %8 %294 -%296 = OpAccessChain %293 %216 %109 -%297 = OpLoad %8 %296 -%298 = OpIAdd %8 %295 %297 -%299 = OpAccessChain %293 %218 %109 -OpStore %299 %298 -%301 = OpAccessChain %116 %215 %117 -%302 = OpLoad %9 %301 -%303 = OpAccessChain %300 %216 %117 -%304 = OpLoad %9 %303 -%305 = OpIAdd %9 %302 %304 -%306 = OpAccessChain %300 %218 %117 -OpStore %306 %305 -%308 = OpAccessChain %124 %215 %125 -%309 = OpLoad %10 %308 -%310 = OpAccessChain %307 %216 %125 -%311 = OpLoad %10 %310 -%312 = OpIAdd %10 %309 %311 -%313 = OpAccessChain %307 %218 %125 -OpStore %313 %312 -%315 = OpAccessChain %314 %217 %43 -%316 = OpLoad %15 %315 -%317 = OpAccessChain %314 %219 %43 -OpStore %317 %316 -%318 = OpLoad %4 %223 -%319 = OpCopyObject %4 %318 -%320 = OpLoad %4 %223 -%321 = OpIAdd %4 %320 %319 -OpStore %223 %321 -%322 = OpLoad %4 %223 -%323 = OpLoad %4 %223 -%324 = OpLoad %4 %223 -%326 = OpExtInst %4 %1 UMax %322 %323 -%325 = OpExtInst %4 %1 UMin %326 %324 -%327 = OpLoad %4 %223 -%328 = OpIAdd %4 %327 %325 -OpStore %223 %328 -%329 = OpLoad %4 %223 -%330 = OpCompositeConstruct %8 %329 %329 -%331 = OpLoad %4 %223 -%332 = OpCompositeConstruct %8 %331 %331 -%335 = OpCompositeExtract %4 %330 0 -%336 = OpCompositeExtract %4 %332 0 -%337 = OpIMul %4 %335 %336 -%338 = OpIAdd %4 %334 %337 -%339 = OpCompositeExtract %4 %330 1 -%340 = OpCompositeExtract %4 %332 1 -%341 = OpIMul %4 %339 %340 -%333 = OpIAdd %4 %338 %341 -%342 = OpLoad %4 %223 -%343 = OpIAdd %4 %342 %333 -OpStore %223 %343 -%344 = OpLoad %4 %223 -%345 = OpLoad %4 %223 -%346 = OpExtInst %4 %1 UMax %344 %345 -%347 = OpLoad %4 %223 -%348 = OpIAdd %4 %347 %346 -OpStore %223 %348 -%349 = OpLoad %4 %223 -%350 = OpLoad %4 %223 -%351 = OpExtInst %4 %1 UMin %349 %350 -%352 = OpLoad %4 %223 -%353 = OpIAdd %4 %352 %351 -OpStore %223 %353 -%354 = OpLoad %4 %223 -OpReturnValue %354 +%215 = OpFunction %4 None %216 +%214 = OpFunctionParameter %4 +%213 = OpLabel +%225 = OpVariable %226 Function %20 +%217 = OpAccessChain %42 %23 %43 +%218 = OpAccessChain %45 %26 %43 +%219 = OpAccessChain %47 %29 %43 +%220 = OpAccessChain %45 %32 %43 +%221 = OpAccessChain %47 %35 %43 +OpBranch %227 +%227 = OpLabel +%228 = OpIAdd %4 %222 %223 +%229 = OpISub %4 %228 %223 +%230 = OpLoad %4 %225 +%231 = OpIAdd %4 %230 %229 +OpStore %225 %231 +%232 = OpLoad %4 %225 +%233 = OpIAdd %4 %232 %224 +%234 = OpLoad %4 %225 +%235 = OpIAdd %4 %234 %233 +OpStore %225 %235 +%236 = OpAccessChain %67 %217 %43 +%237 = OpLoad %5 %236 +%238 = OpLoad %4 %225 +%239 = OpUConvert %5 %238 +%240 = OpIAdd %5 %237 %239 +%241 = OpUConvert %4 %240 +%242 = OpLoad %4 %225 +%243 = OpIAdd %4 %242 %241 +OpStore %225 %243 +%244 = OpAccessChain %76 %217 %77 +%245 = OpLoad %6 %244 +%246 = OpLoad %4 %225 +%247 = OpSConvert %6 %246 +%248 = OpIAdd %6 %245 %247 +%249 = OpUConvert %4 %248 +%250 = OpLoad %4 %225 +%251 = OpIAdd %4 %250 %249 +OpStore %225 %251 +%252 = OpAccessChain %86 %217 %16 +%253 = OpLoad %7 %252 +%254 = OpLoad %4 %225 +%255 = OpConvertUToF %7 %254 +%256 = OpFAdd %7 %253 %255 +%257 = OpConvertFToU %4 %256 +%258 = OpLoad %4 %225 +%259 = OpIAdd %4 %258 %257 +OpStore %225 %259 +%260 = OpAccessChain %103 %217 %104 +%261 = OpLoad %4 %260 +%262 = OpCompositeConstruct %9 %261 %261 %261 +%263 = OpCompositeExtract %4 %262 2 +%264 = OpLoad %4 %225 +%265 = OpIAdd %4 %264 %263 +OpStore %225 %265 +%266 = OpAccessChain %95 %217 %96 +%267 = OpLoad %3 %266 +%268 = OpBitcast %4 %267 +%269 = OpLoad %4 %225 +%270 = OpIAdd %4 %269 %268 +OpStore %225 %270 +%271 = OpAccessChain %142 %217 %143 +%272 = OpLoad %11 %271 +%273 = OpBitcast %8 %272 +%274 = OpCompositeExtract %4 %273 1 +%275 = OpLoad %4 %225 +%276 = OpIAdd %4 %275 %274 +OpStore %225 %276 +%277 = OpAccessChain %151 %217 %152 +%278 = OpLoad %12 %277 +%279 = OpBitcast %9 %278 +%280 = OpCompositeExtract %4 %279 2 +%281 = OpLoad %4 %225 +%282 = OpIAdd %4 %281 %280 +OpStore %225 %282 +%283 = OpAccessChain %160 %217 %161 +%284 = OpLoad %13 %283 +%285 = OpBitcast %10 %284 +%286 = OpCompositeExtract %4 %285 3 +%287 = OpLoad %4 %225 +%288 = OpIAdd %4 %287 %286 +OpStore %225 %288 +%290 = OpAccessChain %103 %217 %104 +%291 = OpLoad %4 %290 +%292 = OpAccessChain %289 %218 %104 +%293 = OpLoad %4 %292 +%294 = OpIAdd %4 %291 %293 +%295 = OpAccessChain %289 %220 %104 +OpStore %295 %294 +%297 = OpAccessChain %110 %217 %111 +%298 = OpLoad %8 %297 +%299 = OpAccessChain %296 %218 %111 +%300 = OpLoad %8 %299 +%301 = OpIAdd %8 %298 %300 +%302 = OpAccessChain %296 %220 %111 +OpStore %302 %301 +%304 = OpAccessChain %118 %217 %119 +%305 = OpLoad %9 %304 +%306 = OpAccessChain %303 %218 %119 +%307 = OpLoad %9 %306 +%308 = OpIAdd %9 %305 %307 +%309 = OpAccessChain %303 %220 %119 +OpStore %309 %308 +%311 = OpAccessChain %126 %217 %127 +%312 = OpLoad %10 %311 +%313 = OpAccessChain %310 %218 %127 +%314 = OpLoad %10 %313 +%315 = OpIAdd %10 %312 %314 +%316 = OpAccessChain %310 %220 %127 +OpStore %316 %315 +%318 = OpAccessChain %317 %219 %43 +%319 = OpLoad %15 %318 +%320 = OpAccessChain %317 %221 %43 +OpStore %320 %319 +%321 = OpLoad %4 %225 +%322 = OpCopyObject %4 %321 +%323 = OpLoad %4 %225 +%324 = OpIAdd %4 %323 %322 +OpStore %225 %324 +%325 = OpLoad %4 %225 +%326 = OpLoad %4 %225 +%327 = OpLoad %4 %225 +%329 = OpExtInst %4 %1 UMax %325 %326 +%328 = OpExtInst %4 %1 UMin %329 %327 +%330 = OpLoad %4 %225 +%331 = OpIAdd %4 %330 %328 +OpStore %225 %331 +%332 = OpLoad %4 %225 +%333 = OpCompositeConstruct %8 %332 %332 +%334 = OpLoad %4 %225 +%335 = OpCompositeConstruct %8 %334 %334 +%338 = OpCompositeExtract %4 %333 0 +%339 = OpCompositeExtract %4 %335 0 +%340 = OpIMul %4 %338 %339 +%341 = OpIAdd %4 %337 %340 +%342 = OpCompositeExtract %4 %333 1 +%343 = OpCompositeExtract %4 %335 1 +%344 = OpIMul %4 %342 %343 +%336 = OpIAdd %4 %341 %344 +%345 = OpLoad %4 %225 +%346 = OpIAdd %4 %345 %336 +OpStore %225 %346 +%347 = OpLoad %4 %225 +%348 = OpLoad %4 %225 +%349 = OpExtInst %4 %1 UMax %347 %348 +%350 = OpLoad %4 %225 +%351 = OpIAdd %4 %350 %349 +OpStore %225 %351 +%352 = OpLoad %4 %225 +%353 = OpLoad %4 %225 +%354 = OpExtInst %4 %1 UMin %352 %353 +%355 = OpLoad %4 %225 +%356 = OpIAdd %4 %355 %354 +OpStore %225 %356 +%357 = OpLoad %4 %225 +OpReturnValue %357 OpFunctionEnd -%356 = OpFunction %2 None %357 -%355 = OpLabel -%358 = OpAccessChain %42 %23 %43 -%359 = OpAccessChain %45 %26 %43 -%360 = OpAccessChain %47 %29 %43 -%361 = OpAccessChain %45 %32 %43 -%362 = OpAccessChain %47 %35 %43 -OpBranch %365 -%365 = OpLabel -%366 = OpFunctionCall %4 %213 %363 -%367 = OpFunctionCall %3 %40 %364 -%368 = OpBitcast %4 %367 -%369 = OpIAdd %4 %366 %368 -%371 = OpAccessChain %286 %361 %370 -OpStore %371 %369 +%359 = OpFunction %2 None %360 +%358 = OpLabel +%361 = OpAccessChain %42 %23 %43 +%362 = OpAccessChain %45 %26 %43 +%363 = OpAccessChain %47 %29 %43 +%364 = OpAccessChain %45 %32 %43 +%365 = OpAccessChain %47 %35 %43 +OpBranch %368 +%368 = OpLabel +%369 = OpFunctionCall %4 %215 %366 +%370 = OpFunctionCall %3 %40 %367 +%371 = OpBitcast %4 %370 +%372 = OpIAdd %4 %369 %371 +%374 = OpAccessChain %289 %364 %373 +OpStore %374 %372 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/int64.wgsl b/naga/tests/out/wgsl/int64.wgsl index f378bef20..19d383eba 100644 --- a/naga/tests/out/wgsl/int64.wgsl +++ b/naga/tests/out/wgsl/int64.wgsl @@ -35,150 +35,150 @@ var output_arrays: StorageCompatible; fn int64_function(x: i64) -> i64 { var val: i64 = 20li; - let _e6 = val; - val = (_e6 + (31li - 1002003004005006li)); let _e8 = val; - let _e11 = val; - val = (_e11 + (_e8 + 5li)); - let _e15 = input_uniform.val_u32_; - let _e16 = val; - let _e20 = val; - val = (_e20 + i64((_e15 + u32(_e16)))); - let _e24 = input_uniform.val_i32_; - let _e25 = val; - let _e29 = val; - val = (_e29 + i64((_e24 + i32(_e25)))); - let _e33 = input_uniform.val_f32_; - let _e34 = val; - let _e38 = val; - val = (_e38 + i64((_e33 + f32(_e34)))); - let _e42 = input_uniform.val_i64_; - let _e45 = val; - val = (_e45 + vec3(_e42).z); - let _e49 = input_uniform.val_u64_; - let _e51 = val; - val = (_e51 + bitcast(_e49)); - let _e55 = input_uniform.val_u64_2_; - let _e58 = val; - val = (_e58 + bitcast>(_e55).y); - let _e62 = input_uniform.val_u64_3_; - let _e65 = val; - val = (_e65 + bitcast>(_e62).z); - let _e69 = input_uniform.val_u64_4_; - let _e72 = val; - val = (_e72 + bitcast>(_e69).w); - let _e78 = input_uniform.val_i64_; - let _e81 = input_storage.val_i64_; - output.val_i64_ = (_e78 + _e81); - let _e87 = input_uniform.val_i64_2_; - let _e90 = input_storage.val_i64_2_; - output.val_i64_2_ = (_e87 + _e90); - let _e96 = input_uniform.val_i64_3_; - let _e99 = input_storage.val_i64_3_; - output.val_i64_3_ = (_e96 + _e99); - let _e105 = input_uniform.val_i64_4_; - let _e108 = input_storage.val_i64_4_; - output.val_i64_4_ = (_e105 + _e108); - let _e114 = input_arrays.val_i64_array_2_; - output_arrays.val_i64_array_2_ = _e114; - let _e115 = val; + val = (_e8 + ((31li - 1002003004005006li) + -9223372036854775807li)); + let _e10 = val; + let _e13 = val; + val = (_e13 + (_e10 + 5li)); + let _e17 = input_uniform.val_u32_; + let _e18 = val; + let _e22 = val; + val = (_e22 + i64((_e17 + u32(_e18)))); + let _e26 = input_uniform.val_i32_; + let _e27 = val; + let _e31 = val; + val = (_e31 + i64((_e26 + i32(_e27)))); + let _e35 = input_uniform.val_f32_; + let _e36 = val; + let _e40 = val; + val = (_e40 + i64((_e35 + f32(_e36)))); + let _e44 = input_uniform.val_i64_; + let _e47 = val; + val = (_e47 + vec3(_e44).z); + let _e51 = input_uniform.val_u64_; + let _e53 = val; + val = (_e53 + bitcast(_e51)); + let _e57 = input_uniform.val_u64_2_; + let _e60 = val; + val = (_e60 + bitcast>(_e57).y); + let _e64 = input_uniform.val_u64_3_; + let _e67 = val; + val = (_e67 + bitcast>(_e64).z); + let _e71 = input_uniform.val_u64_4_; + let _e74 = val; + val = (_e74 + bitcast>(_e71).w); + let _e80 = input_uniform.val_i64_; + let _e83 = input_storage.val_i64_; + output.val_i64_ = (_e80 + _e83); + let _e89 = input_uniform.val_i64_2_; + let _e92 = input_storage.val_i64_2_; + output.val_i64_2_ = (_e89 + _e92); + let _e98 = input_uniform.val_i64_3_; + let _e101 = input_storage.val_i64_3_; + output.val_i64_3_ = (_e98 + _e101); + let _e107 = input_uniform.val_i64_4_; + let _e110 = input_storage.val_i64_4_; + output.val_i64_4_ = (_e107 + _e110); + let _e116 = input_arrays.val_i64_array_2_; + output_arrays.val_i64_array_2_ = _e116; let _e117 = val; - val = (_e117 + abs(_e115)); let _e119 = val; - let _e120 = val; + val = (_e119 + abs(_e117)); let _e121 = val; + let _e122 = val; let _e123 = val; - val = (_e123 + clamp(_e119, _e120, _e121)); let _e125 = val; + val = (_e125 + clamp(_e121, _e122, _e123)); let _e127 = val; - let _e130 = val; - val = (_e130 + dot(vec2(_e125), vec2(_e127))); + let _e129 = val; let _e132 = val; - let _e133 = val; + val = (_e132 + dot(vec2(_e127), vec2(_e129))); + let _e134 = val; let _e135 = val; - val = (_e135 + max(_e132, _e133)); let _e137 = val; - let _e138 = val; + val = (_e137 + max(_e134, _e135)); + let _e139 = val; let _e140 = val; - val = (_e140 + min(_e137, _e138)); let _e142 = val; + val = (_e142 + min(_e139, _e140)); let _e144 = val; - val = (_e144 + sign(_e142)); let _e146 = val; - return _e146; + val = (_e146 + sign(_e144)); + let _e148 = val; + return _e148; } fn uint64_function(x_1: u64) -> u64 { var val_1: u64 = 20lu; - let _e6 = val_1; - val_1 = (_e6 + (31lu + 1002003004005006lu)); let _e8 = val_1; - let _e11 = val_1; - val_1 = (_e11 + (_e8 + 5lu)); - let _e15 = input_uniform.val_u32_; - let _e16 = val_1; - let _e20 = val_1; - val_1 = (_e20 + u64((_e15 + u32(_e16)))); - let _e24 = input_uniform.val_i32_; - let _e25 = val_1; - let _e29 = val_1; - val_1 = (_e29 + u64((_e24 + i32(_e25)))); - let _e33 = input_uniform.val_f32_; - let _e34 = val_1; - let _e38 = val_1; - val_1 = (_e38 + u64((_e33 + f32(_e34)))); - let _e42 = input_uniform.val_u64_; - let _e45 = val_1; - val_1 = (_e45 + vec3(_e42).z); - let _e49 = input_uniform.val_i64_; - let _e51 = val_1; - val_1 = (_e51 + bitcast(_e49)); - let _e55 = input_uniform.val_i64_2_; - let _e58 = val_1; - val_1 = (_e58 + bitcast>(_e55).y); - let _e62 = input_uniform.val_i64_3_; - let _e65 = val_1; - val_1 = (_e65 + bitcast>(_e62).z); - let _e69 = input_uniform.val_i64_4_; - let _e72 = val_1; - val_1 = (_e72 + bitcast>(_e69).w); - let _e78 = input_uniform.val_u64_; - let _e81 = input_storage.val_u64_; - output.val_u64_ = (_e78 + _e81); - let _e87 = input_uniform.val_u64_2_; - let _e90 = input_storage.val_u64_2_; - output.val_u64_2_ = (_e87 + _e90); - let _e96 = input_uniform.val_u64_3_; - let _e99 = input_storage.val_u64_3_; - output.val_u64_3_ = (_e96 + _e99); - let _e105 = input_uniform.val_u64_4_; - let _e108 = input_storage.val_u64_4_; - output.val_u64_4_ = (_e105 + _e108); - let _e114 = input_arrays.val_u64_array_2_; - output_arrays.val_u64_array_2_ = _e114; - let _e115 = val_1; + val_1 = (_e8 + ((31lu + 18446744073709551615lu) - 18446744073709551615lu)); + let _e10 = val_1; + let _e13 = val_1; + val_1 = (_e13 + (_e10 + 5lu)); + let _e17 = input_uniform.val_u32_; + let _e18 = val_1; + let _e22 = val_1; + val_1 = (_e22 + u64((_e17 + u32(_e18)))); + let _e26 = input_uniform.val_i32_; + let _e27 = val_1; + let _e31 = val_1; + val_1 = (_e31 + u64((_e26 + i32(_e27)))); + let _e35 = input_uniform.val_f32_; + let _e36 = val_1; + let _e40 = val_1; + val_1 = (_e40 + u64((_e35 + f32(_e36)))); + let _e44 = input_uniform.val_u64_; + let _e47 = val_1; + val_1 = (_e47 + vec3(_e44).z); + let _e51 = input_uniform.val_i64_; + let _e53 = val_1; + val_1 = (_e53 + bitcast(_e51)); + let _e57 = input_uniform.val_i64_2_; + let _e60 = val_1; + val_1 = (_e60 + bitcast>(_e57).y); + let _e64 = input_uniform.val_i64_3_; + let _e67 = val_1; + val_1 = (_e67 + bitcast>(_e64).z); + let _e71 = input_uniform.val_i64_4_; + let _e74 = val_1; + val_1 = (_e74 + bitcast>(_e71).w); + let _e80 = input_uniform.val_u64_; + let _e83 = input_storage.val_u64_; + output.val_u64_ = (_e80 + _e83); + let _e89 = input_uniform.val_u64_2_; + let _e92 = input_storage.val_u64_2_; + output.val_u64_2_ = (_e89 + _e92); + let _e98 = input_uniform.val_u64_3_; + let _e101 = input_storage.val_u64_3_; + output.val_u64_3_ = (_e98 + _e101); + let _e107 = input_uniform.val_u64_4_; + let _e110 = input_storage.val_u64_4_; + output.val_u64_4_ = (_e107 + _e110); + let _e116 = input_arrays.val_u64_array_2_; + output_arrays.val_u64_array_2_ = _e116; let _e117 = val_1; - val_1 = (_e117 + abs(_e115)); let _e119 = val_1; - let _e120 = val_1; + val_1 = (_e119 + abs(_e117)); let _e121 = val_1; + let _e122 = val_1; let _e123 = val_1; - val_1 = (_e123 + clamp(_e119, _e120, _e121)); let _e125 = val_1; + val_1 = (_e125 + clamp(_e121, _e122, _e123)); let _e127 = val_1; - let _e130 = val_1; - val_1 = (_e130 + dot(vec2(_e125), vec2(_e127))); + let _e129 = val_1; let _e132 = val_1; - let _e133 = val_1; + val_1 = (_e132 + dot(vec2(_e127), vec2(_e129))); + let _e134 = val_1; let _e135 = val_1; - val_1 = (_e135 + max(_e132, _e133)); let _e137 = val_1; - let _e138 = val_1; + val_1 = (_e137 + max(_e134, _e135)); + let _e139 = val_1; let _e140 = val_1; - val_1 = (_e140 + min(_e137, _e138)); let _e142 = val_1; - return _e142; + val_1 = (_e142 + min(_e139, _e140)); + let _e144 = val_1; + return _e144; } @compute @workgroup_size(1, 1, 1)