From c3121eba0d6a6a4c23263dd33a19ba9b65524d76 Mon Sep 17 00:00:00 2001 From: Jamie Nicol Date: Thu, 27 Mar 2025 10:04:09 +0000 Subject: [PATCH 1/2] [naga wgsl-out] Improve representation of minimum int64 literal In #7424 we fixed a bug where the representation of the minimum int64 literal generated by naga was invalid WGSL. It changed us from expressing it as `-9223372036854775808` which was invalid, to `-9223372036854775807li - 1li`. This is valid WGSL. However, as the values are concrete i64 types if the shader is parsed again by naga the expression does not get const evaluated away, leading to suboptimal code generated by the backends. This patch makes us perform the subtraction using abstract integers before casting to i64, solving this problem. Additionally the input WGSL test is updated to use the same construct. --- naga/src/back/wgsl/writer.rs | 10 +- naga/tests/in/wgsl/int64.wgsl | 2 +- naga/tests/out/hlsl/wgsl-int64.hlsl | 56 ++-- naga/tests/out/msl/wgsl-int64.msl | 58 ++-- naga/tests/out/spv/wgsl-int64.spvasm | 284 +++++++++--------- .../wgsl/wgsl-conversion-float-to-int.wgsl | 6 +- naga/tests/out/wgsl/wgsl-int64.wgsl | 56 ++-- 7 files changed, 236 insertions(+), 236 deletions(-) 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/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/msl/wgsl-int64.msl b/naga/tests/out/msl/wgsl-int64.msl index e026cac919..ed5d3da2e4 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(-9223372036854775808L)); + 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/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/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 { From 903f812a7891f99d76bb2b1aa7adda24ae8a6f6e Mon Sep 17 00:00:00 2001 From: Christian Legnitto Date: Wed, 26 Mar 2025 13:15:18 -0400 Subject: [PATCH 2/2] [naga msl-out] invalid code emitted for minimum i32/i64 value literal. The literal `-2147483648` is parsed by Metal as negation of positive 2147483648. As 2147483648 is too large for a int, the expression is silently promoted to a long. Sometimes this does not matter as it will often be implicitly converted back to an int after the negation. However, if the expression is used in a bitcast then we hit a compiler error due to mismatched bitwidths. Similarily for `-9223372036854775808`, as 9223372036854775808 is too large for a long, metal emits a `-Wconstant-conversion` warning and changes the value to -9223372036854775808. This would then be negated again, possibly causing undefined behaviour. In both cases we can avoid the issue by expressing the literals as the second most negative value expressible by the type, minus one. eg `-2147483647 - 1` and `-9223372036854775807L - 1L`. We have added a test which uses the most negative i32 literal in an addition. Because we bitcast addition operands to unsigned in metal, this would cause a validation error without this fix. For the i64 case existing tests already make use of the minimum literal value. Passing the flag `-Werror=constant-conversion` to Metal during validation will therefore catch this issue. --- naga/src/back/msl/writer.rs | 22 +- naga/tests/in/wgsl/operators.wgsl | 8 + .../out/glsl/wgsl-operators.main.Compute.glsl | 4 + naga/tests/out/hlsl/wgsl-operators.hlsl | 5 + .../out/msl/wgsl-abstract-types-operators.msl | 4 +- .../wgsl-conversion-float-to-int-no-f64.msl | 8 +- naga/tests/out/msl/wgsl-int64.msl | 2 +- naga/tests/out/msl/wgsl-operators.msl | 4 + naga/tests/out/spv/wgsl-operators.spvasm | 571 +++++++++--------- naga/tests/out/wgsl/wgsl-operators.wgsl | 5 + naga/xtask/src/validate.rs | 2 + 11 files changed, 344 insertions(+), 291 deletions(-) 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/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-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 ed5d3da2e4..9b3d50d4b4 100644 --- a/naga/tests/out/msl/wgsl-int64.msl +++ b/naga/tests/out/msl/wgsl-int64.msl @@ -84,7 +84,7 @@ long int64_function( long _e74 = val; val = as_type(as_type(_e74) + as_type(as_type(_e71).w)); long _e77 = val; - val = as_type(as_type(_e77) + as_type(-9223372036854775808L)); + 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)); 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-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-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()