diff --git a/src/dawn/tests/BUILD.gn b/src/dawn/tests/BUILD.gn
index 95566b5..e65daa9 100644
--- a/src/dawn/tests/BUILD.gn
+++ b/src/dawn/tests/BUILD.gn
@@ -657,6 +657,7 @@
     "end2end/SamplerFilterAnisotropicTests.cpp",
     "end2end/SamplerTests.cpp",
     "end2end/ScissorTests.cpp",
+    "end2end/ShaderAtomicTests.cpp",
     "end2end/ShaderBuiltinPartialConstArgsErrorTests.cpp",
     "end2end/ShaderF16Tests.cpp",
     "end2end/ShaderTests.cpp",
diff --git a/src/dawn/tests/end2end/ShaderAtomicTests.cpp b/src/dawn/tests/end2end/ShaderAtomicTests.cpp
new file mode 100644
index 0000000..65b7fdf
--- /dev/null
+++ b/src/dawn/tests/end2end/ShaderAtomicTests.cpp
@@ -0,0 +1,208 @@
+// Copyright 2024 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/common/GPUInfo.h"
+#include "dawn/common/GPUInfo_autogen.h"
+#include "dawn/tests/DawnTest.h"
+#include "dawn/utils/ComboRenderPipelineDescriptor.h"
+#include "dawn/utils/WGPUHelpers.h"
+
+// The motivation behind these tests is to investigate the failures in
+// webgpu:shader,execution,expression,call,builtin,atomics,atomic* CTS for mobile gpus. The CAS
+// emulation workaround is enough to avoid the issue on mali but still shows up as an issue on
+// adreno.
+
+namespace dawn {
+namespace {
+
+enum class ShaderAtomicOp : uint8_t {
+    AtomicAdd,
+    AtomicCASFakeAdd,
+};
+
+std::ostream& operator<<(std::ostream& o, ShaderAtomicOp shader_op) {
+    switch (shader_op) {
+        case ShaderAtomicOp::AtomicAdd:
+            o << "AtomicAdd";
+            break;
+        case ShaderAtomicOp::AtomicCASFakeAdd:
+            o << "AtomicCASFakeAdd";
+            break;
+    }
+    return o;
+}
+
+using WorkgroupSizeParameter = int;
+using DispatchSizeParameter = int;
+using ShaderAtomicUseArray = bool;
+DAWN_TEST_PARAM_STRUCT(SubgroupsShaderTestsParams,
+                       ShaderAtomicUseArray,
+                       WorkgroupSizeParameter,
+                       DispatchSizeParameter,
+                       ShaderAtomicOp);
+
+class ShaderAtomicTests : public DawnTestWithParams<SubgroupsShaderTestsParams> {
+  public:
+    using DawnTestWithParams<SubgroupsShaderTestsParams>::GetParam;
+    using DawnTestWithParams<SubgroupsShaderTestsParams>::SupportsFeatures;
+    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(ShaderAtomicTests, WorkgroupAtomicArray) {
+    // TODO(crbug.com/42241359): Work in progress to resolve these tests mobile devices.
+    DAWN_SUPPRESS_TEST_IF(gpu_info::IsQualcomm_PCIAdreno6xx(GetParam().adapterProperties.vendorID,
+                                                            GetParam().adapterProperties.deviceID));
+
+    // Suppression for Mali gpus.
+    DAWN_SUPPRESS_TEST_IF(gpu_info::IsARM(GetParam().adapterProperties.vendorID));
+
+    // Test code only supports up to 256 workgroup size.
+    DAWN_ASSERT(GetParam().mWorkgroupSizeParameter <= 256);
+    std::string kConditionalArray = GetParam().mShaderAtomicUseArray ? "[0]" : "";
+    std::string kAtomicAddOperation =
+        R"(  atomicAdd(&workgroup_buffer)" + kConditionalArray + R"(, 1);)";
+    std::string kAtomicCASFakeAddOperation = R"(
+        var curr = atomicLoad(&workgroup_buffer)" +
+                                             kConditionalArray + R"();
+        var next = curr + 1;
+        var result = atomicCompareExchangeWeak(&workgroup_buffer)" +
+                                             kConditionalArray + R"(, curr, next);
+        while(!result.exchanged){
+            curr = result.old_value;
+            next = curr + 1;
+            result = atomicCompareExchangeWeak(&workgroup_buffer)" +
+                                             kConditionalArray + R"(, curr, next);
+        }
+    )";
+    std::stringstream code;
+    code << R"(
+@binding(0) @group(0) var<storage, read_write> output : array<u32>;
+// The bug requires an array to manifest.
+)"
+         << (GetParam().mShaderAtomicUseArray
+                 ? R"(var<workgroup> workgroup_buffer : array<atomic<u32>, 16>; )"
+                 : R"(var<workgroup> workgroup_buffer : atomic<u32>; )")
+         <<
+        R"(
+@compute @workgroup_size( )"
+         << GetParam().mWorkgroupSizeParameter
+         << R"(  )
+fn main(@builtin(local_invocation_index) local_invocation_index: u32,
+        @builtin(workgroup_id) workgroup_id : vec3<u32>){
+  if (local_invocation_index == 0) {
+     atomicStore(&workgroup_buffer)" +
+                kConditionalArray + R"(, 7);
+  }
+
+  workgroupBarrier();
+  )"
+         << (GetParam().mShaderAtomicOp == ShaderAtomicOp::AtomicAdd ? kAtomicAddOperation
+                                                                     : kAtomicCASFakeAddOperation)
+         <<
+
+        R"(
+
+  workgroupBarrier();
+
+  if (local_invocation_index == 0) {
+     output[workgroup_id.x] = atomicLoad(&workgroup_buffer)" +
+            kConditionalArray + R"();
+  }
+}
+)";
+
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(code.str());
+    constexpr int kInitStartValueShader = 7;
+    std::vector<uint32_t> expected(
+        GetParam().mDispatchSizeParameter,
+        static_cast<uint32_t>(GetParam().mWorkgroupSizeParameter + kInitStartValueShader));
+
+    wgpu::Buffer output = CreateBuffer(GetParam().mDispatchSizeParameter, -1);
+    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(GetParam().mDispatchSizeParameter);
+        pass.End();
+
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
+DAWN_INSTANTIATE_TEST_P(ShaderAtomicTests,
+                        /*Supporting only modern graphics backends for now.*/
+                        {D3D12Backend(), MetalBackend(), VulkanBackend()},
+                        {true, false}, /*use shader array*/
+                        {1,  2,  3,  4,  5,  6,   7,   8,   9,   13, 15,
+                         16, 31, 32, 53, 64, 111, 128, 137, 173, 256}, /* workgroup size*/
+                        {
+                            1,
+                        }, /*dispatch size */
+                        {ShaderAtomicOp::AtomicAdd, ShaderAtomicOp::AtomicCASFakeAdd});
+
+}  // anonymous namespace
+}  // namespace dawn