[tint] Basic scalarization support (max min clamp)
Various platforms seem to have issues with vectorized max,min, and
clamp builtins. For adreno this likely is due to a poor front end
implementation of handling nans.
For Intel (gen9) it is less clear where the implementation goes wrong
but what is clear is that scalarization resolves the issue.
Removed spirv only scalarization as it is duplicate. End2End test
still remains and passes.
Bug: 422144514, 407109052
Change-Id: Ie16617136be6e5993d5cf9c813eac2aa6b5492af
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/246095
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp
index 5103ae4..dd87d12 100644
--- a/src/dawn/native/Toggles.cpp
+++ b/src/dawn/native/Toggles.cpp
@@ -570,6 +570,9 @@
{"disable_polyfills_on_integer_div_and_mod",
"Disable the Tint polyfills on integer division and modulo.", "https://crbug.com/tint/2128",
ToggleStage::Device}},
+ {Toggle::ScalarizeMaxMinClamp,
+ {"scalarize_max_min_clamp", "Scalarize max, min, and clamp builtins.",
+ "https://crbug.com/422144514", ToggleStage::Device}},
{Toggle::MetalEnableModuleConstant,
{"metal_enable_module_constant_transform", "Enable the module constant transform.",
"https://crbug.com/419804339", ToggleStage::Device}},
@@ -642,9 +645,6 @@
{Toggle::UseVulkanMemoryModel,
{"use_vulkan_memory_model", "Use the Vulkan Memory Model if available.",
"https://crbug.com/392606604", ToggleStage::Adapter}},
- {Toggle::VulkanScalarizeClampBuiltin,
- {"vulkan_scalarize_clamp_builtin", "Scalarize calls to the clamp builtin.",
- "https://crbug.com/407109052", ToggleStage::Device}},
{Toggle::VulkanDirectVariableAccessTransformHandle,
{"vulkan_direct_variable_access_transform_handle",
"Transform handles using direct variable access.", "https://crbug.com/387000529",
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h
index 0c7f919..e939459 100644
--- a/src/dawn/native/Toggles.h
+++ b/src/dawn/native/Toggles.h
@@ -139,6 +139,7 @@
ExposeWGSLTestingFeatures,
ExposeWGSLExperimentalFeatures,
DisablePolyfillsOnIntegerDivisonAndModulo,
+ ScalarizeMaxMinClamp,
MetalEnableModuleConstant,
EnableImmediateErrorHandling,
VulkanUseStorageInputOutput16,
@@ -153,7 +154,6 @@
D3D12RelaxMinSubgroupSizeTo8,
D3D12RelaxBufferTextureCopyPitchAndOffsetAlignment,
UseVulkanMemoryModel,
- VulkanScalarizeClampBuiltin,
VulkanDirectVariableAccessTransformHandle,
VulkanAddWorkToEmptyResolvePass,
EnableIntegerRangeAnalysisInRobustness,
diff --git a/src/dawn/native/d3d11/ShaderModuleD3D11.cpp b/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
index d22d324..3da7465 100644
--- a/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
+++ b/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
@@ -234,6 +234,8 @@
req.hlsl.tintOptions.disable_workgroup_init =
device->IsToggleEnabled(Toggle::DisableWorkgroupInit);
req.hlsl.tintOptions.bindings = std::move(bindings);
+ req.hlsl.tintOptions.scalarize_max_min_clamp =
+ device->IsToggleEnabled(Toggle::ScalarizeMaxMinClamp);
// Immediate data available in TintIR only.
if (useTintIR) {
diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
index ac26afc..84a0b8d 100644
--- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -372,6 +372,8 @@
device->IsToggleEnabled(Toggle::PolyFillPacked4x8DotProduct);
req.hlsl.tintOptions.disable_polyfill_integer_div_mod =
device->IsToggleEnabled(Toggle::DisablePolyfillsOnIntegerDivisonAndModulo);
+ req.hlsl.tintOptions.scalarize_max_min_clamp =
+ device->IsToggleEnabled(Toggle::ScalarizeMaxMinClamp);
req.hlsl.tintOptions.polyfill_pack_unpack_4x8 =
device->IsToggleEnabled(Toggle::D3D12PolyFillPackUnpack4x8);
req.hlsl.tintOptions.enable_integer_range_analysis =
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm
index 7c46a13..ec35b39 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -295,6 +295,7 @@
req.tintOptions.bindings = std::move(bindings);
req.tintOptions.disable_polyfill_integer_div_mod =
device->IsToggleEnabled(Toggle::DisablePolyfillsOnIntegerDivisonAndModulo);
+ req.tintOptions.scalarize_max_min_clamp = device->IsToggleEnabled(Toggle::ScalarizeMaxMinClamp);
req.tintOptions.enable_module_constant =
device->IsToggleEnabled(Toggle::MetalEnableModuleConstant);
req.tintOptions.vertex_pulling_config = std::move(vertexPullingTransformConfig);
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
index 48f3c46..3b5d0f8 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -816,7 +816,7 @@
// chromium:407109052: Qualcomm devices have a bug where the spirv extended op NClamp
// modifies other components of a vector when one of the components is nan.
- deviceToggles->Default(Toggle::VulkanScalarizeClampBuiltin, true);
+ deviceToggles->Default(Toggle::ScalarizeMaxMinClamp, true);
}
if (IsAndroidARM()) {
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index da62f09..ba734ad 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -279,10 +279,10 @@
GetDevice()->IsToggleEnabled(Toggle::PolyfillPackUnpack4x8Norm);
req.tintOptions.disable_polyfill_integer_div_mod =
GetDevice()->IsToggleEnabled(Toggle::DisablePolyfillsOnIntegerDivisonAndModulo);
+ req.tintOptions.scalarize_max_min_clamp =
+ GetDevice()->IsToggleEnabled(Toggle::ScalarizeMaxMinClamp);
req.tintOptions.use_vulkan_memory_model =
GetDevice()->IsToggleEnabled(Toggle::UseVulkanMemoryModel);
- req.tintOptions.scalarize_clamp_builtin =
- GetDevice()->IsToggleEnabled(Toggle::VulkanScalarizeClampBuiltin);
req.tintOptions.dva_transform_handle =
GetDevice()->IsToggleEnabled(Toggle::VulkanDirectVariableAccessTransformHandle);
// Pass matrices to user functions by pointer on Qualcomm devices to workaround a known bug.
diff --git a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
index 89d66e4..7eb87f3 100644
--- a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
+++ b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
@@ -33,8 +33,6 @@
#include "dawn/tests/DawnTest.h"
#include "dawn/utils/WGPUHelpers.h"
-//
-
namespace dawn {
namespace {
@@ -69,10 +67,10 @@
}
};
-TEST_P(PolyfillBuiltinSimpleTests, ScalarizeClampBuiltin) {
+TEST_P(PolyfillBuiltinSimpleTests, ScalarizeClampBuiltinNanComponent) {
// Some devices (Adreno) do not handle nan's correctly for the clamp function
- // This test will fail on those devices without the builtin polyfill applied.
- // See: crbug.com/407109052
+ // This test will fail on those devices without the builtin polyfill/scalarize
+ // applied. See: crbug.com/407109052
std::string kShaderCode = R"(
@group(0) @binding(0) var<storage, read_write> in_out : array<u32, 2>;
@compute @workgroup_size(1)
@@ -106,11 +104,85 @@
EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
}
+TEST_P(PolyfillBuiltinSimpleTests, ScalarizeClampBuiltin) {
+ // Basic correctness test for scalariztion of clamp.
+ std::string kShaderCode = R"(
+ @group(0) @binding(0) var<storage, read_write> in_out : array<u32, 2>;
+ @compute @workgroup_size(1)
+ fn main() {
+ var x = vec2(5.0, -2.0);
+ var q = clamp(x, vec2(0.0), vec2(1.0));
+ in_out[0] = u32(q.x);
+ in_out[1] = u32(q.y);
+ }
+ )";
+
+ wgpu::ComputePipeline pipeline = CreateComputePipeline(kShaderCode);
+ uint32_t kDefaultVal = 0;
+ wgpu::Buffer output = CreateBuffer(2, kDefaultVal);
+ wgpu::BindGroup bindGroup =
+ utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, output}});
+
+ wgpu::CommandBuffer commands;
+ {
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetPipeline(pipeline);
+ pass.SetBindGroup(0, bindGroup);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+ commands = encoder.Finish();
+ }
+
+ queue.Submit(1, &commands);
+ std::vector<uint32_t> expected = {1, 0};
+ EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
+TEST_P(PolyfillBuiltinSimpleTests, ScalarizeMinMaxBuiltin) {
+ // Basic correctness test for scalariztion of min and max.
+ std::string kShaderCode = R"(
+ @group(0) @binding(0) var<storage, read_write> in_out : array<u32, 2>;
+ @compute @workgroup_size(1)
+ fn main() {
+ var x = vec2(5.0, -2.0);
+ var q = min(vec2(3.0), max(x, vec2(2.0)));
+ in_out[0] = u32(q.x);
+ in_out[1] = u32(q.y);
+ }
+ )";
+
+ wgpu::ComputePipeline pipeline = CreateComputePipeline(kShaderCode);
+ uint32_t kDefaultVal = 0;
+ wgpu::Buffer output = CreateBuffer(2, kDefaultVal);
+ wgpu::BindGroup bindGroup =
+ utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, output}});
+
+ wgpu::CommandBuffer commands;
+ {
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetPipeline(pipeline);
+ pass.SetBindGroup(0, bindGroup);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+ commands = encoder.Finish();
+ }
+
+ queue.Submit(1, &commands);
+ std::vector<uint32_t> expected = {3, 2};
+ EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
DAWN_INSTANTIATE_TEST(PolyfillBuiltinSimpleTests,
D3D12Backend(),
+ D3D11Backend(),
MetalBackend(),
VulkanBackend(),
- VulkanBackend({"vulkan_scalarize_clamp_builtin"}),
+ D3D12Backend({"scalarize_max_min_clamp"}),
+ MetalBackend({"scalarize_max_min_clamp"}),
+ VulkanBackend({"scalarize_max_min_clamp"}),
+ D3D11Backend({"scalarize_max_min_clamp"}),
OpenGLESBackend());
} // anonymous namespace
diff --git a/src/tint/lang/core/ir/transform/BUILD.bazel b/src/tint/lang/core/ir/transform/BUILD.bazel
index b801ad0..94705ce 100644
--- a/src/tint/lang/core/ir/transform/BUILD.bazel
+++ b/src/tint/lang/core/ir/transform/BUILD.bazel
@@ -46,6 +46,7 @@
"binding_remapper.cc",
"block_decorated_structs.cc",
"builtin_polyfill.cc",
+ "builtin_scalarize.cc",
"combine_access_instructions.cc",
"conversion_polyfill.cc",
"dead_code_elimination.cc",
@@ -76,6 +77,7 @@
"binding_remapper.h",
"block_decorated_structs.h",
"builtin_polyfill.h",
+ "builtin_scalarize.h",
"combine_access_instructions.h",
"conversion_polyfill.h",
"dead_code_elimination.h",
@@ -134,6 +136,7 @@
"binding_remapper_test.cc",
"block_decorated_structs_test.cc",
"builtin_polyfill_test.cc",
+ "builtin_scalarize_test.cc",
"combine_access_instructions_test.cc",
"conversion_polyfill_test.cc",
"dead_code_elimination_test.cc",
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake
index 3f2abf8..c0cb68a 100644
--- a/src/tint/lang/core/ir/transform/BUILD.cmake
+++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -53,6 +53,8 @@
lang/core/ir/transform/block_decorated_structs.h
lang/core/ir/transform/builtin_polyfill.cc
lang/core/ir/transform/builtin_polyfill.h
+ lang/core/ir/transform/builtin_scalarize.cc
+ lang/core/ir/transform/builtin_scalarize.h
lang/core/ir/transform/combine_access_instructions.cc
lang/core/ir/transform/combine_access_instructions.h
lang/core/ir/transform/conversion_polyfill.cc
@@ -135,6 +137,7 @@
lang/core/ir/transform/binding_remapper_test.cc
lang/core/ir/transform/block_decorated_structs_test.cc
lang/core/ir/transform/builtin_polyfill_test.cc
+ lang/core/ir/transform/builtin_scalarize_test.cc
lang/core/ir/transform/combine_access_instructions_test.cc
lang/core/ir/transform/conversion_polyfill_test.cc
lang/core/ir/transform/dead_code_elimination_test.cc
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn
index cdf8820..9e89f0a 100644
--- a/src/tint/lang/core/ir/transform/BUILD.gn
+++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -59,6 +59,8 @@
"block_decorated_structs.h",
"builtin_polyfill.cc",
"builtin_polyfill.h",
+ "builtin_scalarize.cc",
+ "builtin_scalarize.h",
"combine_access_instructions.cc",
"combine_access_instructions.h",
"conversion_polyfill.cc",
@@ -135,6 +137,7 @@
"binding_remapper_test.cc",
"block_decorated_structs_test.cc",
"builtin_polyfill_test.cc",
+ "builtin_scalarize_test.cc",
"combine_access_instructions_test.cc",
"conversion_polyfill_test.cc",
"dead_code_elimination_test.cc",
diff --git a/src/tint/lang/core/ir/transform/builtin_scalarize.cc b/src/tint/lang/core/ir/transform/builtin_scalarize.cc
new file mode 100644
index 0000000..f21b864
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/builtin_scalarize.cc
@@ -0,0 +1,148 @@
+// Copyright 2025 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/core/ir/transform/builtin_scalarize.h"
+#include <cstdint>
+#include <utility>
+
+#include "src/tint/lang/core/builtin_fn.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"
+#include "src/tint/lang/core/type/sampled_texture.h"
+#include "src/tint/lang/core/type/texture.h"
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+namespace tint::core::ir::transform {
+
+namespace {
+
+/// PIMPL state for the transform.
+struct State {
+ /// The polyfill config.
+ const BuiltinScalarizeConfig& config;
+
+ /// The IR module.
+ Module& ir;
+
+ /// The IR builder.
+ Builder b{ir};
+
+ /// The type manager.
+ core::type::Manager& ty{ir.Types()};
+
+ /// The symbol table.
+ SymbolTable& sym{ir.symbols};
+
+ // We cannot arbitrarily allow the config to scalarize any builtin as this might cause
+ // semantically incorrect scalarizations. An example here is 'cross' which is a vec3 for all
+ // input and return parameters but cannot be scalarized.
+ bool ShouldAttemptScalarize(ir::CoreBuiltinCall* builtin) {
+ auto builtin_enum = builtin->Func();
+ if (!builtin->Result()->Type()->Is<core::type::Vector>()) {
+ // No vector found. Already scalar.
+ return false;
+ }
+ switch (builtin_enum) {
+ case core::BuiltinFn::kClamp:
+ return config.scalarize_clamp;
+ case core::BuiltinFn::kMax:
+ return config.scalarize_max;
+ case core::BuiltinFn::kMin:
+ return config.scalarize_min;
+ default:
+ return false;
+ }
+ }
+
+ void Process() {
+ Vector<ir::CoreBuiltinCall*, 4> worklist;
+ for (auto* inst : ir.Instructions()) {
+ if (auto* builtin = inst->As<ir::CoreBuiltinCall>()) {
+ if (ShouldAttemptScalarize(builtin)) {
+ worklist.Push(builtin);
+ }
+ }
+ }
+
+ for (auto* builtin : worklist) {
+ ScalarizeBuiltin(builtin);
+ }
+ }
+
+ void ScalarizeBuiltin(ir::CoreBuiltinCall* builtin) {
+ uint32_t common_vec_width = builtin->Result()->Type()->As<core::type::Vector>()->Width();
+ b.InsertBefore(builtin, [&] {
+ const core::type::Type* scalar_return_type =
+ builtin->Result()->Type()->DeepestElement();
+ Vector<core::ir::Value*, 4> args;
+ for (uint32_t i = 0; i < common_vec_width; i++) {
+ Vector<core::ir::Value*, 4> scalar_args;
+
+ for (auto& e : builtin->Args()) {
+ if (auto* vec = e->Type()->As<core::type::Vector>()) {
+ // It would be an error to scalarize over different sized vectors.
+ TINT_ASSERT(common_vec_width == vec->Width());
+ auto* access_arg = b.Access(vec->DeepestElement(), e, u32(i));
+ scalar_args.Push(access_arg->Result());
+ } else {
+ TINT_ASSERT(e->Type()->IsScalar());
+ // This code generalizes for vector functions that additionally take scalar
+ // inputs. And example of this is the second and third parameters of
+ // 'extract_bits'.
+ scalar_args.Push(e);
+ }
+ }
+
+ auto* scalar_call =
+ b.Call(scalar_return_type, builtin->Func(), std::move(scalar_args));
+ args.Push(scalar_call->Result());
+ }
+ // Places result back into a vector.
+ b.ConstructWithResult(builtin->DetachResult(), std::move(args));
+ });
+ builtin->Destroy();
+ }
+};
+
+} // namespace
+
+Result<SuccessType> BuiltinScalarize(Module& ir, const BuiltinScalarizeConfig& config) {
+ auto result =
+ ValidateAndDumpIfNeeded(ir, "core.BuiltinScalarize", kBuiltinScalarizeCapabilities);
+ if (result != Success) {
+ return result;
+ }
+
+ State{config, ir}.Process();
+
+ return Success;
+}
+
+} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/core/ir/transform/builtin_scalarize.h b/src/tint/lang/core/ir/transform/builtin_scalarize.h
new file mode 100644
index 0000000..cfc9c7a
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/builtin_scalarize.h
@@ -0,0 +1,82 @@
+// Copyright 2025 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_CORE_IR_TRANSFORM_BUILTIN_SCALARIZE_H_
+#define SRC_TINT_LANG_CORE_IR_TRANSFORM_BUILTIN_SCALARIZE_H_
+
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/utils/reflection.h"
+#include "src/tint/utils/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::core::ir::transform {
+
+/// The capabilities that the transform can support.
+const Capabilities kBuiltinScalarizeCapabilities{
+ core::ir::Capability::kAllowDuplicateBindings,
+ core::ir::Capability::kAllow8BitIntegers,
+ core::ir::Capability::kAllow64BitIntegers,
+ core::ir::Capability::kAllowPointersAndHandlesInStructures,
+ core::ir::Capability::kAllowVectorElementPointer,
+ core::ir::Capability::kAllowHandleVarsWithoutBindings,
+ core::ir::Capability::kAllowClipDistancesOnF32,
+ core::ir::Capability::kAllowPrivateVarsInFunctions,
+ core::ir::Capability::kAllowAnyLetType,
+ core::ir::Capability::kAllowWorkspacePointerInputToEntryPoint,
+ core::ir::Capability::kAllowModuleScopeLets,
+ core::ir::Capability::kAllowAnyInputAttachmentIndexType,
+};
+
+/// The scalarizer configuration options
+struct BuiltinScalarizeConfig {
+ // Set to true to scalarize clamp builtin
+ bool scalarize_clamp = false;
+
+ // Set to true to scalarize max builtin
+ bool scalarize_max = false;
+
+ // Set to true to scalarize min builtin
+ bool scalarize_min = false;
+
+ /// Reflection for this class
+ TINT_REFLECT(BuiltinScalarizeConfig, scalarize_clamp, scalarize_max, scalarize_min);
+};
+
+/// BuiltinScalarize is a transform that replaces calls to builtin vector functions with scalar
+/// equivalent alternatives.
+/// @param module the module to transform
+/// @param config the polyfill configuration
+/// @returns success or failure
+Result<SuccessType> BuiltinScalarize(Module& module, const BuiltinScalarizeConfig& config);
+
+} // namespace tint::core::ir::transform
+
+#endif // SRC_TINT_LANG_CORE_IR_TRANSFORM_BUILTIN_SCALARIZE_H_
diff --git a/src/tint/lang/core/ir/transform/builtin_scalarize_test.cc b/src/tint/lang/core/ir/transform/builtin_scalarize_test.cc
new file mode 100644
index 0000000..00a0b76
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/builtin_scalarize_test.cc
@@ -0,0 +1,330 @@
+// Copyright 2025 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/core/ir/transform/builtin_scalarize.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/builtin_fn.h"
+#include "src/tint/lang/core/ir/constant.h"
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/core/type/sampled_texture.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using IR_BuiltinScalarizeTest = TransformTest;
+
+TEST_F(IR_BuiltinScalarizeTest, Clamp_VectorOperands_Scalarize) {
+ auto* x = b.FunctionParam("x", ty.vec2<f32>());
+ auto* low = b.FunctionParam("low", ty.vec2<f32>());
+ auto* high = b.FunctionParam("high", ty.vec2<f32>());
+ auto* func = b.Function("foo", ty.vec2<f32>());
+ func->SetParams({x, low, high});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.vec2<f32>(), core::BuiltinFn::kClamp, x, low, high);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%x:vec2<f32>, %low:vec2<f32>, %high:vec2<f32>):vec2<f32> {
+ $B1: {
+ %5:vec2<f32> = clamp %x, %low, %high
+ ret %5
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%x:vec2<f32>, %low:vec2<f32>, %high:vec2<f32>):vec2<f32> {
+ $B1: {
+ %5:f32 = access %x, 0u
+ %6:f32 = access %low, 0u
+ %7:f32 = access %high, 0u
+ %8:f32 = clamp %5, %6, %7
+ %9:f32 = access %x, 1u
+ %10:f32 = access %low, 1u
+ %11:f32 = access %high, 1u
+ %12:f32 = clamp %9, %10, %11
+ %13:vec2<f32> = construct %8, %12
+ ret %13
+ }
+}
+)";
+
+ BuiltinScalarizeConfig config{.scalarize_clamp = true};
+ Run(BuiltinScalarize, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinScalarizeTest, Clamp_VectorOperands_AlreadyScalarize) {
+ auto* x = b.FunctionParam("x", ty.f32());
+ auto* low = b.FunctionParam("low", ty.f32());
+ auto* high = b.FunctionParam("high", ty.f32());
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({x, low, high});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.f32(), core::BuiltinFn::kClamp, x, low, high);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%x:f32, %low:f32, %high:f32):f32 {
+ $B1: {
+ %5:f32 = clamp %x, %low, %high
+ ret %5
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%x:f32, %low:f32, %high:f32):f32 {
+ $B1: {
+ %5:f32 = clamp %x, %low, %high
+ ret %5
+ }
+}
+)";
+
+ BuiltinScalarizeConfig config{.scalarize_clamp = true};
+ Run(BuiltinScalarize, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinScalarizeTest, Clamp_VectorOperands_LeaveVectorized) {
+ auto* x = b.FunctionParam("x", ty.vec2<f32>());
+ auto* low = b.FunctionParam("low", ty.vec2<f32>());
+ auto* high = b.FunctionParam("high", ty.vec2<f32>());
+ auto* func = b.Function("foo", ty.vec2<f32>());
+ func->SetParams({x, low, high});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.vec2<f32>(), core::BuiltinFn::kClamp, x, low, high);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%x:vec2<f32>, %low:vec2<f32>, %high:vec2<f32>):vec2<f32> {
+ $B1: {
+ %5:vec2<f32> = clamp %x, %low, %high
+ ret %5
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%x:vec2<f32>, %low:vec2<f32>, %high:vec2<f32>):vec2<f32> {
+ $B1: {
+ %5:vec2<f32> = clamp %x, %low, %high
+ ret %5
+ }
+}
+)";
+
+ BuiltinScalarizeConfig config;
+ Run(BuiltinScalarize, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinScalarizeTest, Max_VectorOperands_Scalarize) {
+ auto* x = b.FunctionParam("x", ty.vec3<f32>());
+ auto* y = b.FunctionParam("y", ty.vec3<f32>());
+ auto* func = b.Function("foo", ty.vec3<f32>());
+ func->SetParams({x, y});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.vec3<f32>(), core::BuiltinFn::kMax, x, y);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%x:vec3<f32>, %y:vec3<f32>):vec3<f32> {
+ $B1: {
+ %4:vec3<f32> = max %x, %y
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%x:vec3<f32>, %y:vec3<f32>):vec3<f32> {
+ $B1: {
+ %4:f32 = access %x, 0u
+ %5:f32 = access %y, 0u
+ %6:f32 = max %4, %5
+ %7:f32 = access %x, 1u
+ %8:f32 = access %y, 1u
+ %9:f32 = max %7, %8
+ %10:f32 = access %x, 2u
+ %11:f32 = access %y, 2u
+ %12:f32 = max %10, %11
+ %13:vec3<f32> = construct %6, %9, %12
+ ret %13
+ }
+}
+)";
+
+ BuiltinScalarizeConfig config{.scalarize_max = true};
+ Run(BuiltinScalarize, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinScalarizeTest, Max_VectorOperands_LeaveVectorized) {
+ auto* x = b.FunctionParam("x", ty.vec3<f32>());
+ auto* y = b.FunctionParam("y", ty.vec3<f32>());
+ auto* func = b.Function("foo", ty.vec3<f32>());
+ func->SetParams({x, y});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.vec3<f32>(), core::BuiltinFn::kMax, x, y);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%x:vec3<f32>, %y:vec3<f32>):vec3<f32> {
+ $B1: {
+ %4:vec3<f32> = max %x, %y
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%x:vec3<f32>, %y:vec3<f32>):vec3<f32> {
+ $B1: {
+ %4:vec3<f32> = max %x, %y
+ ret %4
+ }
+}
+)";
+
+ BuiltinScalarizeConfig config;
+ Run(BuiltinScalarize, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinScalarizeTest, Min_VectorOperands_Scalarize) {
+ auto* x = b.FunctionParam("x", ty.vec4<f16>());
+ auto* y = b.FunctionParam("y", ty.vec4<f16>());
+ auto* func = b.Function("foo", ty.vec4<f16>());
+ func->SetParams({x, y});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.vec4<f16>(), core::BuiltinFn::kMin, x, y);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%x:vec4<f16>, %y:vec4<f16>):vec4<f16> {
+ $B1: {
+ %4:vec4<f16> = min %x, %y
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%x:vec4<f16>, %y:vec4<f16>):vec4<f16> {
+ $B1: {
+ %4:f16 = access %x, 0u
+ %5:f16 = access %y, 0u
+ %6:f16 = min %4, %5
+ %7:f16 = access %x, 1u
+ %8:f16 = access %y, 1u
+ %9:f16 = min %7, %8
+ %10:f16 = access %x, 2u
+ %11:f16 = access %y, 2u
+ %12:f16 = min %10, %11
+ %13:f16 = access %x, 3u
+ %14:f16 = access %y, 3u
+ %15:f16 = min %13, %14
+ %16:vec4<f16> = construct %6, %9, %12, %15
+ ret %16
+ }
+}
+)";
+
+ BuiltinScalarizeConfig config{.scalarize_min = true};
+ Run(BuiltinScalarize, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinScalarizeTest, Min_VectorOperands_LeaveVectorized) {
+ auto* x = b.FunctionParam("x", ty.vec4<f16>());
+ auto* y = b.FunctionParam("y", ty.vec4<f16>());
+ auto* func = b.Function("foo", ty.vec4<f16>());
+ func->SetParams({x, y});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.vec4<f16>(), core::BuiltinFn::kMin, x, y);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%x:vec4<f16>, %y:vec4<f16>):vec4<f16> {
+ $B1: {
+ %4:vec4<f16> = min %x, %y
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%x:vec4<f16>, %y:vec4<f16>):vec4<f16> {
+ $B1: {
+ %4:vec4<f16> = min %x, %y
+ ret %4
+ }
+}
+)";
+
+ BuiltinScalarizeConfig config;
+ Run(BuiltinScalarize, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/hlsl/writer/common/options.h b/src/tint/lang/hlsl/writer/common/options.h
index e41f6cf..4c77d67 100644
--- a/src/tint/lang/hlsl/writer/common/options.h
+++ b/src/tint/lang/hlsl/writer/common/options.h
@@ -233,6 +233,9 @@
/// Set to `true` to disable the polyfills on integer division and modulo.
bool disable_polyfill_integer_div_mod = false;
+ /// Set to `true` to scalarize max, min, and clamp builtins.
+ bool scalarize_max_min_clamp = false;
+
/// Set to `true` to generate polyfill for `pack4xI8`, `pack4xU8`, `pack4xI8Clamp`,
/// `unpack4xI8` and `unpack4xU8` builtins
bool polyfill_pack_unpack_4x8 = false;
@@ -282,6 +285,7 @@
polyfill_reflect_vec2_f32,
polyfill_dot_4x8_packed,
disable_polyfill_integer_div_mod,
+ scalarize_max_min_clamp,
polyfill_pack_unpack_4x8,
compiler,
array_length_from_uniform,
diff --git a/src/tint/lang/hlsl/writer/raise/raise.cc b/src/tint/lang/hlsl/writer/raise/raise.cc
index c4d7c3b..773b938 100644
--- a/src/tint/lang/hlsl/writer/raise/raise.cc
+++ b/src/tint/lang/hlsl/writer/raise/raise.cc
@@ -36,6 +36,7 @@
#include "src/tint/lang/core/ir/transform/binary_polyfill.h"
#include "src/tint/lang/core/ir/transform/binding_remapper.h"
#include "src/tint/lang/core/ir/transform/builtin_polyfill.h"
+#include "src/tint/lang/core/ir/transform/builtin_scalarize.h"
#include "src/tint/lang/core/ir/transform/conversion_polyfill.h"
#include "src/tint/lang/core/ir/transform/demote_to_helper.h"
#include "src/tint/lang/core/ir/transform/direct_variable_access.h"
@@ -253,6 +254,12 @@
RUN_TRANSFORM(core::ir::transform::VectorizeScalarMatrixConstructors, module);
RUN_TRANSFORM(core::ir::transform::RemoveContinueInSwitch, module);
+ core::ir::transform::BuiltinScalarizeConfig scalarize_config{
+ .scalarize_clamp = options.scalarize_max_min_clamp,
+ .scalarize_max = options.scalarize_max_min_clamp,
+ .scalarize_min = options.scalarize_max_min_clamp};
+ RUN_TRANSFORM(core::ir::transform::BuiltinScalarize, module, scalarize_config);
+
// These transforms need to be run last as various transforms introduce terminator arguments,
// naming conflicts, and expressions that need to be explicitly not inlined.
RUN_TRANSFORM(core::ir::transform::RemoveTerminatorArgs, module);
diff --git a/src/tint/lang/msl/writer/common/options.h b/src/tint/lang/msl/writer/common/options.h
index 7dac767..925dcfb 100644
--- a/src/tint/lang/msl/writer/common/options.h
+++ b/src/tint/lang/msl/writer/common/options.h
@@ -163,6 +163,9 @@
/// Set to `true` to disable the polyfills on integer division and modulo.
bool disable_polyfill_integer_div_mod = false;
+ /// Set to `true` to scalarize max min and clamp builtins.
+ bool scalarize_max_min_clamp = false;
+
/// Set to `true` to enable the module constant transform
bool enable_module_constant = false;
@@ -200,6 +203,7 @@
disable_demote_to_helper,
emit_vertex_point_size,
disable_polyfill_integer_div_mod,
+ scalarize_max_min_clamp,
enable_module_constant,
use_argument_buffers,
buffer_size_ubo_index,
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index 5a0373f..2a59e32 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -32,6 +32,7 @@
#include "src/tint/lang/core/ir/transform/binary_polyfill.h"
#include "src/tint/lang/core/ir/transform/binding_remapper.h"
#include "src/tint/lang/core/ir/transform/builtin_polyfill.h"
+#include "src/tint/lang/core/ir/transform/builtin_scalarize.h"
#include "src/tint/lang/core/ir/transform/conversion_polyfill.h"
#include "src/tint/lang/core/ir/transform/demote_to_helper.h"
#include "src/tint/lang/core/ir/transform/multiplanar_external_texture.h"
@@ -160,6 +161,13 @@
RUN_TRANSFORM(raise::BinaryPolyfill, module);
RUN_TRANSFORM(raise::BuiltinPolyfill, module);
+ core::ir::transform::BuiltinScalarizeConfig scalarize_config{
+ .scalarize_clamp = options.scalarize_max_min_clamp,
+ .scalarize_max = options.scalarize_max_min_clamp,
+ .scalarize_min = options.scalarize_max_min_clamp,
+ };
+ RUN_TRANSFORM(core::ir::transform::BuiltinScalarize, module, scalarize_config);
+
if (options.enable_module_constant) {
RUN_TRANSFORM(raise::ModuleConstant, module);
}
diff --git a/src/tint/lang/spirv/writer/common/options.h b/src/tint/lang/spirv/writer/common/options.h
index d29bfe5..23d6dfd 100644
--- a/src/tint/lang/spirv/writer/common/options.h
+++ b/src/tint/lang/spirv/writer/common/options.h
@@ -205,12 +205,12 @@
/// Set to `true` to disable the polyfills on integer division and modulo.
bool disable_polyfill_integer_div_mod = false;
+ /// Set to `true` to scalarize max min and clamp builtins.
+ bool scalarize_max_min_clamp = false;
+
/// Set to `true` if the Vulkan Memory Model should be used
bool use_vulkan_memory_model = false;
- /// Set to `true` if the clamp builtin should be scalarized for vector operations
- bool scalarize_clamp_builtin = false;
-
/// Set to `true` if handles should be transformed by direct variable access.
bool dva_transform_handle = false;
@@ -239,8 +239,8 @@
polyfill_dot_4x8_packed,
polyfill_pack_unpack_4x8_norm,
disable_polyfill_integer_div_mod,
+ scalarize_max_min_clamp,
use_vulkan_memory_model,
- scalarize_clamp_builtin,
dva_transform_handle,
depth_range_offsets,
spirv_version);
diff --git a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
index a30117d..08bf6a2 100644
--- a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
@@ -231,12 +231,6 @@
worklist.Push(builtin);
}
break;
- case core::BuiltinFn::kClamp:
- if (config.scalarize_clamp_builtin &&
- builtin->Result()->Type()->Is<core::type::Vector>()) {
- worklist.Push(builtin);
- }
- break;
default:
break;
}
@@ -275,9 +269,6 @@
case core::BuiltinFn::kSelect:
Select(builtin);
break;
- case core::BuiltinFn::kClamp:
- Clamp(builtin);
- break;
case core::BuiltinFn::kSubgroupBroadcast:
SubgroupBroadcast(builtin);
break;
@@ -558,34 +549,6 @@
builtin->Destroy();
}
- /// Handle a `clamp()` builtin when scalarization is required.
- /// @param builtin the builtin call instruction
- void Clamp(core::ir::CoreBuiltinCall* builtin) {
- auto* e = builtin->Args()[0];
- auto* vec = e->Type()->As<core::type::Vector>();
- if (!vec) {
- // Already is a scalar. No change required.
- return;
- }
-
- b.InsertBefore(builtin, [&] {
- auto* low = builtin->Args()[1];
- auto* high = builtin->Args()[2];
- auto* type = vec->DeepestElement();
- Vector<core::ir::Value*, 4> args;
- for (uint32_t i = 0; i < vec->Width(); i++) {
- auto* access_e = b.Access(type, e, u32(i));
- auto* access_low = b.Access(type, low, u32(i));
- auto* access_high = b.Access(type, high, u32(i));
- auto* scalar_call =
- b.Call(type, core::BuiltinFn::kClamp, access_e, access_low, access_high);
- args.Push(scalar_call->Result());
- }
- b.ConstructWithResult(builtin->DetachResult(), std::move(args));
- });
- builtin->Destroy();
- }
-
/// ImageOperands represents the optional image operands for an image instruction.
struct ImageOperands {
/// Bias
diff --git a/src/tint/lang/spirv/writer/raise/builtin_polyfill.h b/src/tint/lang/spirv/writer/raise/builtin_polyfill.h
index 60ea101..98b063d 100644
--- a/src/tint/lang/spirv/writer/raise/builtin_polyfill.h
+++ b/src/tint/lang/spirv/writer/raise/builtin_polyfill.h
@@ -41,7 +41,6 @@
struct PolyfillConfig {
bool use_vulkan_memory_model = false;
- bool scalarize_clamp_builtin = false;
SpvVersion version = SpvVersion::kSpv13;
};
diff --git a/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc
index f01bb01..f881bd1 100644
--- a/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc
@@ -1174,133 +1174,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvWriter_BuiltinPolyfillTest, Clamp_VectorOperands_Scalarize) {
- auto* x = b.FunctionParam("x", ty.vec2<f32>());
- auto* low = b.FunctionParam("low", ty.vec2<f32>());
- auto* high = b.FunctionParam("high", ty.vec2<f32>());
- auto* func = b.Function("foo", ty.vec2<f32>());
- func->SetParams({x, low, high});
-
- b.Append(func->Block(), [&] {
- auto* result = b.Call(ty.vec2<f32>(), core::BuiltinFn::kClamp, x, low, high);
- b.Return(func, result);
- });
-
- auto* src = R"(
-%foo = func(%x:vec2<f32>, %low:vec2<f32>, %high:vec2<f32>):vec2<f32> {
- $B1: {
- %5:vec2<f32> = clamp %x, %low, %high
- ret %5
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-%foo = func(%x:vec2<f32>, %low:vec2<f32>, %high:vec2<f32>):vec2<f32> {
- $B1: {
- %5:f32 = access %x, 0u
- %6:f32 = access %low, 0u
- %7:f32 = access %high, 0u
- %8:f32 = clamp %5, %6, %7
- %9:f32 = access %x, 1u
- %10:f32 = access %low, 1u
- %11:f32 = access %high, 1u
- %12:f32 = clamp %9, %10, %11
- %13:vec2<f32> = construct %8, %12
- ret %13
- }
-}
-)";
-
- PolyfillConfig config{.scalarize_clamp_builtin = true};
- Run(BuiltinPolyfill, config);
-
- EXPECT_EQ(expect, str());
-}
-
-TEST_F(SpirvWriter_BuiltinPolyfillTest, Clamp_VectorOperands_Scalarize_f16) {
- auto* x = b.FunctionParam("x", ty.vec2<f16>());
- auto* low = b.FunctionParam("low", ty.vec2<f16>());
- auto* high = b.FunctionParam("high", ty.vec2<f16>());
- auto* func = b.Function("foo", ty.vec2<f16>());
- func->SetParams({x, low, high});
-
- b.Append(func->Block(), [&] {
- auto* result = b.Call(ty.vec2<f16>(), core::BuiltinFn::kClamp, x, low, high);
- b.Return(func, result);
- });
-
- auto* src = R"(
-%foo = func(%x:vec2<f16>, %low:vec2<f16>, %high:vec2<f16>):vec2<f16> {
- $B1: {
- %5:vec2<f16> = clamp %x, %low, %high
- ret %5
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-%foo = func(%x:vec2<f16>, %low:vec2<f16>, %high:vec2<f16>):vec2<f16> {
- $B1: {
- %5:f16 = access %x, 0u
- %6:f16 = access %low, 0u
- %7:f16 = access %high, 0u
- %8:f16 = clamp %5, %6, %7
- %9:f16 = access %x, 1u
- %10:f16 = access %low, 1u
- %11:f16 = access %high, 1u
- %12:f16 = clamp %9, %10, %11
- %13:vec2<f16> = construct %8, %12
- ret %13
- }
-}
-)";
-
- PolyfillConfig config{.scalarize_clamp_builtin = true};
- Run(BuiltinPolyfill, config);
-
- EXPECT_EQ(expect, str());
-}
-
-TEST_F(SpirvWriter_BuiltinPolyfillTest, Clamp_VectorOperands_Scalarize_AlreadyScalar) {
- auto* x = b.FunctionParam("x", ty.f32());
- auto* low = b.FunctionParam("low", ty.f32());
- auto* high = b.FunctionParam("high", ty.f32());
- auto* func = b.Function("foo", ty.f32());
- func->SetParams({x, low, high});
-
- b.Append(func->Block(), [&] {
- auto* result = b.Call(ty.f32(), core::BuiltinFn::kClamp, x, low, high);
- b.Return(func, result);
- });
-
- auto* src = R"(
-%foo = func(%x:f32, %low:f32, %high:f32):f32 {
- $B1: {
- %5:f32 = clamp %x, %low, %high
- ret %5
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-%foo = func(%x:f32, %low:f32, %high:f32):f32 {
- $B1: {
- %5:f32 = clamp %x, %low, %high
- ret %5
- }
-}
-)";
-
- PolyfillConfig config{.scalarize_clamp_builtin = true};
- Run(BuiltinPolyfill, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(SpirvWriter_BuiltinPolyfillTest, Clamp_VectorOperands_DisabledScalarize) {
auto* x = b.FunctionParam("x", ty.vec2<f32>());
auto* low = b.FunctionParam("low", ty.vec2<f32>());
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc
index c9c6fbd..85fcd0e 100644
--- a/src/tint/lang/spirv/writer/raise/raise.cc
+++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -34,6 +34,7 @@
#include "src/tint/lang/core/ir/transform/binding_remapper.h"
#include "src/tint/lang/core/ir/transform/block_decorated_structs.h"
#include "src/tint/lang/core/ir/transform/builtin_polyfill.h"
+#include "src/tint/lang/core/ir/transform/builtin_scalarize.h"
#include "src/tint/lang/core/ir/transform/combine_access_instructions.h"
#include "src/tint/lang/core/ir/transform/conversion_polyfill.h"
#include "src/tint/lang/core/ir/transform/demote_to_helper.h"
@@ -168,10 +169,16 @@
}
raise::PolyfillConfig config = {.use_vulkan_memory_model = options.use_vulkan_memory_model,
- .scalarize_clamp_builtin = options.scalarize_clamp_builtin,
.version = options.spirv_version};
RUN_TRANSFORM(raise::BuiltinPolyfill, module, config);
RUN_TRANSFORM(raise::ExpandImplicitSplats, module);
+
+ core::ir::transform::BuiltinScalarizeConfig scalarize_config{
+ .scalarize_clamp = options.scalarize_max_min_clamp,
+ .scalarize_max = options.scalarize_max_min_clamp,
+ .scalarize_min = options.scalarize_max_min_clamp};
+ RUN_TRANSFORM(core::ir::transform::BuiltinScalarize, module, scalarize_config);
+
// kAllowAnyInputAttachmentIndexType required after ExpandImplicitSplats
RUN_TRANSFORM(raise::HandleMatrixArithmetic, module);
RUN_TRANSFORM(raise::MergeReturn, module);
