[tint] Abs polyfill for all backends Now that the signed arithmetic polyfill has landed [0] the workaround for abs (as max(x,-x)) will now function correctly. The issue with abs(i32) was only proven to be observed on intel and nvidia but have have chosen to proactively extend this fix to all backends and devices. This decisions was made based on the polyfill being very low overhead and assuming a commonality between all backend compilers (LLVM). [0] https://dawn-review.googlesource.com/c/dawn/+/249135 Bug: 426999765 Change-Id: I2597cbcaacf8e68f93eaa7fb14d67d16cadcf970 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/247896 Commit-Queue: Peter McNeeley <petermcneeley@google.com> Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp index 7eb87f3..53ca6ac 100644 --- a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp +++ b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
@@ -26,7 +26,7 @@ // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include <cstdint> -#include <numeric> +#include <limits> #include <string> #include <vector> @@ -174,6 +174,61 @@ EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size()); } +TEST_P(PolyfillBuiltinSimpleTests, AbsWithBranch) { + // Some backend compilers assume that return value of 'abs' is always positive. This is + // not true for one specific value of i32 (0x8000'0000). + // Operations on the value returned can prove that the compiler is assuming this value is + // positive. See crbug.com/426999765 + std::string kShaderCode = R"( + struct Data { values: array<i32> }; + @group(0) @binding(0) var<storage, read> input_data: Data; + @group(0) @binding(1) var<storage, read_write> output_data: Data; + + @compute @workgroup_size(4) + fn main(@builtin(global_invocation_id) global_id: vec3<u32>) { + var result = input_data.values[global_id.x]; + // Translates to SAbs ext instruction (spriv) + result = abs(result); + // Will translate to SMax ext instruction (spriv) and reproduce the bug. + // result = max(result, 3488); + // Another way to test the compiler is to use a conditional. + // The compiler incorrectly assumes 'result' is positive. + if(result < 0){ + // This branch will (correctly) be taken iff original value was min i32. + result = 1543; + } + // try 2 + output_data.values[global_id.x] = result; + } + )"; + + wgpu::ComputePipeline pipeline = CreateComputePipeline(kShaderCode); + uint32_t kDefaultVal = 0; + std::vector<uint32_t> init_input = {uint32_t(std::numeric_limits<int32_t>::lowest()), + uint32_t(-15), 17, 123}; + + wgpu::Buffer input = CreateBuffer(init_input); + wgpu::Buffer output = CreateBuffer(4, kDefaultVal); + wgpu::BindGroup bindGroup = + utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, input}, {1, output}}); + + wgpu::CommandBuffer commands; + { + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(pipeline); + pass.SetBindGroup(0, bindGroup); + pass.DispatchWorkgroups(64); + pass.End(); + commands = encoder.Finish(); + } + + queue.Submit(1, &commands); + std::vector<uint32_t> expected = {1543, 15, 17, 123}; + + EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size()); +} + DAWN_INSTANTIATE_TEST(PolyfillBuiltinSimpleTests, D3D12Backend(), D3D11Backend(),
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.cc b/src/tint/lang/core/ir/transform/builtin_polyfill.cc index d576f96..19a0b2c 100644 --- a/src/tint/lang/core/ir/transform/builtin_polyfill.cc +++ b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
@@ -27,6 +27,7 @@ #include "src/tint/lang/core/ir/transform/builtin_polyfill.h" +#include "src/tint/lang/core/binary_op.h" #include "src/tint/lang/core/ir/builder.h" #include "src/tint/lang/core/ir/module.h" #include "src/tint/lang/core/ir/validator.h" @@ -76,6 +77,12 @@ worklist.Push(builtin); } break; + case core::BuiltinFn::kAbs: + if (config.abs_signed_int && + builtin->Result()->Type()->IsSignedIntegerScalarOrVector()) { + worklist.Push(builtin); + } + break; case core::BuiltinFn::kCountLeadingZeros: if (config.count_leading_zeros) { worklist.Push(builtin); @@ -194,6 +201,9 @@ case core::BuiltinFn::kClamp: ClampInt(builtin); break; + case core::BuiltinFn::kAbs: + AbsSignedInt(builtin); + break; case core::BuiltinFn::kCountLeadingZeros: CountLeadingZeros(builtin); break; @@ -401,6 +411,17 @@ call->Destroy(); } + /// Polyfill a `abs()` builtin call for signed integers. + /// @param call the builtin call instruction + void AbsSignedInt(ir::CoreBuiltinCall* call) { + auto* type = call->Result()->Type(); + auto* e = call->Args()[0]; + b.InsertBefore(call, [&] { + b.CallWithResult(call->DetachResult(), core::BuiltinFn::kMax, e, b.Negation(type, e)); + }); + call->Destroy(); + } + /// Polyfill a `countLeadingZeros()` builtin call. /// @param call the builtin call instruction void CountLeadingZeros(ir::CoreBuiltinCall* call) {
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.h b/src/tint/lang/core/ir/transform/builtin_polyfill.h index 94bbe0f..c60b777 100644 --- a/src/tint/lang/core/ir/transform/builtin_polyfill.h +++ b/src/tint/lang/core/ir/transform/builtin_polyfill.h
@@ -56,6 +56,8 @@ struct BuiltinPolyfillConfig { /// Should `clamp()` be polyfilled for integer values? bool clamp_int = false; + /// Should `abs()` be polyfilled for signed integer values? + bool abs_signed_int = false; /// Should `countLeadingZeros()` be polyfilled? bool count_leading_zeros = false; /// Should `countTrailingZeros()` be polyfilled?
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc b/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc index 9c63b3f..f9b6188 100644 --- a/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc +++ b/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
@@ -2884,5 +2884,53 @@ EXPECT_EQ(expect, str()); } +TEST_F(IR_BuiltinPolyfillTest, Absi32_NoPolyfill) { + Build(core::BuiltinFn::kAbs, ty.i32(), Vector{ty.i32()}); + auto* src = R"( +%foo = func(%arg:i32):i32 { + $B1: { + %result:i32 = abs %arg + ret %result + } +} +)"; + auto* expect = src; + + EXPECT_EQ(src, str()); + + BuiltinPolyfillConfig config; + config.abs_signed_int = false; + Run(BuiltinPolyfill, config); + EXPECT_EQ(expect, str()); +} + +TEST_F(IR_BuiltinPolyfillTest, Absi32_Enabled) { + Build(core::BuiltinFn::kAbs, ty.i32(), Vector{ty.i32()}); + auto* src = R"( +%foo = func(%arg:i32):i32 { + $B1: { + %result:i32 = abs %arg + ret %result + } +} +)"; + auto* expect = R"( +%foo = func(%arg:i32):i32 { + $B1: { + %3:i32 = negation %arg + %result:i32 = max %arg, %3 + ret %result + } +} +)"; + + EXPECT_EQ(src, str()); + + BuiltinPolyfillConfig config; + config.abs_signed_int = true; + Run(BuiltinPolyfill, config); + EXPECT_EQ(expect, str()); +} + } // namespace } // namespace tint::core::ir::transform
diff --git a/src/tint/lang/glsl/writer/builtin_test.cc b/src/tint/lang/glsl/writer/builtin_test.cc index d4f5b12..a5f3561 100644 --- a/src/tint/lang/glsl/writer/builtin_test.cc +++ b/src/tint/lang/glsl/writer/builtin_test.cc
@@ -62,7 +62,7 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; void main() { int x = 1; - int w = abs(x); + int w = max(x, int((~(uint(x)) + 1u))); } )"); }
diff --git a/src/tint/lang/glsl/writer/raise/raise.cc b/src/tint/lang/glsl/writer/raise/raise.cc index f92a744..b50f90b 100644 --- a/src/tint/lang/glsl/writer/raise/raise.cc +++ b/src/tint/lang/glsl/writer/raise/raise.cc
@@ -144,7 +144,7 @@ core_polyfills.dot_4x8_packed = true; core_polyfills.pack_unpack_4x8 = true; core_polyfills.pack_4xu8_clamp = true; - + core_polyfills.abs_signed_int = true; RUN_TRANSFORM(core::ir::transform::BuiltinPolyfill, module, core_polyfills); }
diff --git a/src/tint/lang/hlsl/writer/raise/raise.cc b/src/tint/lang/hlsl/writer/raise/raise.cc index e263f2c..9a14d3f 100644 --- a/src/tint/lang/hlsl/writer/raise/raise.cc +++ b/src/tint/lang/hlsl/writer/raise/raise.cc
@@ -147,6 +147,7 @@ core_polyfills.radians = true; core_polyfills.reflect_vec2_f32 = options.polyfill_reflect_vec2_f32; core_polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true; + core_polyfills.abs_signed_int = true; RUN_TRANSFORM(core::ir::transform::BuiltinPolyfill, module, core_polyfills); }
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc index 004d8ae..d24e68b 100644 --- a/src/tint/lang/msl/writer/raise/raise.cc +++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -107,6 +107,7 @@ core_polyfills.pack_4xu8_clamp = true; core_polyfills.radians = true; core_polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true; + core_polyfills.abs_signed_int = true; RUN_TRANSFORM(core::ir::transform::BuiltinPolyfill, module, core_polyfills); }
diff --git a/src/tint/lang/spirv/writer/builtin_test.cc b/src/tint/lang/spirv/writer/builtin_test.cc index f0965c5..f7f5ab1 100644 --- a/src/tint/lang/spirv/writer/builtin_test.cc +++ b/src/tint/lang/spirv/writer/builtin_test.cc
@@ -75,8 +75,7 @@ INSTANTIATE_TEST_SUITE_P( SpirvWriterTest, Builtin_1arg, - testing::Values(BuiltinTestCase{kI32, core::BuiltinFn::kAbs, "SAbs"}, - BuiltinTestCase{kF32, core::BuiltinFn::kAbs, "FAbs"}, + testing::Values(BuiltinTestCase{kF32, core::BuiltinFn::kAbs, "FAbs"}, BuiltinTestCase{kF16, core::BuiltinFn::kAbs, "FAbs"}, BuiltinTestCase{kF32, core::BuiltinFn::kAcos, "Acos"}, BuiltinTestCase{kF16, core::BuiltinFn::kAcos, "Acos"}, @@ -165,6 +164,26 @@ )"); } +// Test that abs of a signed integer is implemented as max(x,-x); +TEST_F(SpirvWriterTest, Builtin_Abs_i32) { + auto* func = b.Function("foo", MakeScalarType(kI32)); + b.Append(func->Block(), [&] { + auto* arg = MakeScalarValue(kI32); + auto* result = b.Call(MakeScalarType(kI32), core::BuiltinFn::kAbs, arg); + b.Return(func, result); + mod.SetName(arg, "arg"); + }); + + ASSERT_TRUE(Generate()) << Error() << output_; + EXPECT_INST(R"( + %6 = OpBitcast %uint %arg + %8 = OpNot %uint %6 + %9 = OpIAdd %uint %8 %uint_1 + %11 = OpBitcast %int %9 + %12 = OpExtInst %int %13 SMax %arg %11 +)"); +} + TEST_F(SpirvWriterTest, Builtin_Abs_vec2u) { auto* func = b.Function("foo", MakeVectorType(kU32)); b.Append(func->Block(), [&] {
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc index da2b0a8..18ee17c 100644 --- a/src/tint/lang/spirv/writer/raise/raise.cc +++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -117,6 +117,7 @@ core_polyfills.pack_unpack_4x8 = true; core_polyfills.pack_4xu8_clamp = true; core_polyfills.pack_unpack_4x8_norm = options.polyfill_pack_unpack_4x8_norm; + core_polyfills.abs_signed_int = true; RUN_TRANSFORM(core::ir::transform::BuiltinPolyfill, module, core_polyfills); core::ir::transform::ConversionPolyfillConfig conversion_polyfills;
diff --git a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.dxc.hlsl index 709b2b4..9fcf9db 100644 --- a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.dxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int abs_4ad288() { int arg_0 = int(1); - int res = abs(arg_0); + int v = arg_0; + int res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int abs_4ad288() { int arg_0 = int(1); - int res = abs(arg_0); + int v = arg_0; + int res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int abs_4ad288() { int arg_0 = int(1); - int res = abs(arg_0); + int v = arg_0; + int res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_4ad288(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_4ad288(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.fxc.hlsl index 709b2b4..9fcf9db 100644 --- a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.fxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int abs_4ad288() { int arg_0 = int(1); - int res = abs(arg_0); + int v = arg_0; + int res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int abs_4ad288() { int arg_0 = int(1); - int res = abs(arg_0); + int v = arg_0; + int res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int abs_4ad288() { int arg_0 = int(1); - int res = abs(arg_0); + int v = arg_0; + int res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_4ad288(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_4ad288(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.glsl b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.glsl index 35db619..9a1eb7c 100644 --- a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.glsl
@@ -11,7 +11,8 @@ } v; int abs_4ad288() { int arg_0 = 1; - int res = abs(arg_0); + int v_1 = arg_0; + int res = max(v_1, int((~(uint(v_1)) + 1u))); return res; } void main() { @@ -28,7 +29,8 @@ } v; int abs_4ad288() { int arg_0 = 1; - int res = abs(arg_0); + int v_1 = arg_0; + int res = max(v_1, int((~(uint(v_1)) + 1u))); return res; } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -49,18 +51,19 @@ layout(location = 0) flat out int tint_interstage_location0; int abs_4ad288() { int arg_0 = 1; - int res = abs(arg_0); + int v = arg_0; + int res = max(v, int((~(uint(v)) + 1u))); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = VertexOutput(vec4(0.0f), 0); - v.pos = vec4(0.0f); - v.prevent_dce = abs_4ad288(); - return v; + VertexOutput v_1 = VertexOutput(vec4(0.0f), 0); + v_1.pos = vec4(0.0f); + v_1.prevent_dce = abs_4ad288(); + return v_1; } void main() { - VertexOutput v_1 = vertex_main_inner(); - gl_Position = vec4(v_1.pos.x, -(v_1.pos.y), ((2.0f * v_1.pos.z) - v_1.pos.w), v_1.pos.w); - tint_interstage_location0 = v_1.prevent_dce; + VertexOutput v_2 = vertex_main_inner(); + gl_Position = vec4(v_2.pos.x, -(v_2.pos.y), ((2.0f * v_2.pos.z) - v_2.pos.w), v_2.pos.w); + tint_interstage_location0 = v_2.prevent_dce; gl_PointSize = 1.0f; }
diff --git a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.msl b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.msl index 328b22e..9bcd00b 100644 --- a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.msl
@@ -10,7 +10,8 @@ int abs_4ad288() { int arg_0 = 1; - int res = abs(arg_0); + int const v = arg_0; + int res = max(v, as_type<int>((~(as_type<uint>(v)) + 1u))); return res; } @@ -30,7 +31,8 @@ int abs_4ad288() { int arg_0 = 1; - int res = abs(arg_0); + int const v = arg_0; + int res = max(v, as_type<int>((~(as_type<uint>(v)) + 1u))); return res; } @@ -56,7 +58,8 @@ int abs_4ad288() { int arg_0 = 1; - int res = abs(arg_0); + int const v = arg_0; + int res = max(v, as_type<int>((~(as_type<uint>(v)) + 1u))); return res; } @@ -68,9 +71,9 @@ } vertex vertex_main_outputs vertex_main() { - VertexOutput const v = vertex_main_inner(); + VertexOutput const v_1 = vertex_main_inner(); vertex_main_outputs tint_wrapper_result = {}; - tint_wrapper_result.VertexOutput_pos = v.pos; - tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce; + tint_wrapper_result.VertexOutput_pos = v_1.pos; + tint_wrapper_result.VertexOutput_prevent_dce = v_1.prevent_dce; return tint_wrapper_result; }
diff --git a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.spvasm b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.spvasm index 1719c1f..907d11a 100644 --- a/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/abs/4ad288.wgsl.expected.spvasm
@@ -4,10 +4,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 25 +; Bound: 30 ; Schema: 0 OpCapability Shader - %13 = OpExtInstImport "GLSL.std.450" + %19 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Fragment %fragment_main "fragment_main" OpExecutionMode %fragment_main OriginUpperLeft @@ -29,10 +29,11 @@ %6 = OpTypeFunction %int %_ptr_Function_int = OpTypePointer Function %int %int_1 = OpConstant %int 1 - %void = OpTypeVoid - %18 = OpTypeFunction %void -%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %24 = OpTypeFunction %void +%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %uint_0 = OpConstant %uint 0 %abs_4ad288 = OpFunction %int None %6 %7 = OpLabel @@ -40,16 +41,20 @@ %res = OpVariable %_ptr_Function_int Function OpStore %arg_0 %int_1 %11 = OpLoad %int %arg_0 None - %12 = OpExtInst %int %13 SAbs %11 - OpStore %res %12 - %15 = OpLoad %int %res None - OpReturnValue %15 + %13 = OpBitcast %uint %11 + %14 = OpNot %uint %13 + %15 = OpIAdd %uint %14 %uint_1 + %17 = OpBitcast %int %15 + %18 = OpExtInst %int %19 SMax %11 %17 + OpStore %res %18 + %21 = OpLoad %int %res None + OpReturnValue %21 OpFunctionEnd -%fragment_main = OpFunction %void None %18 - %19 = OpLabel - %20 = OpFunctionCall %int %abs_4ad288 - %21 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0 - OpStore %21 %20 None +%fragment_main = OpFunction %void None %24 + %25 = OpLabel + %26 = OpFunctionCall %int %abs_4ad288 + %27 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0 + OpStore %27 %26 None OpReturn OpFunctionEnd ; @@ -58,10 +63,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 25 +; Bound: 30 ; Schema: 0 OpCapability Shader - %13 = OpExtInstImport "GLSL.std.450" + %19 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %compute_main "compute_main" OpExecutionMode %compute_main LocalSize 1 1 1 @@ -83,10 +88,11 @@ %6 = OpTypeFunction %int %_ptr_Function_int = OpTypePointer Function %int %int_1 = OpConstant %int 1 - %void = OpTypeVoid - %18 = OpTypeFunction %void -%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %24 = OpTypeFunction %void +%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %uint_0 = OpConstant %uint 0 %abs_4ad288 = OpFunction %int None %6 %7 = OpLabel @@ -94,16 +100,20 @@ %res = OpVariable %_ptr_Function_int Function OpStore %arg_0 %int_1 %11 = OpLoad %int %arg_0 None - %12 = OpExtInst %int %13 SAbs %11 - OpStore %res %12 - %15 = OpLoad %int %res None - OpReturnValue %15 + %13 = OpBitcast %uint %11 + %14 = OpNot %uint %13 + %15 = OpIAdd %uint %14 %uint_1 + %17 = OpBitcast %int %15 + %18 = OpExtInst %int %19 SMax %11 %17 + OpStore %res %18 + %21 = OpLoad %int %res None + OpReturnValue %21 OpFunctionEnd -%compute_main = OpFunction %void None %18 - %19 = OpLabel - %20 = OpFunctionCall %int %abs_4ad288 - %21 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0 - OpStore %21 %20 None +%compute_main = OpFunction %void None %24 + %25 = OpLabel + %26 = OpFunctionCall %int %abs_4ad288 + %27 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0 + OpStore %27 %26 None OpReturn OpFunctionEnd ; @@ -112,10 +122,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 45 +; Bound: 49 ; Schema: 0 OpCapability Shader - %18 = OpExtInstImport "GLSL.std.450" + %24 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output OpName %vertex_main_position_Output "vertex_main_position_Output" @@ -146,17 +156,17 @@ %11 = OpTypeFunction %int %_ptr_Function_int = OpTypePointer Function %int %int_1 = OpConstant %int 1 -%VertexOutput = OpTypeStruct %v4float %int - %23 = OpTypeFunction %VertexOutput -%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput - %27 = OpConstantNull %VertexOutput -%_ptr_Function_v4float = OpTypePointer Function %v4float %uint = OpTypeInt 32 0 - %uint_0 = OpConstant %uint 0 - %32 = OpConstantNull %v4float %uint_1 = OpConstant %uint 1 +%VertexOutput = OpTypeStruct %v4float %int + %29 = OpTypeFunction %VertexOutput +%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput + %33 = OpConstantNull %VertexOutput +%_ptr_Function_v4float = OpTypePointer Function %v4float + %uint_0 = OpConstant %uint 0 + %37 = OpConstantNull %v4float %void = OpTypeVoid - %39 = OpTypeFunction %void + %43 = OpTypeFunction %void %float_1 = OpConstant %float 1 %abs_4ad288 = OpFunction %int None %11 %12 = OpLabel @@ -164,29 +174,33 @@ %res = OpVariable %_ptr_Function_int Function OpStore %arg_0 %int_1 %16 = OpLoad %int %arg_0 None - %17 = OpExtInst %int %18 SAbs %16 - OpStore %res %17 - %20 = OpLoad %int %res None - OpReturnValue %20 + %18 = OpBitcast %uint %16 + %19 = OpNot %uint %18 + %20 = OpIAdd %uint %19 %uint_1 + %22 = OpBitcast %int %20 + %23 = OpExtInst %int %24 SMax %16 %22 + OpStore %res %23 + %26 = OpLoad %int %res None + OpReturnValue %26 OpFunctionEnd -%vertex_main_inner = OpFunction %VertexOutput None %23 - %24 = OpLabel - %out = OpVariable %_ptr_Function_VertexOutput Function %27 - %28 = OpAccessChain %_ptr_Function_v4float %out %uint_0 - OpStore %28 %32 None - %33 = OpAccessChain %_ptr_Function_int %out %uint_1 - %35 = OpFunctionCall %int %abs_4ad288 - OpStore %33 %35 None - %36 = OpLoad %VertexOutput %out None - OpReturnValue %36 +%vertex_main_inner = OpFunction %VertexOutput None %29 + %30 = OpLabel + %out = OpVariable %_ptr_Function_VertexOutput Function %33 + %34 = OpAccessChain %_ptr_Function_v4float %out %uint_0 + OpStore %34 %37 None + %38 = OpAccessChain %_ptr_Function_int %out %uint_1 + %39 = OpFunctionCall %int %abs_4ad288 + OpStore %38 %39 None + %40 = OpLoad %VertexOutput %out None + OpReturnValue %40 OpFunctionEnd -%vertex_main = OpFunction %void None %39 - %40 = OpLabel - %41 = OpFunctionCall %VertexOutput %vertex_main_inner - %42 = OpCompositeExtract %v4float %41 0 - OpStore %vertex_main_position_Output %42 None - %43 = OpCompositeExtract %int %41 1 - OpStore %vertex_main_loc0_Output %43 None +%vertex_main = OpFunction %void None %43 + %44 = OpLabel + %45 = OpFunctionCall %VertexOutput %vertex_main_inner + %46 = OpCompositeExtract %v4float %45 0 + OpStore %vertex_main_position_Output %46 None + %47 = OpCompositeExtract %int %45 1 + OpStore %vertex_main_loc0_Output %47 None OpStore %vertex_main___point_size_Output %float_1 None OpReturn OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.dxc.hlsl index 69ef654..4428295 100644 --- a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.dxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int3 abs_5ad50a() { int3 arg_0 = (int(1)).xxx; - int3 res = abs(arg_0); + int3 v = arg_0; + int3 res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int3 abs_5ad50a() { int3 arg_0 = (int(1)).xxx; - int3 res = abs(arg_0); + int3 v = arg_0; + int3 res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int3 abs_5ad50a() { int3 arg_0 = (int(1)).xxx; - int3 res = abs(arg_0); + int3 v = arg_0; + int3 res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_5ad50a(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_5ad50a(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.fxc.hlsl index 69ef654..4428295 100644 --- a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.fxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int3 abs_5ad50a() { int3 arg_0 = (int(1)).xxx; - int3 res = abs(arg_0); + int3 v = arg_0; + int3 res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int3 abs_5ad50a() { int3 arg_0 = (int(1)).xxx; - int3 res = abs(arg_0); + int3 v = arg_0; + int3 res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int3 abs_5ad50a() { int3 arg_0 = (int(1)).xxx; - int3 res = abs(arg_0); + int3 v = arg_0; + int3 res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_5ad50a(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_5ad50a(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.glsl b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.glsl index a50f658..cd82a9d 100644 --- a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.glsl
@@ -11,7 +11,8 @@ } v; ivec3 abs_5ad50a() { ivec3 arg_0 = ivec3(1); - ivec3 res = abs(arg_0); + ivec3 v_1 = arg_0; + ivec3 res = max(v_1, ivec3((~(uvec3(v_1)) + uvec3(1u)))); return res; } void main() { @@ -28,7 +29,8 @@ } v; ivec3 abs_5ad50a() { ivec3 arg_0 = ivec3(1); - ivec3 res = abs(arg_0); + ivec3 v_1 = arg_0; + ivec3 res = max(v_1, ivec3((~(uvec3(v_1)) + uvec3(1u)))); return res; } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -49,18 +51,19 @@ layout(location = 0) flat out ivec3 tint_interstage_location0; ivec3 abs_5ad50a() { ivec3 arg_0 = ivec3(1); - ivec3 res = abs(arg_0); + ivec3 v = arg_0; + ivec3 res = max(v, ivec3((~(uvec3(v)) + uvec3(1u)))); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = VertexOutput(vec4(0.0f), ivec3(0)); - v.pos = vec4(0.0f); - v.prevent_dce = abs_5ad50a(); - return v; + VertexOutput v_1 = VertexOutput(vec4(0.0f), ivec3(0)); + v_1.pos = vec4(0.0f); + v_1.prevent_dce = abs_5ad50a(); + return v_1; } void main() { - VertexOutput v_1 = vertex_main_inner(); - gl_Position = vec4(v_1.pos.x, -(v_1.pos.y), ((2.0f * v_1.pos.z) - v_1.pos.w), v_1.pos.w); - tint_interstage_location0 = v_1.prevent_dce; + VertexOutput v_2 = vertex_main_inner(); + gl_Position = vec4(v_2.pos.x, -(v_2.pos.y), ((2.0f * v_2.pos.z) - v_2.pos.w), v_2.pos.w); + tint_interstage_location0 = v_2.prevent_dce; gl_PointSize = 1.0f; }
diff --git a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.msl b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.msl index 8d64713..e6dca6c 100644 --- a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.msl
@@ -10,7 +10,8 @@ int3 abs_5ad50a() { int3 arg_0 = int3(1); - int3 res = abs(arg_0); + int3 const v = arg_0; + int3 res = max(v, as_type<int3>((~(as_type<uint3>(v)) + uint3(1u)))); return res; } @@ -30,7 +31,8 @@ int3 abs_5ad50a() { int3 arg_0 = int3(1); - int3 res = abs(arg_0); + int3 const v = arg_0; + int3 res = max(v, as_type<int3>((~(as_type<uint3>(v)) + uint3(1u)))); return res; } @@ -56,7 +58,8 @@ int3 abs_5ad50a() { int3 arg_0 = int3(1); - int3 res = abs(arg_0); + int3 const v = arg_0; + int3 res = max(v, as_type<int3>((~(as_type<uint3>(v)) + uint3(1u)))); return res; } @@ -68,9 +71,9 @@ } vertex vertex_main_outputs vertex_main() { - VertexOutput const v = vertex_main_inner(); + VertexOutput const v_1 = vertex_main_inner(); vertex_main_outputs tint_wrapper_result = {}; - tint_wrapper_result.VertexOutput_pos = v.pos; - tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce; + tint_wrapper_result.VertexOutput_pos = v_1.pos; + tint_wrapper_result.VertexOutput_prevent_dce = v_1.prevent_dce; return tint_wrapper_result; }
diff --git a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.spvasm b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.spvasm index 953e6f9..037efc1 100644 --- a/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/abs/5ad50a.wgsl.expected.spvasm
@@ -4,10 +4,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 27 +; Bound: 34 ; Schema: 0 OpCapability Shader - %15 = OpExtInstImport "GLSL.std.450" + %23 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Fragment %fragment_main "fragment_main" OpExecutionMode %fragment_main OriginUpperLeft @@ -31,10 +31,13 @@ %_ptr_Function_v3int = OpTypePointer Function %v3int %int_1 = OpConstant %int 1 %11 = OpConstantComposite %v3int %int_1 %int_1 %int_1 - %void = OpTypeVoid - %20 = OpTypeFunction %void -%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %void = OpTypeVoid + %28 = OpTypeFunction %void +%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int %uint_0 = OpConstant %uint 0 %abs_5ad50a = OpFunction %v3int None %7 %8 = OpLabel @@ -42,16 +45,20 @@ %res = OpVariable %_ptr_Function_v3int Function OpStore %arg_0 %11 %13 = OpLoad %v3int %arg_0 None - %14 = OpExtInst %v3int %15 SAbs %13 - OpStore %res %14 - %17 = OpLoad %v3int %res None - OpReturnValue %17 + %16 = OpBitcast %v3uint %13 + %17 = OpNot %v3uint %16 + %18 = OpIAdd %v3uint %17 %19 + %21 = OpBitcast %v3int %18 + %22 = OpExtInst %v3int %23 SMax %13 %21 + OpStore %res %22 + %25 = OpLoad %v3int %res None + OpReturnValue %25 OpFunctionEnd -%fragment_main = OpFunction %void None %20 - %21 = OpLabel - %22 = OpFunctionCall %v3int %abs_5ad50a - %23 = OpAccessChain %_ptr_StorageBuffer_v3int %1 %uint_0 - OpStore %23 %22 None +%fragment_main = OpFunction %void None %28 + %29 = OpLabel + %30 = OpFunctionCall %v3int %abs_5ad50a + %31 = OpAccessChain %_ptr_StorageBuffer_v3int %1 %uint_0 + OpStore %31 %30 None OpReturn OpFunctionEnd ; @@ -60,10 +67,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 27 +; Bound: 34 ; Schema: 0 OpCapability Shader - %15 = OpExtInstImport "GLSL.std.450" + %23 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %compute_main "compute_main" OpExecutionMode %compute_main LocalSize 1 1 1 @@ -87,10 +94,13 @@ %_ptr_Function_v3int = OpTypePointer Function %v3int %int_1 = OpConstant %int 1 %11 = OpConstantComposite %v3int %int_1 %int_1 %int_1 - %void = OpTypeVoid - %20 = OpTypeFunction %void -%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %void = OpTypeVoid + %28 = OpTypeFunction %void +%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int %uint_0 = OpConstant %uint 0 %abs_5ad50a = OpFunction %v3int None %7 %8 = OpLabel @@ -98,16 +108,20 @@ %res = OpVariable %_ptr_Function_v3int Function OpStore %arg_0 %11 %13 = OpLoad %v3int %arg_0 None - %14 = OpExtInst %v3int %15 SAbs %13 - OpStore %res %14 - %17 = OpLoad %v3int %res None - OpReturnValue %17 + %16 = OpBitcast %v3uint %13 + %17 = OpNot %v3uint %16 + %18 = OpIAdd %v3uint %17 %19 + %21 = OpBitcast %v3int %18 + %22 = OpExtInst %v3int %23 SMax %13 %21 + OpStore %res %22 + %25 = OpLoad %v3int %res None + OpReturnValue %25 OpFunctionEnd -%compute_main = OpFunction %void None %20 - %21 = OpLabel - %22 = OpFunctionCall %v3int %abs_5ad50a - %23 = OpAccessChain %_ptr_StorageBuffer_v3int %1 %uint_0 - OpStore %23 %22 None +%compute_main = OpFunction %void None %28 + %29 = OpLabel + %30 = OpFunctionCall %v3int %abs_5ad50a + %31 = OpAccessChain %_ptr_StorageBuffer_v3int %1 %uint_0 + OpStore %31 %30 None OpReturn OpFunctionEnd ; @@ -116,10 +130,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 47 +; Bound: 53 ; Schema: 0 OpCapability Shader - %20 = OpExtInstImport "GLSL.std.450" + %28 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output OpName %vertex_main_position_Output "vertex_main_position_Output" @@ -152,17 +166,19 @@ %_ptr_Function_v3int = OpTypePointer Function %v3int %int_1 = OpConstant %int 1 %16 = OpConstantComposite %v3int %int_1 %int_1 %int_1 -%VertexOutput = OpTypeStruct %v4float %v3int - %25 = OpTypeFunction %VertexOutput -%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput - %29 = OpConstantNull %VertexOutput -%_ptr_Function_v4float = OpTypePointer Function %v4float %uint = OpTypeInt 32 0 - %uint_0 = OpConstant %uint 0 - %34 = OpConstantNull %v4float + %v3uint = OpTypeVector %uint 3 %uint_1 = OpConstant %uint 1 + %24 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 +%VertexOutput = OpTypeStruct %v4float %v3int + %33 = OpTypeFunction %VertexOutput +%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput + %37 = OpConstantNull %VertexOutput +%_ptr_Function_v4float = OpTypePointer Function %v4float + %uint_0 = OpConstant %uint 0 + %41 = OpConstantNull %v4float %void = OpTypeVoid - %41 = OpTypeFunction %void + %47 = OpTypeFunction %void %float_1 = OpConstant %float 1 %abs_5ad50a = OpFunction %v3int None %12 %13 = OpLabel @@ -170,29 +186,33 @@ %res = OpVariable %_ptr_Function_v3int Function OpStore %arg_0 %16 %18 = OpLoad %v3int %arg_0 None - %19 = OpExtInst %v3int %20 SAbs %18 - OpStore %res %19 - %22 = OpLoad %v3int %res None - OpReturnValue %22 + %21 = OpBitcast %v3uint %18 + %22 = OpNot %v3uint %21 + %23 = OpIAdd %v3uint %22 %24 + %26 = OpBitcast %v3int %23 + %27 = OpExtInst %v3int %28 SMax %18 %26 + OpStore %res %27 + %30 = OpLoad %v3int %res None + OpReturnValue %30 OpFunctionEnd -%vertex_main_inner = OpFunction %VertexOutput None %25 - %26 = OpLabel - %out = OpVariable %_ptr_Function_VertexOutput Function %29 - %30 = OpAccessChain %_ptr_Function_v4float %out %uint_0 - OpStore %30 %34 None - %35 = OpAccessChain %_ptr_Function_v3int %out %uint_1 - %37 = OpFunctionCall %v3int %abs_5ad50a - OpStore %35 %37 None - %38 = OpLoad %VertexOutput %out None - OpReturnValue %38 +%vertex_main_inner = OpFunction %VertexOutput None %33 + %34 = OpLabel + %out = OpVariable %_ptr_Function_VertexOutput Function %37 + %38 = OpAccessChain %_ptr_Function_v4float %out %uint_0 + OpStore %38 %41 None + %42 = OpAccessChain %_ptr_Function_v3int %out %uint_1 + %43 = OpFunctionCall %v3int %abs_5ad50a + OpStore %42 %43 None + %44 = OpLoad %VertexOutput %out None + OpReturnValue %44 OpFunctionEnd -%vertex_main = OpFunction %void None %41 - %42 = OpLabel - %43 = OpFunctionCall %VertexOutput %vertex_main_inner - %44 = OpCompositeExtract %v4float %43 0 - OpStore %vertex_main_position_Output %44 None - %45 = OpCompositeExtract %v3int %43 1 - OpStore %vertex_main_loc0_Output %45 None +%vertex_main = OpFunction %void None %47 + %48 = OpLabel + %49 = OpFunctionCall %VertexOutput %vertex_main_inner + %50 = OpCompositeExtract %v4float %49 0 + OpStore %vertex_main_position_Output %50 None + %51 = OpCompositeExtract %v3int %49 1 + OpStore %vertex_main_loc0_Output %51 None OpStore %vertex_main___point_size_Output %float_1 None OpReturn OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.dxc.hlsl index 513011d..e6f6136 100644 --- a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.dxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int2 abs_7faa9e() { int2 arg_0 = (int(1)).xx; - int2 res = abs(arg_0); + int2 v = arg_0; + int2 res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int2 abs_7faa9e() { int2 arg_0 = (int(1)).xx; - int2 res = abs(arg_0); + int2 v = arg_0; + int2 res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int2 abs_7faa9e() { int2 arg_0 = (int(1)).xx; - int2 res = abs(arg_0); + int2 v = arg_0; + int2 res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_7faa9e(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_7faa9e(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.fxc.hlsl index 513011d..e6f6136 100644 --- a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.fxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int2 abs_7faa9e() { int2 arg_0 = (int(1)).xx; - int2 res = abs(arg_0); + int2 v = arg_0; + int2 res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int2 abs_7faa9e() { int2 arg_0 = (int(1)).xx; - int2 res = abs(arg_0); + int2 v = arg_0; + int2 res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int2 abs_7faa9e() { int2 arg_0 = (int(1)).xx; - int2 res = abs(arg_0); + int2 v = arg_0; + int2 res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_7faa9e(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_7faa9e(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.glsl b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.glsl index be1cca3..33e7773 100644 --- a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.glsl
@@ -11,7 +11,8 @@ } v; ivec2 abs_7faa9e() { ivec2 arg_0 = ivec2(1); - ivec2 res = abs(arg_0); + ivec2 v_1 = arg_0; + ivec2 res = max(v_1, ivec2((~(uvec2(v_1)) + uvec2(1u)))); return res; } void main() { @@ -28,7 +29,8 @@ } v; ivec2 abs_7faa9e() { ivec2 arg_0 = ivec2(1); - ivec2 res = abs(arg_0); + ivec2 v_1 = arg_0; + ivec2 res = max(v_1, ivec2((~(uvec2(v_1)) + uvec2(1u)))); return res; } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -49,18 +51,19 @@ layout(location = 0) flat out ivec2 tint_interstage_location0; ivec2 abs_7faa9e() { ivec2 arg_0 = ivec2(1); - ivec2 res = abs(arg_0); + ivec2 v = arg_0; + ivec2 res = max(v, ivec2((~(uvec2(v)) + uvec2(1u)))); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = VertexOutput(vec4(0.0f), ivec2(0)); - v.pos = vec4(0.0f); - v.prevent_dce = abs_7faa9e(); - return v; + VertexOutput v_1 = VertexOutput(vec4(0.0f), ivec2(0)); + v_1.pos = vec4(0.0f); + v_1.prevent_dce = abs_7faa9e(); + return v_1; } void main() { - VertexOutput v_1 = vertex_main_inner(); - gl_Position = vec4(v_1.pos.x, -(v_1.pos.y), ((2.0f * v_1.pos.z) - v_1.pos.w), v_1.pos.w); - tint_interstage_location0 = v_1.prevent_dce; + VertexOutput v_2 = vertex_main_inner(); + gl_Position = vec4(v_2.pos.x, -(v_2.pos.y), ((2.0f * v_2.pos.z) - v_2.pos.w), v_2.pos.w); + tint_interstage_location0 = v_2.prevent_dce; gl_PointSize = 1.0f; }
diff --git a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.msl b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.msl index 01e9f7c..03022a1 100644 --- a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.msl
@@ -10,7 +10,8 @@ int2 abs_7faa9e() { int2 arg_0 = int2(1); - int2 res = abs(arg_0); + int2 const v = arg_0; + int2 res = max(v, as_type<int2>((~(as_type<uint2>(v)) + uint2(1u)))); return res; } @@ -30,7 +31,8 @@ int2 abs_7faa9e() { int2 arg_0 = int2(1); - int2 res = abs(arg_0); + int2 const v = arg_0; + int2 res = max(v, as_type<int2>((~(as_type<uint2>(v)) + uint2(1u)))); return res; } @@ -56,7 +58,8 @@ int2 abs_7faa9e() { int2 arg_0 = int2(1); - int2 res = abs(arg_0); + int2 const v = arg_0; + int2 res = max(v, as_type<int2>((~(as_type<uint2>(v)) + uint2(1u)))); return res; } @@ -68,9 +71,9 @@ } vertex vertex_main_outputs vertex_main() { - VertexOutput const v = vertex_main_inner(); + VertexOutput const v_1 = vertex_main_inner(); vertex_main_outputs tint_wrapper_result = {}; - tint_wrapper_result.VertexOutput_pos = v.pos; - tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce; + tint_wrapper_result.VertexOutput_pos = v_1.pos; + tint_wrapper_result.VertexOutput_prevent_dce = v_1.prevent_dce; return tint_wrapper_result; }
diff --git a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.spvasm b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.spvasm index 087b066..4cb4293 100644 --- a/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/abs/7faa9e.wgsl.expected.spvasm
@@ -4,10 +4,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 27 +; Bound: 34 ; Schema: 0 OpCapability Shader - %15 = OpExtInstImport "GLSL.std.450" + %23 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Fragment %fragment_main "fragment_main" OpExecutionMode %fragment_main OriginUpperLeft @@ -31,10 +31,13 @@ %_ptr_Function_v2int = OpTypePointer Function %v2int %int_1 = OpConstant %int 1 %11 = OpConstantComposite %v2int %int_1 %int_1 - %void = OpTypeVoid - %20 = OpTypeFunction %void -%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int %uint = OpTypeInt 32 0 + %v2uint = OpTypeVector %uint 2 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v2uint %uint_1 %uint_1 + %void = OpTypeVoid + %28 = OpTypeFunction %void +%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int %uint_0 = OpConstant %uint 0 %abs_7faa9e = OpFunction %v2int None %7 %8 = OpLabel @@ -42,16 +45,20 @@ %res = OpVariable %_ptr_Function_v2int Function OpStore %arg_0 %11 %13 = OpLoad %v2int %arg_0 None - %14 = OpExtInst %v2int %15 SAbs %13 - OpStore %res %14 - %17 = OpLoad %v2int %res None - OpReturnValue %17 + %16 = OpBitcast %v2uint %13 + %17 = OpNot %v2uint %16 + %18 = OpIAdd %v2uint %17 %19 + %21 = OpBitcast %v2int %18 + %22 = OpExtInst %v2int %23 SMax %13 %21 + OpStore %res %22 + %25 = OpLoad %v2int %res None + OpReturnValue %25 OpFunctionEnd -%fragment_main = OpFunction %void None %20 - %21 = OpLabel - %22 = OpFunctionCall %v2int %abs_7faa9e - %23 = OpAccessChain %_ptr_StorageBuffer_v2int %1 %uint_0 - OpStore %23 %22 None +%fragment_main = OpFunction %void None %28 + %29 = OpLabel + %30 = OpFunctionCall %v2int %abs_7faa9e + %31 = OpAccessChain %_ptr_StorageBuffer_v2int %1 %uint_0 + OpStore %31 %30 None OpReturn OpFunctionEnd ; @@ -60,10 +67,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 27 +; Bound: 34 ; Schema: 0 OpCapability Shader - %15 = OpExtInstImport "GLSL.std.450" + %23 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %compute_main "compute_main" OpExecutionMode %compute_main LocalSize 1 1 1 @@ -87,10 +94,13 @@ %_ptr_Function_v2int = OpTypePointer Function %v2int %int_1 = OpConstant %int 1 %11 = OpConstantComposite %v2int %int_1 %int_1 - %void = OpTypeVoid - %20 = OpTypeFunction %void -%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int %uint = OpTypeInt 32 0 + %v2uint = OpTypeVector %uint 2 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v2uint %uint_1 %uint_1 + %void = OpTypeVoid + %28 = OpTypeFunction %void +%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int %uint_0 = OpConstant %uint 0 %abs_7faa9e = OpFunction %v2int None %7 %8 = OpLabel @@ -98,16 +108,20 @@ %res = OpVariable %_ptr_Function_v2int Function OpStore %arg_0 %11 %13 = OpLoad %v2int %arg_0 None - %14 = OpExtInst %v2int %15 SAbs %13 - OpStore %res %14 - %17 = OpLoad %v2int %res None - OpReturnValue %17 + %16 = OpBitcast %v2uint %13 + %17 = OpNot %v2uint %16 + %18 = OpIAdd %v2uint %17 %19 + %21 = OpBitcast %v2int %18 + %22 = OpExtInst %v2int %23 SMax %13 %21 + OpStore %res %22 + %25 = OpLoad %v2int %res None + OpReturnValue %25 OpFunctionEnd -%compute_main = OpFunction %void None %20 - %21 = OpLabel - %22 = OpFunctionCall %v2int %abs_7faa9e - %23 = OpAccessChain %_ptr_StorageBuffer_v2int %1 %uint_0 - OpStore %23 %22 None +%compute_main = OpFunction %void None %28 + %29 = OpLabel + %30 = OpFunctionCall %v2int %abs_7faa9e + %31 = OpAccessChain %_ptr_StorageBuffer_v2int %1 %uint_0 + OpStore %31 %30 None OpReturn OpFunctionEnd ; @@ -116,10 +130,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 47 +; Bound: 53 ; Schema: 0 OpCapability Shader - %20 = OpExtInstImport "GLSL.std.450" + %28 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output OpName %vertex_main_position_Output "vertex_main_position_Output" @@ -152,17 +166,19 @@ %_ptr_Function_v2int = OpTypePointer Function %v2int %int_1 = OpConstant %int 1 %16 = OpConstantComposite %v2int %int_1 %int_1 -%VertexOutput = OpTypeStruct %v4float %v2int - %25 = OpTypeFunction %VertexOutput -%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput - %29 = OpConstantNull %VertexOutput -%_ptr_Function_v4float = OpTypePointer Function %v4float %uint = OpTypeInt 32 0 - %uint_0 = OpConstant %uint 0 - %34 = OpConstantNull %v4float + %v2uint = OpTypeVector %uint 2 %uint_1 = OpConstant %uint 1 + %24 = OpConstantComposite %v2uint %uint_1 %uint_1 +%VertexOutput = OpTypeStruct %v4float %v2int + %33 = OpTypeFunction %VertexOutput +%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput + %37 = OpConstantNull %VertexOutput +%_ptr_Function_v4float = OpTypePointer Function %v4float + %uint_0 = OpConstant %uint 0 + %41 = OpConstantNull %v4float %void = OpTypeVoid - %41 = OpTypeFunction %void + %47 = OpTypeFunction %void %float_1 = OpConstant %float 1 %abs_7faa9e = OpFunction %v2int None %12 %13 = OpLabel @@ -170,29 +186,33 @@ %res = OpVariable %_ptr_Function_v2int Function OpStore %arg_0 %16 %18 = OpLoad %v2int %arg_0 None - %19 = OpExtInst %v2int %20 SAbs %18 - OpStore %res %19 - %22 = OpLoad %v2int %res None - OpReturnValue %22 + %21 = OpBitcast %v2uint %18 + %22 = OpNot %v2uint %21 + %23 = OpIAdd %v2uint %22 %24 + %26 = OpBitcast %v2int %23 + %27 = OpExtInst %v2int %28 SMax %18 %26 + OpStore %res %27 + %30 = OpLoad %v2int %res None + OpReturnValue %30 OpFunctionEnd -%vertex_main_inner = OpFunction %VertexOutput None %25 - %26 = OpLabel - %out = OpVariable %_ptr_Function_VertexOutput Function %29 - %30 = OpAccessChain %_ptr_Function_v4float %out %uint_0 - OpStore %30 %34 None - %35 = OpAccessChain %_ptr_Function_v2int %out %uint_1 - %37 = OpFunctionCall %v2int %abs_7faa9e - OpStore %35 %37 None - %38 = OpLoad %VertexOutput %out None - OpReturnValue %38 +%vertex_main_inner = OpFunction %VertexOutput None %33 + %34 = OpLabel + %out = OpVariable %_ptr_Function_VertexOutput Function %37 + %38 = OpAccessChain %_ptr_Function_v4float %out %uint_0 + OpStore %38 %41 None + %42 = OpAccessChain %_ptr_Function_v2int %out %uint_1 + %43 = OpFunctionCall %v2int %abs_7faa9e + OpStore %42 %43 None + %44 = OpLoad %VertexOutput %out None + OpReturnValue %44 OpFunctionEnd -%vertex_main = OpFunction %void None %41 - %42 = OpLabel - %43 = OpFunctionCall %VertexOutput %vertex_main_inner - %44 = OpCompositeExtract %v4float %43 0 - OpStore %vertex_main_position_Output %44 None - %45 = OpCompositeExtract %v2int %43 1 - OpStore %vertex_main_loc0_Output %45 None +%vertex_main = OpFunction %void None %47 + %48 = OpLabel + %49 = OpFunctionCall %VertexOutput %vertex_main_inner + %50 = OpCompositeExtract %v4float %49 0 + OpStore %vertex_main_position_Output %50 None + %51 = OpCompositeExtract %v2int %49 1 + OpStore %vertex_main_loc0_Output %51 None OpStore %vertex_main___point_size_Output %float_1 None OpReturn OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.dxc.hlsl index 61f16c8..73e26e3 100644 --- a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.dxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int4 abs_9c80a6() { int4 arg_0 = (int(1)).xxxx; - int4 res = abs(arg_0); + int4 v = arg_0; + int4 res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int4 abs_9c80a6() { int4 arg_0 = (int(1)).xxxx; - int4 res = abs(arg_0); + int4 v = arg_0; + int4 res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int4 abs_9c80a6() { int4 arg_0 = (int(1)).xxxx; - int4 res = abs(arg_0); + int4 v = arg_0; + int4 res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_9c80a6(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_9c80a6(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.fxc.hlsl index 61f16c8..73e26e3 100644 --- a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.fxc.hlsl
@@ -5,7 +5,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int4 abs_9c80a6() { int4 arg_0 = (int(1)).xxxx; - int4 res = abs(arg_0); + int4 v = arg_0; + int4 res = max(v, -(v)); return res; } @@ -20,7 +21,8 @@ RWByteAddressBuffer prevent_dce : register(u0); int4 abs_9c80a6() { int4 arg_0 = (int(1)).xxxx; - int4 res = abs(arg_0); + int4 v = arg_0; + int4 res = max(v, -(v)); return res; } @@ -45,21 +47,22 @@ int4 abs_9c80a6() { int4 arg_0 = (int(1)).xxxx; - int4 res = abs(arg_0); + int4 v = arg_0; + int4 res = max(v, -(v)); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = (VertexOutput)0; - v.pos = (0.0f).xxxx; - v.prevent_dce = abs_9c80a6(); - VertexOutput v_1 = v; - return v_1; + VertexOutput v_1 = (VertexOutput)0; + v_1.pos = (0.0f).xxxx; + v_1.prevent_dce = abs_9c80a6(); + VertexOutput v_2 = v_1; + return v_2; } vertex_main_outputs vertex_main() { - VertexOutput v_2 = vertex_main_inner(); - vertex_main_outputs v_3 = {v_2.prevent_dce, v_2.pos}; - return v_3; + VertexOutput v_3 = vertex_main_inner(); + vertex_main_outputs v_4 = {v_3.prevent_dce, v_3.pos}; + return v_4; }
diff --git a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.glsl b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.glsl index b954350..860c92d 100644 --- a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.glsl
@@ -11,7 +11,8 @@ } v; ivec4 abs_9c80a6() { ivec4 arg_0 = ivec4(1); - ivec4 res = abs(arg_0); + ivec4 v_1 = arg_0; + ivec4 res = max(v_1, ivec4((~(uvec4(v_1)) + uvec4(1u)))); return res; } void main() { @@ -28,7 +29,8 @@ } v; ivec4 abs_9c80a6() { ivec4 arg_0 = ivec4(1); - ivec4 res = abs(arg_0); + ivec4 v_1 = arg_0; + ivec4 res = max(v_1, ivec4((~(uvec4(v_1)) + uvec4(1u)))); return res; } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -49,18 +51,19 @@ layout(location = 0) flat out ivec4 tint_interstage_location0; ivec4 abs_9c80a6() { ivec4 arg_0 = ivec4(1); - ivec4 res = abs(arg_0); + ivec4 v = arg_0; + ivec4 res = max(v, ivec4((~(uvec4(v)) + uvec4(1u)))); return res; } VertexOutput vertex_main_inner() { - VertexOutput v = VertexOutput(vec4(0.0f), ivec4(0)); - v.pos = vec4(0.0f); - v.prevent_dce = abs_9c80a6(); - return v; + VertexOutput v_1 = VertexOutput(vec4(0.0f), ivec4(0)); + v_1.pos = vec4(0.0f); + v_1.prevent_dce = abs_9c80a6(); + return v_1; } void main() { - VertexOutput v_1 = vertex_main_inner(); - gl_Position = vec4(v_1.pos.x, -(v_1.pos.y), ((2.0f * v_1.pos.z) - v_1.pos.w), v_1.pos.w); - tint_interstage_location0 = v_1.prevent_dce; + VertexOutput v_2 = vertex_main_inner(); + gl_Position = vec4(v_2.pos.x, -(v_2.pos.y), ((2.0f * v_2.pos.z) - v_2.pos.w), v_2.pos.w); + tint_interstage_location0 = v_2.prevent_dce; gl_PointSize = 1.0f; }
diff --git a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.msl b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.msl index 237e8cf..0e82aac 100644 --- a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.msl
@@ -10,7 +10,8 @@ int4 abs_9c80a6() { int4 arg_0 = int4(1); - int4 res = abs(arg_0); + int4 const v = arg_0; + int4 res = max(v, as_type<int4>((~(as_type<uint4>(v)) + uint4(1u)))); return res; } @@ -30,7 +31,8 @@ int4 abs_9c80a6() { int4 arg_0 = int4(1); - int4 res = abs(arg_0); + int4 const v = arg_0; + int4 res = max(v, as_type<int4>((~(as_type<uint4>(v)) + uint4(1u)))); return res; } @@ -56,7 +58,8 @@ int4 abs_9c80a6() { int4 arg_0 = int4(1); - int4 res = abs(arg_0); + int4 const v = arg_0; + int4 res = max(v, as_type<int4>((~(as_type<uint4>(v)) + uint4(1u)))); return res; } @@ -68,9 +71,9 @@ } vertex vertex_main_outputs vertex_main() { - VertexOutput const v = vertex_main_inner(); + VertexOutput const v_1 = vertex_main_inner(); vertex_main_outputs tint_wrapper_result = {}; - tint_wrapper_result.VertexOutput_pos = v.pos; - tint_wrapper_result.VertexOutput_prevent_dce = v.prevent_dce; + tint_wrapper_result.VertexOutput_pos = v_1.pos; + tint_wrapper_result.VertexOutput_prevent_dce = v_1.prevent_dce; return tint_wrapper_result; }
diff --git a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.spvasm b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.spvasm index 95b4e00..da67cee 100644 --- a/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/abs/9c80a6.wgsl.expected.spvasm
@@ -4,10 +4,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 27 +; Bound: 34 ; Schema: 0 OpCapability Shader - %15 = OpExtInstImport "GLSL.std.450" + %23 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Fragment %fragment_main "fragment_main" OpExecutionMode %fragment_main OriginUpperLeft @@ -31,10 +31,13 @@ %_ptr_Function_v4int = OpTypePointer Function %v4int %int_1 = OpConstant %int 1 %11 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1 - %void = OpTypeVoid - %20 = OpTypeFunction %void -%_ptr_StorageBuffer_v4int = OpTypePointer StorageBuffer %v4int %uint = OpTypeInt 32 0 + %v4uint = OpTypeVector %uint 4 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 + %void = OpTypeVoid + %28 = OpTypeFunction %void +%_ptr_StorageBuffer_v4int = OpTypePointer StorageBuffer %v4int %uint_0 = OpConstant %uint 0 %abs_9c80a6 = OpFunction %v4int None %7 %8 = OpLabel @@ -42,16 +45,20 @@ %res = OpVariable %_ptr_Function_v4int Function OpStore %arg_0 %11 %13 = OpLoad %v4int %arg_0 None - %14 = OpExtInst %v4int %15 SAbs %13 - OpStore %res %14 - %17 = OpLoad %v4int %res None - OpReturnValue %17 + %16 = OpBitcast %v4uint %13 + %17 = OpNot %v4uint %16 + %18 = OpIAdd %v4uint %17 %19 + %21 = OpBitcast %v4int %18 + %22 = OpExtInst %v4int %23 SMax %13 %21 + OpStore %res %22 + %25 = OpLoad %v4int %res None + OpReturnValue %25 OpFunctionEnd -%fragment_main = OpFunction %void None %20 - %21 = OpLabel - %22 = OpFunctionCall %v4int %abs_9c80a6 - %23 = OpAccessChain %_ptr_StorageBuffer_v4int %1 %uint_0 - OpStore %23 %22 None +%fragment_main = OpFunction %void None %28 + %29 = OpLabel + %30 = OpFunctionCall %v4int %abs_9c80a6 + %31 = OpAccessChain %_ptr_StorageBuffer_v4int %1 %uint_0 + OpStore %31 %30 None OpReturn OpFunctionEnd ; @@ -60,10 +67,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 27 +; Bound: 34 ; Schema: 0 OpCapability Shader - %15 = OpExtInstImport "GLSL.std.450" + %23 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %compute_main "compute_main" OpExecutionMode %compute_main LocalSize 1 1 1 @@ -87,10 +94,13 @@ %_ptr_Function_v4int = OpTypePointer Function %v4int %int_1 = OpConstant %int 1 %11 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1 - %void = OpTypeVoid - %20 = OpTypeFunction %void -%_ptr_StorageBuffer_v4int = OpTypePointer StorageBuffer %v4int %uint = OpTypeInt 32 0 + %v4uint = OpTypeVector %uint 4 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 + %void = OpTypeVoid + %28 = OpTypeFunction %void +%_ptr_StorageBuffer_v4int = OpTypePointer StorageBuffer %v4int %uint_0 = OpConstant %uint 0 %abs_9c80a6 = OpFunction %v4int None %7 %8 = OpLabel @@ -98,16 +108,20 @@ %res = OpVariable %_ptr_Function_v4int Function OpStore %arg_0 %11 %13 = OpLoad %v4int %arg_0 None - %14 = OpExtInst %v4int %15 SAbs %13 - OpStore %res %14 - %17 = OpLoad %v4int %res None - OpReturnValue %17 + %16 = OpBitcast %v4uint %13 + %17 = OpNot %v4uint %16 + %18 = OpIAdd %v4uint %17 %19 + %21 = OpBitcast %v4int %18 + %22 = OpExtInst %v4int %23 SMax %13 %21 + OpStore %res %22 + %25 = OpLoad %v4int %res None + OpReturnValue %25 OpFunctionEnd -%compute_main = OpFunction %void None %20 - %21 = OpLabel - %22 = OpFunctionCall %v4int %abs_9c80a6 - %23 = OpAccessChain %_ptr_StorageBuffer_v4int %1 %uint_0 - OpStore %23 %22 None +%compute_main = OpFunction %void None %28 + %29 = OpLabel + %30 = OpFunctionCall %v4int %abs_9c80a6 + %31 = OpAccessChain %_ptr_StorageBuffer_v4int %1 %uint_0 + OpStore %31 %30 None OpReturn OpFunctionEnd ; @@ -116,10 +130,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 1 -; Bound: 47 +; Bound: 53 ; Schema: 0 OpCapability Shader - %20 = OpExtInstImport "GLSL.std.450" + %28 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %vertex_main "vertex_main" %vertex_main_position_Output %vertex_main_loc0_Output %vertex_main___point_size_Output OpName %vertex_main_position_Output "vertex_main_position_Output" @@ -152,17 +166,19 @@ %_ptr_Function_v4int = OpTypePointer Function %v4int %int_1 = OpConstant %int 1 %16 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1 -%VertexOutput = OpTypeStruct %v4float %v4int - %25 = OpTypeFunction %VertexOutput -%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput - %29 = OpConstantNull %VertexOutput -%_ptr_Function_v4float = OpTypePointer Function %v4float %uint = OpTypeInt 32 0 - %uint_0 = OpConstant %uint 0 - %34 = OpConstantNull %v4float + %v4uint = OpTypeVector %uint 4 %uint_1 = OpConstant %uint 1 + %24 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 +%VertexOutput = OpTypeStruct %v4float %v4int + %33 = OpTypeFunction %VertexOutput +%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput + %37 = OpConstantNull %VertexOutput +%_ptr_Function_v4float = OpTypePointer Function %v4float + %uint_0 = OpConstant %uint 0 + %41 = OpConstantNull %v4float %void = OpTypeVoid - %41 = OpTypeFunction %void + %47 = OpTypeFunction %void %float_1 = OpConstant %float 1 %abs_9c80a6 = OpFunction %v4int None %12 %13 = OpLabel @@ -170,29 +186,33 @@ %res = OpVariable %_ptr_Function_v4int Function OpStore %arg_0 %16 %18 = OpLoad %v4int %arg_0 None - %19 = OpExtInst %v4int %20 SAbs %18 - OpStore %res %19 - %22 = OpLoad %v4int %res None - OpReturnValue %22 + %21 = OpBitcast %v4uint %18 + %22 = OpNot %v4uint %21 + %23 = OpIAdd %v4uint %22 %24 + %26 = OpBitcast %v4int %23 + %27 = OpExtInst %v4int %28 SMax %18 %26 + OpStore %res %27 + %30 = OpLoad %v4int %res None + OpReturnValue %30 OpFunctionEnd -%vertex_main_inner = OpFunction %VertexOutput None %25 - %26 = OpLabel - %out = OpVariable %_ptr_Function_VertexOutput Function %29 - %30 = OpAccessChain %_ptr_Function_v4float %out %uint_0 - OpStore %30 %34 None - %35 = OpAccessChain %_ptr_Function_v4int %out %uint_1 - %37 = OpFunctionCall %v4int %abs_9c80a6 - OpStore %35 %37 None - %38 = OpLoad %VertexOutput %out None - OpReturnValue %38 +%vertex_main_inner = OpFunction %VertexOutput None %33 + %34 = OpLabel + %out = OpVariable %_ptr_Function_VertexOutput Function %37 + %38 = OpAccessChain %_ptr_Function_v4float %out %uint_0 + OpStore %38 %41 None + %42 = OpAccessChain %_ptr_Function_v4int %out %uint_1 + %43 = OpFunctionCall %v4int %abs_9c80a6 + OpStore %42 %43 None + %44 = OpLoad %VertexOutput %out None + OpReturnValue %44 OpFunctionEnd -%vertex_main = OpFunction %void None %41 - %42 = OpLabel - %43 = OpFunctionCall %VertexOutput %vertex_main_inner - %44 = OpCompositeExtract %v4float %43 0 - OpStore %vertex_main_position_Output %44 None - %45 = OpCompositeExtract %v4int %43 1 - OpStore %vertex_main_loc0_Output %45 None +%vertex_main = OpFunction %void None %47 + %48 = OpLabel + %49 = OpFunctionCall %VertexOutput %vertex_main_inner + %50 = OpCompositeExtract %v4float %49 0 + OpStore %vertex_main_position_Output %50 None + %51 = OpCompositeExtract %v4int %49 1 + OpStore %vertex_main_loc0_Output %51 None OpStore %vertex_main___point_size_Output %float_1 None OpReturn OpFunctionEnd