[dawn] Polyfill vector clamp via scalarization (part 2)
See original:
https://dawn-review.googlesource.com/c/dawn/+/240495
This change enables and tests the polyfill.
The change was first tested with the polyfill disabled to prove that
the test will indeed fail on adreno bots.
Bug:407109052
Change-Id: Iaec8cdf2c3688f0305571473356f870a7d8a5f2a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/240894
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
index ee8945c..6881162 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -812,6 +812,10 @@
// resolve target doesn't perform the resolve. To work around it, add a small amount of work
// to the pass to force it to execute.
deviceToggles->Default(Toggle::VulkanAddWorkToEmptyResolvePass, true);
+
+ // 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);
}
if (IsAndroidARM()) {
diff --git a/src/dawn/tests/BUILD.gn b/src/dawn/tests/BUILD.gn
index 9cde65f..f56e9ec 100644
--- a/src/dawn/tests/BUILD.gn
+++ b/src/dawn/tests/BUILD.gn
@@ -656,6 +656,7 @@
"end2end/PipelineCachingTests.cpp",
"end2end/PipelineLayoutTests.cpp",
"end2end/PixelLocalStorageTests.cpp",
+ "end2end/PolyfillBuiltinSimpleTests.cpp",
"end2end/PrimitiveStateTests.cpp",
"end2end/PrimitiveTopologyTests.cpp",
"end2end/QueryTests.cpp",
diff --git a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
new file mode 100644
index 0000000..89d66e4
--- /dev/null
+++ b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
@@ -0,0 +1,117 @@
+// 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 <cstdint>
+#include <numeric>
+#include <string>
+#include <vector>
+
+#include "dawn/tests/DawnTest.h"
+#include "dawn/utils/WGPUHelpers.h"
+
+//
+
+namespace dawn {
+namespace {
+
+class PolyfillBuiltinSimpleTests : public DawnTest {
+ public:
+ wgpu::Buffer CreateBuffer(const std::vector<uint32_t>& data,
+ wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
+ wgpu::BufferUsage::CopySrc) {
+ uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(uint32_t));
+ return utils::CreateBufferFromData(device, data.data(), bufferSize, usage);
+ }
+
+ wgpu::Buffer CreateBuffer(const uint32_t count,
+ const uint32_t default_val = 0,
+ wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
+ wgpu::BufferUsage::CopySrc) {
+ return CreateBuffer(std::vector<uint32_t>(count, default_val), usage);
+ }
+
+ wgpu::ComputePipeline CreateComputePipeline(
+ const std::string& shader,
+ const char* entryPoint = nullptr,
+ const std::vector<wgpu::ConstantEntry>* constants = nullptr) {
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
+ csDesc.compute.entryPoint = entryPoint;
+ if (constants) {
+ csDesc.compute.constants = constants->data();
+ csDesc.compute.constantCount = constants->size();
+ }
+ return device.CreateComputePipeline(&csDesc);
+ }
+};
+
+TEST_P(PolyfillBuiltinSimpleTests, ScalarizeClampBuiltin) {
+ // 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
+ std::string kShaderCode = R"(
+ @group(0) @binding(0) var<storage, read_write> in_out : array<u32, 2>;
+ @compute @workgroup_size(1)
+ fn main() {
+ var zero = f32(in_out[0]);
+ var x = vec2(0.0/zero, 1.0);
+ var q = clamp(x, vec2(0.0), vec2(1.0));
+ 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 = {0, 1};
+ EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
+DAWN_INSTANTIATE_TEST(PolyfillBuiltinSimpleTests,
+ D3D12Backend(),
+ MetalBackend(),
+ VulkanBackend(),
+ VulkanBackend({"vulkan_scalarize_clamp_builtin"}),
+ OpenGLESBackend());
+
+} // anonymous namespace
+} // namespace dawn
