[ir] Validate intermediate workgroup size products

When checking that the total workgroup size is less that UINT32_MAX,
we were only checking the final x*y*z product. This may overflow a
uint64_t and wrap around to be a valid uint32_t value, so we need to
check the intermediate products instead.

Fixed: 463283605
Change-Id: Ie4cb2354bc6693230b5b591d152ec3d95b0469c4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/277935
Reviewed-by: Peter McNeeley <petermcneeley@google.com>
Commit-Queue: James Price <jrprice@google.com>
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
Auto-Submit: James Price <jrprice@google.com>

Author: jrprice

src/tint/lang/core/ir/validator.cc

@@ -2970,6 +2970,12 @@ void Validator::CheckWorkgroupSize(const Function* func) {
                 return;
             }
             total_size *= c->Value()->ValueAs<uint64_t>();
+
+            constexpr uint64_t kMaxGridSize = 0xffffffff;
+            if (total_size > kMaxGridSize) {
+                AddError(func) << "workgroup grid size cannot exceed 0x" << std::hex
+                               << kMaxGridSize;
+            }
             continue;
         }
 
@@ -3007,11 +3013,6 @@ void Validator::CheckWorkgroupSize(const Function* func) {
 
         AddError(func) << "@workgroup_size must be an InstructionResult or a Constant";
     }
-
-    constexpr uint64_t kMaxGridSize = 0xffffffff;
-    if (total_size > kMaxGridSize) {
-        AddError(func) << "workgroup grid size cannot exceed 0x" << std::hex << kMaxGridSize;
-    }
 }
 
 void Validator::CheckPositionPresentForVertexOutput(const Function* ep) {

src/tint/lang/core/ir/validator_function_test.cc

@@ -2498,6 +2498,24 @@ TEST_F(IR_ValidatorTest, Function_WorkgroupSize_ParamsTooLarge) {
 )")) << res.Failure();
 }
 
+// Test the case where the intermediate workgroup product overflows a uint64_t and wraps back around
+// to be a valid uint32_t value.
+TEST_F(IR_ValidatorTest, Function_WorkgroupSize_ParamsTooLarge_U64Overflow) {
+    auto* f = ComputeEntryPoint();
+    f->SetWorkgroupSize(
+        {b.Constant(1526726656_i), b.Constant(1526726656_i), b.Constant(1526726656_i)});
+
+    b.Append(f->Block(), [&] { b.Unreachable(); });
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_THAT(res.Failure().reason,
+                testing::HasSubstr(R"(:1:1 error: workgroup grid size cannot exceed 0xffffffff
+%f = @compute @workgroup_size(1526726656i, 1526726656i, 1526726656i) func():void {
+^^
+)")) << res.Failure();
+}
+
 TEST_F(IR_ValidatorTest, Function_WorkgroupSize_OverrideWithoutAllowOverrides) {
     auto* o = b.Override(ty.u32());
     auto* f = ComputeEntryPoint();