diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 4dfb5b8892..2935a6ad3c 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -1494,13 +1494,31 @@ impl Writer { write!(self.out, "{value}u")?; } crate::Literal::I32(value) => { - write!(self.out, "{value}")?; + // `-2147483648` is parsed as unary negation of positive 2147483648. + // 2147483648 is too large for int32_t meaning the expression gets + // promoted to a int64_t which is not our intention. Avoid this by instead + // using `-2147483647 - 1`. + if value == i32::MIN { + write!(self.out, "({} - 1)", value + 1)?; + } else { + write!(self.out, "{value}")?; + } } crate::Literal::U64(value) => { write!(self.out, "{value}uL")?; } crate::Literal::I64(value) => { - write!(self.out, "{value}L")?; + // `-9223372036854775808` is parsed as unary negation of positive + // 9223372036854775808. 9223372036854775808 is too large for int64_t + // causing Metal to emit a `-Wconstant-conversion` warning, and change the + // value to `-9223372036854775808`. Which would then be negated, possibly + // causing undefined behaviour. Avoid this by instead using + // `-9223372036854775808L - 1L`. + if value == i64::MIN { + write!(self.out, "({}L - 1L)", value + 1)?; + } else { + write!(self.out, "{value}L")?; + } } crate::Literal::Bool(value) => { write!(self.out, "{value}")?; diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index af4923c65a..a43c320553 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -1139,12 +1139,12 @@ impl Writer { crate::Literal::Bool(value) => write!(self.out, "{value}")?, crate::Literal::F64(value) => write!(self.out, "{value:?}lf")?, crate::Literal::I64(value) => { - // `-9223372036854775808li` is not valid WGSL. Nor can we use the AbstractInt - // trick above, as AbstractInt also cannot represent `9223372036854775808`. - // The most negative `i64` value can only be expressed in WGSL using - // subtracting 1 from the second most negative value. + // `-9223372036854775808li` is not valid WGSL. Nor can we simply use the + // AbstractInt trick above, as AbstractInt also cannot represent + // `9223372036854775808`. Instead construct the second most negative + // AbstractInt, subtract one from it, then cast to i64. if value == i64::MIN { - write!(self.out, "{}li - 1li", value + 1)?; + write!(self.out, "i64({} - 1)", value + 1)?; } else { write!(self.out, "{value}li")?; } diff --git a/naga/tests/in/wgsl/int64.wgsl b/naga/tests/in/wgsl/int64.wgsl index f9766cfdc7..86ce73261c 100644 --- a/naga/tests/in/wgsl/int64.wgsl +++ b/naga/tests/in/wgsl/int64.wgsl @@ -56,7 +56,7 @@ fn int64_function(x: i64) -> i64 { val += bitcast>(input_uniform.val_u64_3).z; val += bitcast>(input_uniform.val_u64_4).w; // Most negative i64 - val += -9223372036854775807li - 1li; + val += i64(-9223372036854775807 - 1); // Reading/writing to a uniform/storage buffer output.val_i64 = input_uniform.val_i64 + input_storage.val_i64; diff --git a/naga/tests/in/wgsl/operators.wgsl b/naga/tests/in/wgsl/operators.wgsl index 5b621988e9..6190dd47f8 100644 --- a/naga/tests/in/wgsl/operators.wgsl +++ b/naga/tests/in/wgsl/operators.wgsl @@ -160,6 +160,14 @@ fn arithmetic() { let mul_vector1 = vec3f(two_f) * mat4x3f(); let mul = mat4x3() * mat3x4(); + + // Arithmetic involving the minimum value i32 literal. What we're really testing here + // is how this literal is expressed by Naga backends. eg in Metal, `-2147483648` is + // silently promoted to a `long` which we don't want. The addition ensures this would + // be caught as a compiler error, as we bitcast the operands to unsigned which fails + // if the expression's type has an unexpected width. + var prevent_const_eval: i32; + var wgpu_7437 = prevent_const_eval + -2147483648; } fn bit() { diff --git a/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl b/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl index 3e38d43202..7189ace52a 100644 --- a/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl +++ b/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl @@ -59,6 +59,8 @@ void logical() { } void arithmetic() { + int prevent_const_eval = 0; + int wgpu_7437_ = 0; float neg0_1 = -(1.0); ivec2 neg1_1 = -(ivec2(1)); vec2 neg2_ = -(vec2(1.0)); @@ -131,6 +133,8 @@ void arithmetic() { vec3 mul_vector0_ = (mat4x3(0.0) * vec4(1.0)); vec4 mul_vector1_ = (vec3(2.0) * mat4x3(0.0)); mat3x3 mul = (mat4x3(0.0) * mat3x4(0.0)); + int _e175 = prevent_const_eval; + wgpu_7437_ = (_e175 + -2147483648); return; } diff --git a/naga/tests/out/hlsl/wgsl-int64.hlsl b/naga/tests/out/hlsl/wgsl-int64.hlsl index ef13f9c623..6705e77e71 100644 --- a/naga/tests/out/hlsl/wgsl-int64.hlsl +++ b/naga/tests/out/hlsl/wgsl-int64.hlsl @@ -99,51 +99,51 @@ int64_t int64_function(int64_t x) uint64_t4 _e71 = input_uniform.val_u64_4_; int64_t _e74 = val; val = (_e74 + _e71.w); - int64_t _e79 = val; - val = (_e79 + (-9223372036854775807L - 1L)); - int64_t _e85 = input_uniform.val_i64_; - int64_t _e88 = input_storage.Load(128); - output.Store(128, (_e85 + _e88)); - int64_t2 _e94 = input_uniform.val_i64_2_; - int64_t2 _e97 = input_storage.Load(144); - output.Store(144, (_e94 + _e97)); - int64_t3 _e103 = input_uniform.val_i64_3_; - int64_t3 _e106 = input_storage.Load(160); - output.Store(160, (_e103 + _e106)); - int64_t4 _e112 = input_uniform.val_i64_4_; - int64_t4 _e115 = input_storage.Load(192); - output.Store(192, (_e112 + _e115)); - int64_t _e121[2] = Constructarray2_int64_t_(input_arrays.Load(16+0), input_arrays.Load(16+8)); + int64_t _e77 = val; + val = (_e77 + -9223372036854775808L); + int64_t _e83 = input_uniform.val_i64_; + int64_t _e86 = input_storage.Load(128); + output.Store(128, (_e83 + _e86)); + int64_t2 _e92 = input_uniform.val_i64_2_; + int64_t2 _e95 = input_storage.Load(144); + output.Store(144, (_e92 + _e95)); + int64_t3 _e101 = input_uniform.val_i64_3_; + int64_t3 _e104 = input_storage.Load(160); + output.Store(160, (_e101 + _e104)); + int64_t4 _e110 = input_uniform.val_i64_4_; + int64_t4 _e113 = input_storage.Load(192); + output.Store(192, (_e110 + _e113)); + int64_t _e119[2] = Constructarray2_int64_t_(input_arrays.Load(16+0), input_arrays.Load(16+8)); { - int64_t _value2[2] = _e121; + int64_t _value2[2] = _e119; output_arrays.Store(16+0, _value2[0]); output_arrays.Store(16+8, _value2[1]); } + int64_t _e120 = val; int64_t _e122 = val; + val = (_e122 + abs(_e120)); int64_t _e124 = val; - val = (_e124 + abs(_e122)); + int64_t _e125 = val; int64_t _e126 = val; - int64_t _e127 = val; int64_t _e128 = val; + val = (_e128 + clamp(_e124, _e125, _e126)); int64_t _e130 = val; - val = (_e130 + clamp(_e126, _e127, _e128)); int64_t _e132 = val; - int64_t _e134 = val; + int64_t _e135 = val; + val = (_e135 + dot((_e130).xx, (_e132).xx)); int64_t _e137 = val; - val = (_e137 + dot((_e132).xx, (_e134).xx)); - int64_t _e139 = val; + int64_t _e138 = val; int64_t _e140 = val; + val = (_e140 + max(_e137, _e138)); int64_t _e142 = val; - val = (_e142 + max(_e139, _e140)); - int64_t _e144 = val; + int64_t _e143 = val; int64_t _e145 = val; + val = (_e145 + min(_e142, _e143)); int64_t _e147 = val; - val = (_e147 + min(_e144, _e145)); int64_t _e149 = val; + val = (_e149 + sign(_e147)); int64_t _e151 = val; - val = (_e151 + sign(_e149)); - int64_t _e153 = val; - return _e153; + return _e151; } uint64_t naga_f2u64(float value) { diff --git a/naga/tests/out/hlsl/wgsl-operators.hlsl b/naga/tests/out/hlsl/wgsl-operators.hlsl index 0e74dc54fa..b19e526ff0 100644 --- a/naga/tests/out/hlsl/wgsl-operators.hlsl +++ b/naga/tests/out/hlsl/wgsl-operators.hlsl @@ -121,6 +121,9 @@ float3x4 ZeroValuefloat3x4() { void arithmetic() { + int prevent_const_eval = (int)0; + int wgpu_7437_ = (int)0; + float neg0_1 = -(1.0); int2 neg1_1 = naga_neg((int(1)).xx); float2 neg2_ = -((1.0).xx); @@ -193,6 +196,8 @@ void arithmetic() float3 mul_vector0_ = mul((1.0).xxxx, ZeroValuefloat4x3()); float4 mul_vector1_ = mul(ZeroValuefloat4x3(), (2.0).xxx); float3x3 mul_ = mul(ZeroValuefloat3x4(), ZeroValuefloat4x3()); + int _e175 = prevent_const_eval; + wgpu_7437_ = asint(asuint(_e175) + asuint(int(-2147483648))); return; } diff --git a/naga/tests/out/msl/wgsl-abstract-types-operators.msl b/naga/tests/out/msl/wgsl-abstract-types-operators.msl index e2603bedcc..b8c66aa590 100644 --- a/naga/tests/out/msl/wgsl-abstract-types-operators.msl +++ b/naga/tests/out/msl/wgsl-abstract-types-operators.msl @@ -26,7 +26,7 @@ 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 int least_i32_ = (-2147483647 - 1); constant float least_f32_ = -340282350000000000000000000000000000000.0; constant int shl_iaiai = 4; constant int shl_iai_u_1 = 4; @@ -36,7 +36,7 @@ constant int shr_iaiai = 0; constant int shr_iai_u_1 = 0; constant uint shr_uaiai = 0u; constant uint shr_uai_u = 0u; -constant int wgpu_4492_ = -2147483648; +constant int wgpu_4492_ = (-2147483647 - 1); void runtime_values( ) { diff --git a/naga/tests/out/msl/wgsl-conversion-float-to-int-no-f64.msl b/naga/tests/out/msl/wgsl-conversion-float-to-int-no-f64.msl index b6a5f7521f..8b103e7842 100644 --- a/naga/tests/out/msl/wgsl-conversion-float-to-int-no-f64.msl +++ b/naga/tests/out/msl/wgsl-conversion-float-to-int-no-f64.msl @@ -19,19 +19,19 @@ void test_const_eval( long max_f16_to_i64_ = 65504L; ulong min_f16_to_u64_ = 0uL; ulong max_f16_to_u64_ = 65504uL; - int min_f32_to_i32_ = -2147483648; + int min_f32_to_i32_ = (-2147483647 - 1); int max_f32_to_i32_ = 2147483520; uint min_f32_to_u32_ = 0u; uint max_f32_to_u32_ = 4294967040u; - long min_f32_to_i64_ = -9223372036854775808L; + long min_f32_to_i64_ = (-9223372036854775807L - 1L); long max_f32_to_i64_ = 9223371487098961920L; ulong min_f32_to_u64_ = 0uL; ulong max_f32_to_u64_ = 18446742974197923840uL; - int min_abstract_float_to_i32_ = -2147483648; + int min_abstract_float_to_i32_ = (-2147483647 - 1); int max_abstract_float_to_i32_ = 2147483647; uint min_abstract_float_to_u32_ = 0u; uint max_abstract_float_to_u32_ = 4294967295u; - long min_abstract_float_to_i64_ = -9223372036854775808L; + long min_abstract_float_to_i64_ = (-9223372036854775807L - 1L); long max_abstract_float_to_i64_ = 9223372036854774784L; ulong min_abstract_float_to_u64_ = 0uL; ulong max_abstract_float_to_u64_ = 18446744073709549568uL; diff --git a/naga/tests/out/msl/wgsl-int64.msl b/naga/tests/out/msl/wgsl-int64.msl index e026cac919..9b3d50d4b4 100644 --- a/naga/tests/out/msl/wgsl-int64.msl +++ b/naga/tests/out/msl/wgsl-int64.msl @@ -83,49 +83,49 @@ long int64_function( metal::ulong4 _e71 = input_uniform.val_u64_4_; long _e74 = val; val = as_type(as_type(_e74) + as_type(as_type(_e71).w)); - long _e79 = val; - val = as_type(as_type(_e79) + as_type(as_type(as_type(-9223372036854775807L) - as_type(1L)))); - long _e85 = input_uniform.val_i64_; - long _e88 = input_storage.val_i64_; - output.val_i64_ = as_type(as_type(_e85) + as_type(_e88)); - metal::long2 _e94 = input_uniform.val_i64_2_; - metal::long2 _e97 = input_storage.val_i64_2_; - output.val_i64_2_ = as_type(as_type(_e94) + as_type(_e97)); - metal::long3 _e103 = input_uniform.val_i64_3_; - metal::long3 _e106 = input_storage.val_i64_3_; - output.val_i64_3_ = as_type(as_type(_e103) + as_type(_e106)); - metal::long4 _e112 = input_uniform.val_i64_4_; - metal::long4 _e115 = input_storage.val_i64_4_; - output.val_i64_4_ = as_type(as_type(_e112) + as_type(_e115)); - type_12 _e121 = input_arrays.val_i64_array_2_; - output_arrays.val_i64_array_2_ = _e121; + long _e77 = val; + val = as_type(as_type(_e77) + as_type((-9223372036854775807L - 1L))); + long _e83 = input_uniform.val_i64_; + long _e86 = input_storage.val_i64_; + output.val_i64_ = as_type(as_type(_e83) + as_type(_e86)); + metal::long2 _e92 = input_uniform.val_i64_2_; + metal::long2 _e95 = input_storage.val_i64_2_; + output.val_i64_2_ = as_type(as_type(_e92) + as_type(_e95)); + metal::long3 _e101 = input_uniform.val_i64_3_; + metal::long3 _e104 = input_storage.val_i64_3_; + output.val_i64_3_ = as_type(as_type(_e101) + as_type(_e104)); + metal::long4 _e110 = input_uniform.val_i64_4_; + metal::long4 _e113 = input_storage.val_i64_4_; + output.val_i64_4_ = as_type(as_type(_e110) + as_type(_e113)); + type_12 _e119 = input_arrays.val_i64_array_2_; + output_arrays.val_i64_array_2_ = _e119; + long _e120 = val; long _e122 = val; + val = as_type(as_type(_e122) + as_type(naga_abs(_e120))); long _e124 = val; - val = as_type(as_type(_e124) + as_type(naga_abs(_e122))); + long _e125 = val; long _e126 = val; - long _e127 = val; long _e128 = val; + val = as_type(as_type(_e128) + as_type(metal::clamp(_e124, _e125, _e126))); long _e130 = val; - val = as_type(as_type(_e130) + as_type(metal::clamp(_e126, _e127, _e128))); + metal::long2 _e131 = metal::long2(_e130); long _e132 = val; metal::long2 _e133 = metal::long2(_e132); - long _e134 = val; - metal::long2 _e135 = metal::long2(_e134); + long _e135 = val; + val = as_type(as_type(_e135) + as_type(( + _e131.x * _e133.x + _e131.y * _e133.y))); long _e137 = val; - val = as_type(as_type(_e137) + as_type(( + _e133.x * _e135.x + _e133.y * _e135.y))); - long _e139 = val; + long _e138 = val; long _e140 = val; + val = as_type(as_type(_e140) + as_type(metal::max(_e137, _e138))); long _e142 = val; - val = as_type(as_type(_e142) + as_type(metal::max(_e139, _e140))); - long _e144 = val; + long _e143 = val; long _e145 = val; + val = as_type(as_type(_e145) + as_type(metal::min(_e142, _e143))); long _e147 = val; - val = as_type(as_type(_e147) + as_type(metal::min(_e144, _e145))); long _e149 = val; + val = as_type(as_type(_e149) + as_type(metal::select(metal::select(long(-1), long(1), (_e147 > 0)), long(0), (_e147 == 0)))); long _e151 = val; - val = as_type(as_type(_e151) + as_type(metal::select(metal::select(long(-1), long(1), (_e149 > 0)), long(0), (_e149 == 0)))); - long _e153 = val; - return _e153; + return _e151; } ulong naga_f2u64(float value) { diff --git a/naga/tests/out/msl/wgsl-operators.msl b/naga/tests/out/msl/wgsl-operators.msl index 1d2d81834a..383c65713a 100644 --- a/naga/tests/out/msl/wgsl-operators.msl +++ b/naga/tests/out/msl/wgsl-operators.msl @@ -117,6 +117,8 @@ metal::uint2 naga_mod(metal::uint2 lhs, metal::uint2 rhs) { void arithmetic( ) { + int prevent_const_eval = {}; + int wgpu_7437_ = {}; float neg0_1 = -(1.0); metal::int2 neg1_1 = naga_neg(metal::int2(1)); metal::float2 neg2_ = -(metal::float2(1.0)); @@ -189,6 +191,8 @@ void arithmetic( metal::float3 mul_vector0_ = metal::float4x3 {} * metal::float4(1.0); metal::float4 mul_vector1_ = metal::float3(2.0) * metal::float4x3 {}; metal::float3x3 mul = metal::float4x3 {} * metal::float3x4 {}; + int _e175 = prevent_const_eval; + wgpu_7437_ = as_type(as_type(_e175) + as_type((-2147483647 - 1))); return; } diff --git a/naga/tests/out/spv/wgsl-int64.spvasm b/naga/tests/out/spv/wgsl-int64.spvasm index 377247a215..883eb157f5 100644 --- a/naga/tests/out/spv/wgsl-int64.spvasm +++ b/naga/tests/out/spv/wgsl-int64.spvasm @@ -93,23 +93,24 @@ OpMemberDecorate %36 0 Offset 0 %53 = OpConstant %3 1002003004005006 %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 -%92 = OpConstant %7 -9.223372e18 -%93 = OpConstant %7 9.2233715e18 -%98 = OpTypePointer Uniform %3 -%99 = OpConstant %5 7 -%106 = OpTypePointer Uniform %4 -%107 = OpConstant %5 3 -%113 = OpTypePointer Uniform %8 -%114 = OpConstant %5 4 -%121 = OpTypePointer Uniform %9 -%122 = OpConstant %5 5 -%129 = OpTypePointer Uniform %10 -%130 = OpConstant %5 6 +%56 = OpConstant %3 -9223372036854775808 +%58 = OpTypePointer Function %3 +%68 = OpTypePointer Uniform %5 +%77 = OpTypePointer Uniform %6 +%78 = OpConstant %5 1 +%87 = OpTypePointer Uniform %7 +%93 = OpConstant %7 -9.223372e18 +%94 = OpConstant %7 9.2233715e18 +%99 = OpTypePointer Uniform %3 +%100 = OpConstant %5 7 +%107 = OpTypePointer Uniform %4 +%108 = OpConstant %5 3 +%114 = OpTypePointer Uniform %8 +%115 = OpConstant %5 4 +%122 = OpTypePointer Uniform %9 +%123 = OpConstant %5 5 +%130 = OpTypePointer Uniform %10 +%131 = OpConstant %5 6 %140 = OpTypePointer StorageBuffer %3 %147 = OpTypePointer StorageBuffer %11 %148 = OpTypePointer Uniform %11 @@ -142,96 +143,95 @@ OpMemberDecorate %36 0 Offset 0 %40 = OpFunction %3 None %41 %39 = OpFunctionParameter %3 %38 = OpLabel -%56 = OpVariable %57 Function %51 +%57 = OpVariable %58 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 %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 -%94 = OpExtInst %7 %1 FClamp %91 %92 %93 -%95 = OpConvertFToS %3 %94 -%96 = OpLoad %3 %56 -%97 = OpIAdd %3 %96 %95 -OpStore %56 %97 -%100 = OpAccessChain %98 %44 %99 -%101 = OpLoad %3 %100 -%102 = OpCompositeConstruct %12 %101 %101 %101 -%103 = OpCompositeExtract %3 %102 2 -%104 = OpLoad %3 %56 -%105 = OpIAdd %3 %104 %103 -OpStore %56 %105 -%108 = OpAccessChain %106 %44 %107 -%109 = OpLoad %4 %108 -%110 = OpBitcast %3 %109 -%111 = OpLoad %3 %56 -%112 = OpIAdd %3 %111 %110 -OpStore %56 %112 -%115 = OpAccessChain %113 %44 %114 -%116 = OpLoad %8 %115 -%117 = OpBitcast %11 %116 -%118 = OpCompositeExtract %3 %117 1 -%119 = OpLoad %3 %56 -%120 = OpIAdd %3 %119 %118 -OpStore %56 %120 -%123 = OpAccessChain %121 %44 %122 -%124 = OpLoad %9 %123 -%125 = OpBitcast %12 %124 -%126 = OpCompositeExtract %3 %125 2 -%127 = OpLoad %3 %56 -%128 = OpIAdd %3 %127 %126 -OpStore %56 %128 -%131 = OpAccessChain %129 %44 %130 -%132 = OpLoad %10 %131 -%133 = OpBitcast %13 %132 -%134 = OpCompositeExtract %3 %133 3 -%135 = OpLoad %3 %56 -%136 = OpIAdd %3 %135 %134 -OpStore %56 %136 -%137 = OpISub %3 %54 %19 -%138 = OpLoad %3 %56 -%139 = OpIAdd %3 %138 %137 -OpStore %56 %139 -%141 = OpAccessChain %98 %44 %99 +OpBranch %59 +%59 = OpLabel +%60 = OpISub %3 %52 %53 +%61 = OpIAdd %3 %60 %54 +%62 = OpLoad %3 %57 +%63 = OpIAdd %3 %62 %61 +OpStore %57 %63 +%64 = OpLoad %3 %57 +%65 = OpIAdd %3 %64 %55 +%66 = OpLoad %3 %57 +%67 = OpIAdd %3 %66 %65 +OpStore %57 %67 +%69 = OpAccessChain %68 %44 %43 +%70 = OpLoad %5 %69 +%71 = OpLoad %3 %57 +%72 = OpUConvert %5 %71 +%73 = OpIAdd %5 %70 %72 +%74 = OpSConvert %3 %73 +%75 = OpLoad %3 %57 +%76 = OpIAdd %3 %75 %74 +OpStore %57 %76 +%79 = OpAccessChain %77 %44 %78 +%80 = OpLoad %6 %79 +%81 = OpLoad %3 %57 +%82 = OpSConvert %6 %81 +%83 = OpIAdd %6 %80 %82 +%84 = OpSConvert %3 %83 +%85 = OpLoad %3 %57 +%86 = OpIAdd %3 %85 %84 +OpStore %57 %86 +%88 = OpAccessChain %87 %44 %16 +%89 = OpLoad %7 %88 +%90 = OpLoad %3 %57 +%91 = OpConvertSToF %7 %90 +%92 = OpFAdd %7 %89 %91 +%95 = OpExtInst %7 %1 FClamp %92 %93 %94 +%96 = OpConvertFToS %3 %95 +%97 = OpLoad %3 %57 +%98 = OpIAdd %3 %97 %96 +OpStore %57 %98 +%101 = OpAccessChain %99 %44 %100 +%102 = OpLoad %3 %101 +%103 = OpCompositeConstruct %12 %102 %102 %102 +%104 = OpCompositeExtract %3 %103 2 +%105 = OpLoad %3 %57 +%106 = OpIAdd %3 %105 %104 +OpStore %57 %106 +%109 = OpAccessChain %107 %44 %108 +%110 = OpLoad %4 %109 +%111 = OpBitcast %3 %110 +%112 = OpLoad %3 %57 +%113 = OpIAdd %3 %112 %111 +OpStore %57 %113 +%116 = OpAccessChain %114 %44 %115 +%117 = OpLoad %8 %116 +%118 = OpBitcast %11 %117 +%119 = OpCompositeExtract %3 %118 1 +%120 = OpLoad %3 %57 +%121 = OpIAdd %3 %120 %119 +OpStore %57 %121 +%124 = OpAccessChain %122 %44 %123 +%125 = OpLoad %9 %124 +%126 = OpBitcast %12 %125 +%127 = OpCompositeExtract %3 %126 2 +%128 = OpLoad %3 %57 +%129 = OpIAdd %3 %128 %127 +OpStore %57 %129 +%132 = OpAccessChain %130 %44 %131 +%133 = OpLoad %10 %132 +%134 = OpBitcast %13 %133 +%135 = OpCompositeExtract %3 %134 3 +%136 = OpLoad %3 %57 +%137 = OpIAdd %3 %136 %135 +OpStore %57 %137 +%138 = OpLoad %3 %57 +%139 = OpIAdd %3 %138 %56 +OpStore %57 %139 +%141 = OpAccessChain %99 %44 %100 %142 = OpLoad %3 %141 -%143 = OpAccessChain %140 %46 %99 +%143 = OpAccessChain %140 %46 %100 %144 = OpLoad %3 %143 %145 = OpIAdd %3 %142 %144 -%146 = OpAccessChain %140 %49 %99 +%146 = OpAccessChain %140 %49 %100 OpStore %146 %145 %150 = OpAccessChain %148 %44 %149 %151 = OpLoad %11 %150 @@ -254,26 +254,26 @@ OpStore %164 %163 %172 = OpIAdd %13 %169 %171 %173 = OpAccessChain %165 %49 %167 OpStore %173 %172 -%175 = OpAccessChain %174 %48 %77 +%175 = OpAccessChain %174 %48 %78 %176 = OpLoad %17 %175 -%177 = OpAccessChain %174 %50 %77 +%177 = OpAccessChain %174 %50 %78 OpStore %177 %176 -%178 = OpLoad %3 %56 +%178 = OpLoad %3 %57 %179 = OpExtInst %3 %1 SAbs %178 -%180 = OpLoad %3 %56 +%180 = OpLoad %3 %57 %181 = OpIAdd %3 %180 %179 -OpStore %56 %181 -%182 = OpLoad %3 %56 -%183 = OpLoad %3 %56 -%184 = OpLoad %3 %56 +OpStore %57 %181 +%182 = OpLoad %3 %57 +%183 = OpLoad %3 %57 +%184 = OpLoad %3 %57 %186 = OpExtInst %3 %1 SMax %182 %183 %185 = OpExtInst %3 %1 SMin %186 %184 -%187 = OpLoad %3 %56 +%187 = OpLoad %3 %57 %188 = OpIAdd %3 %187 %185 -OpStore %56 %188 -%189 = OpLoad %3 %56 +OpStore %57 %188 +%189 = OpLoad %3 %57 %190 = OpCompositeConstruct %11 %189 %189 -%191 = OpLoad %3 %56 +%191 = OpLoad %3 %57 %192 = OpCompositeConstruct %11 %191 %191 %195 = OpCompositeExtract %3 %190 0 %196 = OpCompositeExtract %3 %192 0 @@ -283,27 +283,27 @@ OpStore %56 %188 %200 = OpCompositeExtract %3 %192 1 %201 = OpIMul %3 %199 %200 %193 = OpIAdd %3 %198 %201 -%202 = OpLoad %3 %56 +%202 = OpLoad %3 %57 %203 = OpIAdd %3 %202 %193 -OpStore %56 %203 -%204 = OpLoad %3 %56 -%205 = OpLoad %3 %56 +OpStore %57 %203 +%204 = OpLoad %3 %57 +%205 = OpLoad %3 %57 %206 = OpExtInst %3 %1 SMax %204 %205 -%207 = OpLoad %3 %56 +%207 = OpLoad %3 %57 %208 = OpIAdd %3 %207 %206 -OpStore %56 %208 -%209 = OpLoad %3 %56 -%210 = OpLoad %3 %56 +OpStore %57 %208 +%209 = OpLoad %3 %57 +%210 = OpLoad %3 %57 %211 = OpExtInst %3 %1 SMin %209 %210 -%212 = OpLoad %3 %56 +%212 = OpLoad %3 %57 %213 = OpIAdd %3 %212 %211 -OpStore %56 %213 -%214 = OpLoad %3 %56 +OpStore %57 %213 +%214 = OpLoad %3 %57 %215 = OpExtInst %3 %1 SSign %214 -%216 = OpLoad %3 %56 +%216 = OpLoad %3 %57 %217 = OpIAdd %3 %216 %215 -OpStore %56 %217 -%218 = OpLoad %3 %56 +OpStore %57 %217 +%218 = OpLoad %3 %57 OpReturnValue %218 OpFunctionEnd %221 = OpFunction %4 None %222 @@ -327,7 +327,7 @@ OpStore %231 %237 %240 = OpLoad %4 %231 %241 = OpIAdd %4 %240 %239 OpStore %231 %241 -%242 = OpAccessChain %67 %223 %43 +%242 = OpAccessChain %68 %223 %43 %243 = OpLoad %5 %242 %244 = OpLoad %4 %231 %245 = OpUConvert %5 %244 @@ -336,7 +336,7 @@ OpStore %231 %241 %248 = OpLoad %4 %231 %249 = OpIAdd %4 %248 %247 OpStore %231 %249 -%250 = OpAccessChain %76 %223 %77 +%250 = OpAccessChain %77 %223 %78 %251 = OpLoad %6 %250 %252 = OpLoad %4 %231 %253 = OpSConvert %6 %252 @@ -345,7 +345,7 @@ OpStore %231 %249 %256 = OpLoad %4 %231 %257 = OpIAdd %4 %256 %255 OpStore %231 %257 -%258 = OpAccessChain %86 %223 %16 +%258 = OpAccessChain %87 %223 %16 %259 = OpLoad %7 %258 %260 = OpLoad %4 %231 %261 = OpConvertUToF %7 %260 @@ -355,14 +355,14 @@ OpStore %231 %257 %267 = OpLoad %4 %231 %268 = OpIAdd %4 %267 %266 OpStore %231 %268 -%269 = OpAccessChain %106 %223 %107 +%269 = OpAccessChain %107 %223 %108 %270 = OpLoad %4 %269 %271 = OpCompositeConstruct %9 %270 %270 %270 %272 = OpCompositeExtract %4 %271 2 %273 = OpLoad %4 %231 %274 = OpIAdd %4 %273 %272 OpStore %231 %274 -%275 = OpAccessChain %98 %223 %99 +%275 = OpAccessChain %99 %223 %100 %276 = OpLoad %3 %275 %277 = OpBitcast %4 %276 %278 = OpLoad %4 %231 @@ -389,33 +389,33 @@ OpStore %231 %291 %296 = OpLoad %4 %231 %297 = OpIAdd %4 %296 %295 OpStore %231 %297 -%299 = OpAccessChain %106 %223 %107 +%299 = OpAccessChain %107 %223 %108 %300 = OpLoad %4 %299 -%301 = OpAccessChain %298 %224 %107 +%301 = OpAccessChain %298 %224 %108 %302 = OpLoad %4 %301 %303 = OpIAdd %4 %300 %302 -%304 = OpAccessChain %298 %226 %107 +%304 = OpAccessChain %298 %226 %108 OpStore %304 %303 -%306 = OpAccessChain %113 %223 %114 +%306 = OpAccessChain %114 %223 %115 %307 = OpLoad %8 %306 -%308 = OpAccessChain %305 %224 %114 +%308 = OpAccessChain %305 %224 %115 %309 = OpLoad %8 %308 %310 = OpIAdd %8 %307 %309 -%311 = OpAccessChain %305 %226 %114 +%311 = OpAccessChain %305 %226 %115 OpStore %311 %310 -%313 = OpAccessChain %121 %223 %122 +%313 = OpAccessChain %122 %223 %123 %314 = OpLoad %9 %313 -%315 = OpAccessChain %312 %224 %122 +%315 = OpAccessChain %312 %224 %123 %316 = OpLoad %9 %315 %317 = OpIAdd %9 %314 %316 -%318 = OpAccessChain %312 %226 %122 +%318 = OpAccessChain %312 %226 %123 OpStore %318 %317 -%320 = OpAccessChain %129 %223 %130 +%320 = OpAccessChain %130 %223 %131 %321 = OpLoad %10 %320 -%322 = OpAccessChain %319 %224 %130 +%322 = OpAccessChain %319 %224 %131 %323 = OpLoad %10 %322 %324 = OpIAdd %10 %321 %323 -%325 = OpAccessChain %319 %226 %130 +%325 = OpAccessChain %319 %226 %131 OpStore %325 %324 %327 = OpAccessChain %326 %225 %43 %328 = OpLoad %15 %327 diff --git a/naga/tests/out/spv/wgsl-operators.spvasm b/naga/tests/out/spv/wgsl-operators.spvasm index 687f8eddce..f3cfe57c8a 100644 --- a/naga/tests/out/spv/wgsl-operators.spvasm +++ b/naga/tests/out/spv/wgsl-operators.spvasm @@ -1,13 +1,13 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 527 +; Bound: 533 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %512 "main" %509 -OpExecutionMode %512 LocalSize 1 1 1 -OpDecorate %509 BuiltIn WorkgroupId +OpEntryPoint GLCompute %518 "main" %515 +OpExecutionMode %518 LocalSize 1 1 1 +OpDecorate %515 BuiltIn WorkgroupId %2 = OpTypeVoid %3 = OpTypeFloat 32 %4 = OpTypeVector %3 4 @@ -97,13 +97,15 @@ OpDecorate %509 BuiltIn WorkgroupId %250 = OpConstantNull %14 %251 = OpConstantComposite %10 %76 %76 %76 %252 = OpConstantNull %15 -%416 = OpConstantNull %16 -%418 = OpTypePointer Function %5 -%419 = OpConstantNull %5 -%421 = OpTypePointer Function %16 -%510 = OpTypePointer Input %11 -%509 = OpVariable %510 Input -%513 = OpConstantComposite %10 %17 %17 %17 +%254 = OpTypePointer Function %5 +%255 = OpConstantNull %5 +%257 = OpConstantNull %5 +%423 = OpConstantNull %16 +%425 = OpConstantNull %5 +%427 = OpTypePointer Function %16 +%516 = OpTypePointer Input %11 +%515 = OpVariable %516 Input +%519 = OpConstantComposite %10 %17 %17 %17 %26 = OpFunction %4 None %27 %25 = OpLabel OpBranch %34 @@ -306,305 +308,310 @@ OpReturnValue %240 OpFunctionEnd %242 = OpFunction %2 None %122 %241 = OpLabel -OpBranch %253 -%253 = OpLabel -%254 = OpFNegate %3 %17 -%255 = OpSNegate %160 %174 -%256 = OpFNegate %9 %97 -%257 = OpIAdd %5 %82 %23 -%258 = OpIAdd %12 %243 %157 -%259 = OpFAdd %3 %76 %17 -%260 = OpIAdd %160 %244 %174 -%261 = OpIAdd %11 %245 %184 -%262 = OpFAdd %4 %246 %247 -%263 = OpISub %5 %82 %23 -%264 = OpISub %12 %243 %157 -%265 = OpFSub %3 %76 %17 -%266 = OpISub %160 %244 %174 -%267 = OpISub %11 %245 %184 -%268 = OpFSub %4 %246 %247 -%269 = OpIMul %5 %82 %23 -%270 = OpIMul %12 %243 %157 -%271 = OpFMul %3 %76 %17 -%272 = OpIMul %160 %244 %174 -%273 = OpIMul %11 %245 %184 -%274 = OpFMul %4 %246 %247 -%275 = OpFunctionCall %5 %138 %82 %23 -%276 = OpFunctionCall %12 %150 %243 %157 -%277 = OpFDiv %3 %76 %17 -%278 = OpFunctionCall %160 %161 %244 %174 -%279 = OpFunctionCall %11 %177 %245 %184 -%280 = OpFDiv %4 %246 %247 -%281 = OpFunctionCall %5 %187 %82 %23 -%282 = OpFunctionCall %12 %198 %243 %157 -%283 = OpFRem %3 %76 %17 -%284 = OpFunctionCall %160 %205 %244 %174 -%285 = OpFunctionCall %11 %216 %245 %184 -%286 = OpFRem %4 %246 %247 -OpBranch %287 -%287 = OpLabel -%289 = OpIAdd %160 %244 %174 -%290 = OpIAdd %160 %244 %174 -%291 = OpIAdd %223 %248 %231 -%292 = OpIAdd %223 %248 %231 -%293 = OpFAdd %9 %77 %97 -%294 = OpFAdd %9 %77 %97 -%295 = OpISub %160 %244 %174 -%296 = OpISub %160 %244 %174 -%297 = OpISub %223 %248 %231 -%298 = OpISub %223 %248 %231 -%299 = OpFSub %9 %77 %97 -%300 = OpFSub %9 %77 %97 -%302 = OpCompositeConstruct %160 %23 %23 -%301 = OpIMul %160 %244 %302 -%304 = OpCompositeConstruct %160 %82 %82 -%303 = OpIMul %160 %174 %304 -%306 = OpCompositeConstruct %223 %157 %157 -%305 = OpIMul %223 %248 %306 -%308 = OpCompositeConstruct %223 %243 %243 -%307 = OpIMul %223 %231 %308 -%309 = OpVectorTimesScalar %9 %77 %17 -%310 = OpVectorTimesScalar %9 %97 %76 -%311 = OpFunctionCall %160 %161 %244 %174 -%312 = OpFunctionCall %160 %161 %244 %174 -%313 = OpFunctionCall %223 %224 %248 %231 -%314 = OpFunctionCall %223 %224 %248 %231 -%315 = OpFDiv %9 %77 %97 -%316 = OpFDiv %9 %77 %97 -%317 = OpFunctionCall %160 %205 %244 %174 -%318 = OpFunctionCall %160 %205 %244 %174 -%319 = OpFunctionCall %223 %234 %248 %231 -%320 = OpFunctionCall %223 %234 %248 %231 -%321 = OpFRem %9 %77 %97 -%322 = OpFRem %9 %77 %97 -OpBranch %288 -%288 = OpLabel -%324 = OpCompositeExtract %10 %249 0 -%325 = OpCompositeExtract %10 %249 0 -%326 = OpFAdd %10 %324 %325 -%327 = OpCompositeExtract %10 %249 1 -%328 = OpCompositeExtract %10 %249 1 -%329 = OpFAdd %10 %327 %328 -%330 = OpCompositeExtract %10 %249 2 -%331 = OpCompositeExtract %10 %249 2 -%332 = OpFAdd %10 %330 %331 -%323 = OpCompositeConstruct %13 %326 %329 %332 -%334 = OpCompositeExtract %10 %249 0 -%335 = OpCompositeExtract %10 %249 0 -%336 = OpFSub %10 %334 %335 -%337 = OpCompositeExtract %10 %249 1 -%338 = OpCompositeExtract %10 %249 1 -%339 = OpFSub %10 %337 %338 -%340 = OpCompositeExtract %10 %249 2 -%341 = OpCompositeExtract %10 %249 2 -%342 = OpFSub %10 %340 %341 -%333 = OpCompositeConstruct %13 %336 %339 %342 -%343 = OpMatrixTimesScalar %13 %249 %17 -%344 = OpMatrixTimesScalar %13 %249 %76 -%345 = OpMatrixTimesVector %10 %250 %247 -%346 = OpVectorTimesMatrix %4 %251 %250 -%347 = OpMatrixTimesMatrix %13 %250 %252 +%253 = OpVariable %254 Function %255 +%256 = OpVariable %254 Function %257 +OpBranch %258 +%258 = OpLabel +%259 = OpFNegate %3 %17 +%260 = OpSNegate %160 %174 +%261 = OpFNegate %9 %97 +%262 = OpIAdd %5 %82 %23 +%263 = OpIAdd %12 %243 %157 +%264 = OpFAdd %3 %76 %17 +%265 = OpIAdd %160 %244 %174 +%266 = OpIAdd %11 %245 %184 +%267 = OpFAdd %4 %246 %247 +%268 = OpISub %5 %82 %23 +%269 = OpISub %12 %243 %157 +%270 = OpFSub %3 %76 %17 +%271 = OpISub %160 %244 %174 +%272 = OpISub %11 %245 %184 +%273 = OpFSub %4 %246 %247 +%274 = OpIMul %5 %82 %23 +%275 = OpIMul %12 %243 %157 +%276 = OpFMul %3 %76 %17 +%277 = OpIMul %160 %244 %174 +%278 = OpIMul %11 %245 %184 +%279 = OpFMul %4 %246 %247 +%280 = OpFunctionCall %5 %138 %82 %23 +%281 = OpFunctionCall %12 %150 %243 %157 +%282 = OpFDiv %3 %76 %17 +%283 = OpFunctionCall %160 %161 %244 %174 +%284 = OpFunctionCall %11 %177 %245 %184 +%285 = OpFDiv %4 %246 %247 +%286 = OpFunctionCall %5 %187 %82 %23 +%287 = OpFunctionCall %12 %198 %243 %157 +%288 = OpFRem %3 %76 %17 +%289 = OpFunctionCall %160 %205 %244 %174 +%290 = OpFunctionCall %11 %216 %245 %184 +%291 = OpFRem %4 %246 %247 +OpBranch %292 +%292 = OpLabel +%294 = OpIAdd %160 %244 %174 +%295 = OpIAdd %160 %244 %174 +%296 = OpIAdd %223 %248 %231 +%297 = OpIAdd %223 %248 %231 +%298 = OpFAdd %9 %77 %97 +%299 = OpFAdd %9 %77 %97 +%300 = OpISub %160 %244 %174 +%301 = OpISub %160 %244 %174 +%302 = OpISub %223 %248 %231 +%303 = OpISub %223 %248 %231 +%304 = OpFSub %9 %77 %97 +%305 = OpFSub %9 %77 %97 +%307 = OpCompositeConstruct %160 %23 %23 +%306 = OpIMul %160 %244 %307 +%309 = OpCompositeConstruct %160 %82 %82 +%308 = OpIMul %160 %174 %309 +%311 = OpCompositeConstruct %223 %157 %157 +%310 = OpIMul %223 %248 %311 +%313 = OpCompositeConstruct %223 %243 %243 +%312 = OpIMul %223 %231 %313 +%314 = OpVectorTimesScalar %9 %77 %17 +%315 = OpVectorTimesScalar %9 %97 %76 +%316 = OpFunctionCall %160 %161 %244 %174 +%317 = OpFunctionCall %160 %161 %244 %174 +%318 = OpFunctionCall %223 %224 %248 %231 +%319 = OpFunctionCall %223 %224 %248 %231 +%320 = OpFDiv %9 %77 %97 +%321 = OpFDiv %9 %77 %97 +%322 = OpFunctionCall %160 %205 %244 %174 +%323 = OpFunctionCall %160 %205 %244 %174 +%324 = OpFunctionCall %223 %234 %248 %231 +%325 = OpFunctionCall %223 %234 %248 %231 +%326 = OpFRem %9 %77 %97 +%327 = OpFRem %9 %77 %97 +OpBranch %293 +%293 = OpLabel +%329 = OpCompositeExtract %10 %249 0 +%330 = OpCompositeExtract %10 %249 0 +%331 = OpFAdd %10 %329 %330 +%332 = OpCompositeExtract %10 %249 1 +%333 = OpCompositeExtract %10 %249 1 +%334 = OpFAdd %10 %332 %333 +%335 = OpCompositeExtract %10 %249 2 +%336 = OpCompositeExtract %10 %249 2 +%337 = OpFAdd %10 %335 %336 +%328 = OpCompositeConstruct %13 %331 %334 %337 +%339 = OpCompositeExtract %10 %249 0 +%340 = OpCompositeExtract %10 %249 0 +%341 = OpFSub %10 %339 %340 +%342 = OpCompositeExtract %10 %249 1 +%343 = OpCompositeExtract %10 %249 1 +%344 = OpFSub %10 %342 %343 +%345 = OpCompositeExtract %10 %249 2 +%346 = OpCompositeExtract %10 %249 2 +%347 = OpFSub %10 %345 %346 +%338 = OpCompositeConstruct %13 %341 %344 %347 +%348 = OpMatrixTimesScalar %13 %249 %17 +%349 = OpMatrixTimesScalar %13 %249 %76 +%350 = OpMatrixTimesVector %10 %250 %247 +%351 = OpVectorTimesMatrix %4 %251 %250 +%352 = OpMatrixTimesMatrix %13 %250 %252 +%353 = OpLoad %5 %253 +%354 = OpIAdd %5 %353 %60 +OpStore %256 %354 OpReturn OpFunctionEnd -%349 = OpFunction %2 None %122 -%348 = OpLabel -OpBranch %350 -%350 = OpLabel -%351 = OpNot %5 %23 -%352 = OpNot %12 %157 -%353 = OpNot %160 %174 -%354 = OpNot %11 %184 -%355 = OpBitwiseOr %5 %82 %23 -%356 = OpBitwiseOr %12 %243 %157 -%357 = OpBitwiseOr %160 %244 %174 -%358 = OpBitwiseOr %11 %245 %184 -%359 = OpBitwiseAnd %5 %82 %23 -%360 = OpBitwiseAnd %12 %243 %157 -%361 = OpBitwiseAnd %160 %244 %174 -%362 = OpBitwiseAnd %11 %245 %184 -%363 = OpBitwiseXor %5 %82 %23 -%364 = OpBitwiseXor %12 %243 %157 -%365 = OpBitwiseXor %160 %244 %174 -%366 = OpBitwiseXor %11 %245 %184 -%367 = OpShiftLeftLogical %5 %82 %157 -%368 = OpShiftLeftLogical %12 %243 %157 -%369 = OpShiftLeftLogical %160 %244 %231 -%370 = OpShiftLeftLogical %11 %245 %184 -%371 = OpShiftRightArithmetic %5 %82 %157 -%372 = OpShiftRightLogical %12 %243 %157 -%373 = OpShiftRightArithmetic %160 %244 %231 -%374 = OpShiftRightLogical %11 %245 %184 +%356 = OpFunction %2 None %122 +%355 = OpLabel +OpBranch %357 +%357 = OpLabel +%358 = OpNot %5 %23 +%359 = OpNot %12 %157 +%360 = OpNot %160 %174 +%361 = OpNot %11 %184 +%362 = OpBitwiseOr %5 %82 %23 +%363 = OpBitwiseOr %12 %243 %157 +%364 = OpBitwiseOr %160 %244 %174 +%365 = OpBitwiseOr %11 %245 %184 +%366 = OpBitwiseAnd %5 %82 %23 +%367 = OpBitwiseAnd %12 %243 %157 +%368 = OpBitwiseAnd %160 %244 %174 +%369 = OpBitwiseAnd %11 %245 %184 +%370 = OpBitwiseXor %5 %82 %23 +%371 = OpBitwiseXor %12 %243 %157 +%372 = OpBitwiseXor %160 %244 %174 +%373 = OpBitwiseXor %11 %245 %184 +%374 = OpShiftLeftLogical %5 %82 %157 +%375 = OpShiftLeftLogical %12 %243 %157 +%376 = OpShiftLeftLogical %160 %244 %231 +%377 = OpShiftLeftLogical %11 %245 %184 +%378 = OpShiftRightArithmetic %5 %82 %157 +%379 = OpShiftRightLogical %12 %243 %157 +%380 = OpShiftRightArithmetic %160 %244 %231 +%381 = OpShiftRightLogical %11 %245 %184 OpReturn OpFunctionEnd -%376 = OpFunction %2 None %122 -%375 = OpLabel -OpBranch %377 -%377 = OpLabel -%378 = OpIEqual %8 %82 %23 -%379 = OpIEqual %8 %243 %157 -%380 = OpFOrdEqual %8 %76 %17 -%381 = OpIEqual %123 %244 %174 -%382 = OpIEqual %115 %245 %184 -%383 = OpFOrdEqual %7 %246 %247 -%384 = OpINotEqual %8 %82 %23 -%385 = OpINotEqual %8 %243 %157 -%386 = OpFOrdNotEqual %8 %76 %17 -%387 = OpINotEqual %123 %244 %174 -%388 = OpINotEqual %115 %245 %184 -%389 = OpFOrdNotEqual %7 %246 %247 -%390 = OpSLessThan %8 %82 %23 -%391 = OpULessThan %8 %243 %157 -%392 = OpFOrdLessThan %8 %76 %17 -%393 = OpSLessThan %123 %244 %174 -%394 = OpULessThan %115 %245 %184 -%395 = OpFOrdLessThan %7 %246 %247 -%396 = OpSLessThanEqual %8 %82 %23 -%397 = OpULessThanEqual %8 %243 %157 -%398 = OpFOrdLessThanEqual %8 %76 %17 -%399 = OpSLessThanEqual %123 %244 %174 -%400 = OpULessThanEqual %115 %245 %184 -%401 = OpFOrdLessThanEqual %7 %246 %247 -%402 = OpSGreaterThan %8 %82 %23 -%403 = OpUGreaterThan %8 %243 %157 -%404 = OpFOrdGreaterThan %8 %76 %17 -%405 = OpSGreaterThan %123 %244 %174 -%406 = OpUGreaterThan %115 %245 %184 -%407 = OpFOrdGreaterThan %7 %246 %247 -%408 = OpSGreaterThanEqual %8 %82 %23 -%409 = OpUGreaterThanEqual %8 %243 %157 -%410 = OpFOrdGreaterThanEqual %8 %76 %17 -%411 = OpSGreaterThanEqual %123 %244 %174 -%412 = OpUGreaterThanEqual %115 %245 %184 -%413 = OpFOrdGreaterThanEqual %7 %246 %247 +%383 = OpFunction %2 None %122 +%382 = OpLabel +OpBranch %384 +%384 = OpLabel +%385 = OpIEqual %8 %82 %23 +%386 = OpIEqual %8 %243 %157 +%387 = OpFOrdEqual %8 %76 %17 +%388 = OpIEqual %123 %244 %174 +%389 = OpIEqual %115 %245 %184 +%390 = OpFOrdEqual %7 %246 %247 +%391 = OpINotEqual %8 %82 %23 +%392 = OpINotEqual %8 %243 %157 +%393 = OpFOrdNotEqual %8 %76 %17 +%394 = OpINotEqual %123 %244 %174 +%395 = OpINotEqual %115 %245 %184 +%396 = OpFOrdNotEqual %7 %246 %247 +%397 = OpSLessThan %8 %82 %23 +%398 = OpULessThan %8 %243 %157 +%399 = OpFOrdLessThan %8 %76 %17 +%400 = OpSLessThan %123 %244 %174 +%401 = OpULessThan %115 %245 %184 +%402 = OpFOrdLessThan %7 %246 %247 +%403 = OpSLessThanEqual %8 %82 %23 +%404 = OpULessThanEqual %8 %243 %157 +%405 = OpFOrdLessThanEqual %8 %76 %17 +%406 = OpSLessThanEqual %123 %244 %174 +%407 = OpULessThanEqual %115 %245 %184 +%408 = OpFOrdLessThanEqual %7 %246 %247 +%409 = OpSGreaterThan %8 %82 %23 +%410 = OpUGreaterThan %8 %243 %157 +%411 = OpFOrdGreaterThan %8 %76 %17 +%412 = OpSGreaterThan %123 %244 %174 +%413 = OpUGreaterThan %115 %245 %184 +%414 = OpFOrdGreaterThan %7 %246 %247 +%415 = OpSGreaterThanEqual %8 %82 %23 +%416 = OpUGreaterThanEqual %8 %243 %157 +%417 = OpFOrdGreaterThanEqual %8 %76 %17 +%418 = OpSGreaterThanEqual %123 %244 %174 +%419 = OpUGreaterThanEqual %115 %245 %184 +%420 = OpFOrdGreaterThanEqual %7 %246 %247 OpReturn OpFunctionEnd -%415 = OpFunction %2 None %122 -%414 = OpLabel -%417 = OpVariable %418 Function %419 -%420 = OpVariable %421 Function %416 -OpBranch %422 -%422 = OpLabel -OpStore %417 %23 -%423 = OpLoad %5 %417 -%424 = OpIAdd %5 %423 %23 -OpStore %417 %424 -%425 = OpLoad %5 %417 -%426 = OpISub %5 %425 %23 -OpStore %417 %426 -%427 = OpLoad %5 %417 -%428 = OpLoad %5 %417 -%429 = OpIMul %5 %428 %427 -OpStore %417 %429 -%430 = OpLoad %5 %417 -%431 = OpLoad %5 %417 -%432 = OpFunctionCall %5 %138 %431 %430 -OpStore %417 %432 -%433 = OpLoad %5 %417 -%434 = OpFunctionCall %5 %187 %433 %23 -OpStore %417 %434 -%435 = OpLoad %5 %417 -%436 = OpBitwiseAnd %5 %435 %29 -OpStore %417 %436 -%437 = OpLoad %5 %417 -%438 = OpBitwiseOr %5 %437 %29 -OpStore %417 %438 -%439 = OpLoad %5 %417 -%440 = OpBitwiseXor %5 %439 %29 -OpStore %417 %440 -%441 = OpLoad %5 %417 -%442 = OpShiftLeftLogical %5 %441 %243 -OpStore %417 %442 -%443 = OpLoad %5 %417 -%444 = OpShiftRightArithmetic %5 %443 %157 -OpStore %417 %444 -%445 = OpLoad %5 %417 -%446 = OpIAdd %5 %445 %23 -OpStore %417 %446 -%447 = OpLoad %5 %417 -%448 = OpISub %5 %447 %23 -OpStore %417 %448 -%449 = OpAccessChain %418 %420 %157 -%450 = OpLoad %5 %449 -%451 = OpIAdd %5 %450 %23 -%452 = OpAccessChain %418 %420 %157 -OpStore %452 %451 -%453 = OpAccessChain %418 %420 %157 -%454 = OpLoad %5 %453 -%455 = OpISub %5 %454 %23 -%456 = OpAccessChain %418 %420 %157 -OpStore %456 %455 +%422 = OpFunction %2 None %122 +%421 = OpLabel +%424 = OpVariable %254 Function %425 +%426 = OpVariable %427 Function %423 +OpBranch %428 +%428 = OpLabel +OpStore %424 %23 +%429 = OpLoad %5 %424 +%430 = OpIAdd %5 %429 %23 +OpStore %424 %430 +%431 = OpLoad %5 %424 +%432 = OpISub %5 %431 %23 +OpStore %424 %432 +%433 = OpLoad %5 %424 +%434 = OpLoad %5 %424 +%435 = OpIMul %5 %434 %433 +OpStore %424 %435 +%436 = OpLoad %5 %424 +%437 = OpLoad %5 %424 +%438 = OpFunctionCall %5 %138 %437 %436 +OpStore %424 %438 +%439 = OpLoad %5 %424 +%440 = OpFunctionCall %5 %187 %439 %23 +OpStore %424 %440 +%441 = OpLoad %5 %424 +%442 = OpBitwiseAnd %5 %441 %29 +OpStore %424 %442 +%443 = OpLoad %5 %424 +%444 = OpBitwiseOr %5 %443 %29 +OpStore %424 %444 +%445 = OpLoad %5 %424 +%446 = OpBitwiseXor %5 %445 %29 +OpStore %424 %446 +%447 = OpLoad %5 %424 +%448 = OpShiftLeftLogical %5 %447 %243 +OpStore %424 %448 +%449 = OpLoad %5 %424 +%450 = OpShiftRightArithmetic %5 %449 %157 +OpStore %424 %450 +%451 = OpLoad %5 %424 +%452 = OpIAdd %5 %451 %23 +OpStore %424 %452 +%453 = OpLoad %5 %424 +%454 = OpISub %5 %453 %23 +OpStore %424 %454 +%455 = OpAccessChain %254 %426 %157 +%456 = OpLoad %5 %455 +%457 = OpIAdd %5 %456 %23 +%458 = OpAccessChain %254 %426 %157 +OpStore %458 %457 +%459 = OpAccessChain %254 %426 %157 +%460 = OpLoad %5 %459 +%461 = OpISub %5 %460 %23 +%462 = OpAccessChain %254 %426 %157 +OpStore %462 %461 OpReturn OpFunctionEnd -%458 = OpFunction %2 None %122 -%457 = OpLabel -OpBranch %459 -%459 = OpLabel -%460 = OpSNegate %5 %23 -%461 = OpSNegate %5 %23 -%462 = OpSNegate %5 %461 -%463 = OpSNegate %5 %23 -%464 = OpSNegate %5 %463 -%465 = OpSNegate %5 %23 -%466 = OpSNegate %5 %465 +%464 = OpFunction %2 None %122 +%463 = OpLabel +OpBranch %465 +%465 = OpLabel +%466 = OpSNegate %5 %23 %467 = OpSNegate %5 %23 %468 = OpSNegate %5 %467 -%469 = OpSNegate %5 %468 -%470 = OpSNegate %5 %23 -%471 = OpSNegate %5 %470 +%469 = OpSNegate %5 %23 +%470 = OpSNegate %5 %469 +%471 = OpSNegate %5 %23 %472 = OpSNegate %5 %471 -%473 = OpSNegate %5 %472 -%474 = OpSNegate %5 %23 +%473 = OpSNegate %5 %23 +%474 = OpSNegate %5 %473 %475 = OpSNegate %5 %474 -%476 = OpSNegate %5 %475 +%476 = OpSNegate %5 %23 %477 = OpSNegate %5 %476 %478 = OpSNegate %5 %477 -%479 = OpSNegate %5 %23 -%480 = OpSNegate %5 %479 +%479 = OpSNegate %5 %478 +%480 = OpSNegate %5 %23 %481 = OpSNegate %5 %480 %482 = OpSNegate %5 %481 %483 = OpSNegate %5 %482 -%484 = OpFNegate %3 %17 -%485 = OpFNegate %3 %17 -%486 = OpFNegate %3 %485 -%487 = OpFNegate %3 %17 -%488 = OpFNegate %3 %487 -%489 = OpFNegate %3 %17 -%490 = OpFNegate %3 %489 +%484 = OpSNegate %5 %483 +%485 = OpSNegate %5 %23 +%486 = OpSNegate %5 %485 +%487 = OpSNegate %5 %486 +%488 = OpSNegate %5 %487 +%489 = OpSNegate %5 %488 +%490 = OpFNegate %3 %17 %491 = OpFNegate %3 %17 %492 = OpFNegate %3 %491 -%493 = OpFNegate %3 %492 -%494 = OpFNegate %3 %17 -%495 = OpFNegate %3 %494 +%493 = OpFNegate %3 %17 +%494 = OpFNegate %3 %493 +%495 = OpFNegate %3 %17 %496 = OpFNegate %3 %495 -%497 = OpFNegate %3 %496 -%498 = OpFNegate %3 %17 +%497 = OpFNegate %3 %17 +%498 = OpFNegate %3 %497 %499 = OpFNegate %3 %498 -%500 = OpFNegate %3 %499 +%500 = OpFNegate %3 %17 %501 = OpFNegate %3 %500 %502 = OpFNegate %3 %501 -%503 = OpFNegate %3 %17 -%504 = OpFNegate %3 %503 +%503 = OpFNegate %3 %502 +%504 = OpFNegate %3 %17 %505 = OpFNegate %3 %504 %506 = OpFNegate %3 %505 %507 = OpFNegate %3 %506 +%508 = OpFNegate %3 %507 +%509 = OpFNegate %3 %17 +%510 = OpFNegate %3 %509 +%511 = OpFNegate %3 %510 +%512 = OpFNegate %3 %511 +%513 = OpFNegate %3 %512 OpReturn OpFunctionEnd -%512 = OpFunction %2 None %122 -%508 = OpLabel -%511 = OpLoad %11 %509 -OpBranch %514 +%518 = OpFunction %2 None %122 %514 = OpLabel -%515 = OpFunctionCall %4 %26 -%516 = OpCompositeExtract %12 %511 0 -%517 = OpConvertUToF %3 %516 -%518 = OpCompositeExtract %12 %511 1 -%519 = OpBitcast %5 %518 -%520 = OpFunctionCall %4 %74 %517 %519 -%521 = OpFunctionCall %10 %112 %513 -%522 = OpFunctionCall %2 %121 -%523 = OpFunctionCall %2 %242 -%524 = OpFunctionCall %2 %349 -%525 = OpFunctionCall %2 %376 -%526 = OpFunctionCall %2 %415 +%517 = OpLoad %11 %515 +OpBranch %520 +%520 = OpLabel +%521 = OpFunctionCall %4 %26 +%522 = OpCompositeExtract %12 %517 0 +%523 = OpConvertUToF %3 %522 +%524 = OpCompositeExtract %12 %517 1 +%525 = OpBitcast %5 %524 +%526 = OpFunctionCall %4 %74 %523 %525 +%527 = OpFunctionCall %10 %112 %519 +%528 = OpFunctionCall %2 %121 +%529 = OpFunctionCall %2 %242 +%530 = OpFunctionCall %2 %356 +%531 = OpFunctionCall %2 %383 +%532 = OpFunctionCall %2 %422 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/wgsl-conversion-float-to-int.wgsl b/naga/tests/out/wgsl/wgsl-conversion-float-to-int.wgsl index a185b81dea..8435f01bf5 100644 --- a/naga/tests/out/wgsl/wgsl-conversion-float-to-int.wgsl +++ b/naga/tests/out/wgsl/wgsl-conversion-float-to-int.wgsl @@ -20,11 +20,11 @@ fn test_const_eval() { var max_f32_to_i32_: i32 = 2147483520i; var min_f32_to_u32_: u32 = 0u; var max_f32_to_u32_: u32 = 4294967040u; - var min_f32_to_i64_: i64 = -9223372036854775807li - 1li; + var min_f32_to_i64_: i64 = i64(-9223372036854775807 - 1); var max_f32_to_i64_: i64 = 9223371487098961920li; var min_f32_to_u64_: u64 = 0lu; var max_f32_to_u64_: u64 = 18446742974197923840lu; - var min_f64_to_i64_: i64 = -9223372036854775807li - 1li; + var min_f64_to_i64_: i64 = i64(-9223372036854775807 - 1); var max_f64_to_i64_: i64 = 9223372036854774784li; var min_f64_to_u64_: u64 = 0lu; var max_f64_to_u64_: u64 = 18446744073709549568lu; @@ -32,7 +32,7 @@ fn test_const_eval() { var max_abstract_float_to_i32_: i32 = 2147483647i; var min_abstract_float_to_u32_: u32 = 0u; var max_abstract_float_to_u32_: u32 = 4294967295u; - var min_abstract_float_to_i64_: i64 = -9223372036854775807li - 1li; + var min_abstract_float_to_i64_: i64 = i64(-9223372036854775807 - 1); var max_abstract_float_to_i64_: i64 = 9223372036854774784li; var min_abstract_float_to_u64_: u64 = 0lu; var max_abstract_float_to_u64_: u64 = 18446744073709549568lu; diff --git a/naga/tests/out/wgsl/wgsl-int64.wgsl b/naga/tests/out/wgsl/wgsl-int64.wgsl index c339d8d4c3..b554457700 100644 --- a/naga/tests/out/wgsl/wgsl-int64.wgsl +++ b/naga/tests/out/wgsl/wgsl-int64.wgsl @@ -67,47 +67,47 @@ fn int64_function(x: i64) -> i64 { let _e71 = input_uniform.val_u64_4_; let _e74 = val; val = (_e74 + bitcast>(_e71).w); - let _e79 = val; - val = (_e79 + (-9223372036854775807li - 1li)); - let _e85 = input_uniform.val_i64_; - let _e88 = input_storage.val_i64_; - output.val_i64_ = (_e85 + _e88); - let _e94 = input_uniform.val_i64_2_; - let _e97 = input_storage.val_i64_2_; - output.val_i64_2_ = (_e94 + _e97); - let _e103 = input_uniform.val_i64_3_; - let _e106 = input_storage.val_i64_3_; - output.val_i64_3_ = (_e103 + _e106); - let _e112 = input_uniform.val_i64_4_; - let _e115 = input_storage.val_i64_4_; - output.val_i64_4_ = (_e112 + _e115); - let _e121 = input_arrays.val_i64_array_2_; - output_arrays.val_i64_array_2_ = _e121; + let _e77 = val; + val = (_e77 + i64(-9223372036854775807 - 1)); + let _e83 = input_uniform.val_i64_; + let _e86 = input_storage.val_i64_; + output.val_i64_ = (_e83 + _e86); + let _e92 = input_uniform.val_i64_2_; + let _e95 = input_storage.val_i64_2_; + output.val_i64_2_ = (_e92 + _e95); + let _e101 = input_uniform.val_i64_3_; + let _e104 = input_storage.val_i64_3_; + output.val_i64_3_ = (_e101 + _e104); + let _e110 = input_uniform.val_i64_4_; + let _e113 = input_storage.val_i64_4_; + output.val_i64_4_ = (_e110 + _e113); + let _e119 = input_arrays.val_i64_array_2_; + output_arrays.val_i64_array_2_ = _e119; + let _e120 = val; let _e122 = val; + val = (_e122 + abs(_e120)); let _e124 = val; - val = (_e124 + abs(_e122)); + let _e125 = val; let _e126 = val; - let _e127 = val; let _e128 = val; + val = (_e128 + clamp(_e124, _e125, _e126)); let _e130 = val; - val = (_e130 + clamp(_e126, _e127, _e128)); let _e132 = val; - let _e134 = val; + let _e135 = val; + val = (_e135 + dot(vec2(_e130), vec2(_e132))); let _e137 = val; - val = (_e137 + dot(vec2(_e132), vec2(_e134))); - let _e139 = val; + let _e138 = val; let _e140 = val; + val = (_e140 + max(_e137, _e138)); let _e142 = val; - val = (_e142 + max(_e139, _e140)); - let _e144 = val; + let _e143 = val; let _e145 = val; + val = (_e145 + min(_e142, _e143)); let _e147 = val; - val = (_e147 + min(_e144, _e145)); let _e149 = val; + val = (_e149 + sign(_e147)); let _e151 = val; - val = (_e151 + sign(_e149)); - let _e153 = val; - return _e153; + return _e151; } fn uint64_function(x_1: u64) -> u64 { diff --git a/naga/tests/out/wgsl/wgsl-operators.wgsl b/naga/tests/out/wgsl/wgsl-operators.wgsl index eaae0a429a..b009650e32 100644 --- a/naga/tests/out/wgsl/wgsl-operators.wgsl +++ b/naga/tests/out/wgsl/wgsl-operators.wgsl @@ -52,6 +52,9 @@ fn logical() { } fn arithmetic() { + var prevent_const_eval: i32; + var wgpu_7437_: i32; + let neg0_1 = -(1f); let neg1_1 = -(vec2(1i)); let neg2_ = -(vec2(1f)); @@ -124,6 +127,8 @@ fn arithmetic() { let mul_vector0_ = (mat4x3() * vec4(1f)); let mul_vector1_ = (vec3(2f) * mat4x3()); let mul = (mat4x3() * mat3x4()); + let _e175 = prevent_const_eval; + wgpu_7437_ = (_e175 + i32(-2147483648)); return; } diff --git a/naga/xtask/src/validate.rs b/naga/xtask/src/validate.rs index 8eb72b3ea8..d1d75783ea 100644 --- a/naga/xtask/src/validate.rs +++ b/naga/xtask/src/validate.rs @@ -280,9 +280,11 @@ fn validate_metal(path: &Path, xcrun: &str) -> anyhow::Result<()> { } else { format!("-std={language}") }; + let warnings_as_errors = ["-Werror=constant-conversion"]; EasyCommand::new(xcrun, |cmd| { cmd.args(["-sdk", "macosx", "metal", "-mmacosx-version-min=10.11"]) .arg(std_arg) + .args(warnings_as_errors) .args(["-x", "metal", &*path.to_string_lossy(), "-o", "/dev/null"]) }) .success()