From 218d48890ab7886258e2ac2dbc9bfcc4de746fcb Mon Sep 17 00:00:00 2001
From: Ben Clayton <bclayton@google.com>
Date: Wed, 24 Mar 2021 09:48:11 +0000
Subject: [PATCH] Port ComputeSharedMemoryTests to WGSL

Tint now supports barriers

Bug: dawn:572
Change-Id: I1b5cdd9026e0ead7ecd743295c81ef147bdc5080
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/45602
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
---
 .../end2end/ComputeSharedMemoryTests.cpp      | 40 ++++++++++---------
 1 file changed, 21 insertions(+), 19 deletions(-)

diff --git a/src/tests/end2end/ComputeSharedMemoryTests.cpp b/src/tests/end2end/ComputeSharedMemoryTests.cpp
index 527f28d2b0..25844f2a5e 100644
--- a/src/tests/end2end/ComputeSharedMemoryTests.cpp
+++ b/src/tests/end2end/ComputeSharedMemoryTests.cpp
@@ -27,7 +27,7 @@ class ComputeSharedMemoryTests : public DawnTest {
 
 void ComputeSharedMemoryTests::BasicTest(const char* shader) {
     // Set up shader and pipeline
-    auto module = utils::CreateShaderModule(device, utils::SingleShaderStage::Compute, shader);
+    auto module = utils::CreateShaderModuleFromWGSL(device, shader);
 
     wgpu::ComputePipelineDescriptor csDesc;
     csDesc.computeStage.module = module;
@@ -70,31 +70,33 @@ void ComputeSharedMemoryTests::BasicTest(const char* shader) {
 
 // Basic shared memory test
 TEST_P(ComputeSharedMemoryTests, Basic) {
-    // TODO(crbug.com/tint/375): Implement barriers in Tint.
-    DAWN_SKIP_TEST_IF(HasToggleEnabled("use_tint_generator"));
-
     BasicTest(R"(
-        #version 450
-        const uint kTileSize = 4;
-        const uint kInstances = 11;
+        const kTileSize : u32 = 4;
+        const kInstances : u32 = 11;
 
-        layout(local_size_x = kTileSize, local_size_y = kTileSize, local_size_z = 1) in;
-        layout(std140, set = 0, binding = 0) buffer Dst { uint x; } dst;
-        shared uint tmp;
+        [[builtin(local_invocation_id)]] var<in> LocalInvocationID : vec3<u32>;
 
-        void main() {
-            uint index = gl_LocalInvocationID.y * kTileSize + gl_LocalInvocationID.x;
-            if (index == 0) {
-                tmp = 0;
+        [[block]] struct Dst {
+            x : u32;
+        };
+
+        [[group(0), binding(0)]] var<storage> dst : Dst;
+        var<workgroup> tmp : u32;
+
+        [[stage(compute), workgroup_size(4,4,1)]]
+        fn main() -> void {
+            var index : u32 = LocalInvocationID.y * kTileSize + LocalInvocationID.x;
+            if (index == 0u) {
+                tmp = 0u;
             }
-            barrier();
-            for (uint i = 0; i < kInstances; ++i) {
+            workgroupBarrier();
+            for (var i : u32 = 0u; i < kInstances; i = i + 1u) {
                 if (i == index) {
-                    tmp++;
+                    tmp = tmp + 1u;
                 }
-                barrier();
+                workgroupBarrier();
             }
-            if (index == 0) {
+            if (index == 0u) {
                 dst.x = tmp;
             }
         })");