From 8e199f643a4bd7e735eff8ffbb10af1c7f75738b Mon Sep 17 00:00:00 2001
From: Greggman <github@greggman.com>
Date: Fri, 2 Feb 2024 02:49:08 +0900
Subject: [PATCH] Fix maxComputeWorkgroupStorageSize test (#3357)

The spec says each var<workgroup> is roundUp(16, SizeOf(T)).
---
 .../limits/maxComputeWorkgroupStorageSize.spec.ts        | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/src/webgpu/api/validation/capability_checks/limits/maxComputeWorkgroupStorageSize.spec.ts b/src/webgpu/api/validation/capability_checks/limits/maxComputeWorkgroupStorageSize.spec.ts
index cb26e18ebeb1..8a9176983d38 100644
--- a/src/webgpu/api/validation/capability_checks/limits/maxComputeWorkgroupStorageSize.spec.ts
+++ b/src/webgpu/api/validation/capability_checks/limits/maxComputeWorkgroupStorageSize.spec.ts
@@ -12,7 +12,8 @@ import {
 const limit = 'maxComputeWorkgroupStorageSize';
 export const { g, description } = makeLimitTestGroup(limit);
 
-const kSmallestWorkgroupVarSize = 4;
+// Each var is roundUp(16, SizeOf(T))
+const kSmallestWorkgroupVarSize = 16;
 
 const wgslF16Types = {
   f16: { alignOf: 2, sizeOf: 2, requireF16: true },
@@ -71,7 +72,9 @@ function getModuleForWorkgroupStorageSize(device: GPUDevice, wgslType: WGSLType,
   const { sizeOf, alignOf, requireF16 } = wgslTypes[wgslType];
   const unitSize = align(sizeOf, alignOf);
   const units = Math.floor(size / unitSize);
-  const extra = (size - units * unitSize) / kSmallestWorkgroupVarSize;
+  const sizeUsed = align(units * unitSize, 16);
+  const sizeLeft = size - sizeUsed;
+  const extra = Math.floor(sizeLeft / kSmallestWorkgroupVarSize);
 
   const code =
     (requireF16 ? 'enable f16;\n' : '') +
@@ -89,7 +92,7 @@ function getModuleForWorkgroupStorageSize(device: GPUDevice, wgslType: WGSLType,
       b: vec2f,
     };
     var<workgroup> d0: array<${wgslType}, ${units}>;
-    ${extra ? `var<workgroup> d1: array<f32, ${extra}>;` : ''}
+    ${extra ? `var<workgroup> d1: array<vec4<f32>, ${extra}>;` : ''}
     @compute @workgroup_size(1) fn main() {
       _ = d0;
       ${extra ? '_ = d1;' : ''}