[Impeller] Compute dispatch takes 3D workgroup counts and honors the shader local_size by bdero · Pull Request #188601 · flutter/flutter

bdero · 2026-06-26T00:53:49Z

Reworks the Impeller compute dispatch contract so it means the same thing on every backend and respects the workgroup size declared by the shader.

Before, ComputePass::Compute(ISize) was 2D and diverged per backend. Vulkan treated the argument as a workgroup count, while Metal treated it as a total invocation count and ignored the shader's local_size, sizing threadgroups to the device maximum. That made workgroup shared memory and 3D dispatch impossible, and produced very different invocation counts for the same call.

What changed:

Compute now takes 3D workgroup counts (Compute(x, y, z)), matching the dispatch model used by Metal, Vulkan, and the wider ecosystem. The per-workgroup invocation count is the shader's local_size.
The compiler reflects the shader's local_size into the generated header as kWorkgroupSize, and ComputePipelineDescriptor carries it. A dimension of 0 is a sentinel for "sized by a specialization constant", which the backend resolves at dispatch.
Metal passes the reflected local_size as threadsPerThreadgroup and dispatches the given workgroup counts, dropping the old halving heuristic.
Vulkan dispatches the workgroup counts directly and drops a bogus clamp of the count to the device's maximum workgroup size.

The Vulkan specialization-constant entry for workgroup size is retained. It only affects shaders that opt into a device-adaptive size via local_size_x_id (such as the prefix sum test). For shaders with a literal local_size it is a no-op, since Vulkan ignores specialization entries for constant IDs the shader does not use.

Compute is not yet exposed through Flutter GPU, so the only callers are the renderer compute unit tests, which are updated to dispatch workgroup counts. threadgroup_sizing_test.comp is converted to a literal local_size so it deterministically exercises the literal-size path.

Pre-launch Checklist

gemini-code-assist

Code Review

This pull request refactors Impeller's compute dispatching to use explicit workgroup counts instead of a grid size, utilizing the shader's declared workgroup size. Feedback on these changes identifies a potential issue in the Metal backend where fallback logic for zero-sized workgroups overrides non-zero Y and Z dimensions, which could cause dispatch failures or incorrect indexing. Additionally, a division-by-zero risk was noted in the test helper WorkgroupCount when the local size is zero.

gaaclarke · 2026-06-26T16:42:48Z


-  // Give all compute shaders a specialization constant entry for the
-  // workgroup/threadgroup size.
+  // Size specialization constant 0 to the device's maximum workgroup size.


incomplete sentence, awkward

Reworded in d69d06f.

gaaclarke · 2026-06-26T16:43:31Z

+  virtual fml::Status Compute(uint32_t workgroup_count_x,
+                              uint32_t workgroup_count_y = 1u,
+                              uint32_t workgroup_count_z = 1u) = 0;


nit: should the arg be a std::array<uint32_t,3>? That would make it easier for someone to make a constant for this.

Good idea, switched to std::array<uint32_t, 3> in d69d06f so callers can pass a constant and it matches GetWorkgroupSize.

gaaclarke · 2026-06-26T16:49:19Z

+  return *this;
+}
+
+const std::array<uint32_t, 3>& ComputePipelineDescriptor::GetWorkgroupSize()


nit: returning a copy for 12 byte payload is typically faster than returning a reference on a 64bit machine

Done in d69d06f, it returns by value now.

gaaclarke · 2026-06-26T16:52:37Z

+  // Intentionally making the workgroup count zero in one dimension. No GPU will
  // tolerate this.
-  EXPECT_FALSE(pass->Compute(ISize(0, 1)).ok());
+  EXPECT_FALSE(pass->Compute(0, 1, 1).ok());


instead of expecting false, let's have some positive assertions about the error

Done in d69d06f, it now asserts the kCancelled status code.

gaaclarke · 2026-06-26T16:54:11Z

I think we need a new test that takes advantage of the 3rd dimension, no? This patch is just updating the existing tests.

Added compute_3d_test.comp and a 3D dispatch test in d69d06f that checks indexing across all three dimensions and that the local size is honored per axis.

gaaclarke

LGTM, thanks brandon

gaaclarke · 2026-06-29T16:33:27Z

@walley892 I know you are interested in the compute shaders, what to be a second reviewer on this?

walley892 · 2026-06-30T02:35:29Z

+  // Unlike Vulkan and GLES, Metal does not bake the threadgroup size into the
+  // shader; it is supplied here at dispatch. Honor the shader's declared local
+  // size. A dimension of 0 means the shader sized it with a specialization
+  // constant. In that case fill the remaining threadgroup capacity into x while
+  // honoring any literal y and z, keeping the total within the device maximum.


Do we need to validate that workgroup_size[0] *workgroup_size[1]*workgroup_size[2] <= maxTotalThreadsPerThreadgroup?

Wasn't validated on Metal before. Vulkan rejects an oversized local size at pipeline creation; Metal supplies it at dispatch, so I added a maxTotalThreadsPerThreadgroup check that fails gracefully. Done in 58f41fd.

…12135) Manual roll Flutter from ca9f874f5284 to 6995038d96ef (44 revisions) Manual roll requested by [email protected] flutter/flutter@ca9f874...6995038 2026-07-03 [email protected] Roll Fuchsia Linux SDK from sx_eN0J_f2BV6jqjW... to Jr08vyeibMSv3Oxst... (flutter/flutter#188946) 2026-07-03 [email protected] Roll Packages from 420e135 to 2fbe873 (1 revision) (flutter/flutter#188945) 2026-07-03 [email protected] Roll Dart SDK from 786212a2ce0d to 1f9a08ce0638 (4 revisions) (flutter/flutter#188944) 2026-07-03 [email protected] Roll Skia from 5358ab75b840 to 919956953af6 (7 revisions) (flutter/flutter#188943) 2026-07-02 [email protected] [Windows] Keep regular windows in place when another is activated (flutter/flutter#188016) 2026-07-02 [email protected] [flutter_tools] Provide guided message when iOS/macOS build fails due to low minimum version (flutter/flutter#188812) 2026-07-02 [email protected] [Flutter GPU] Load a ShaderLibrary from shader bundle bytes (flutter/flutter#188596) 2026-07-02 [email protected] [tool] Enable record_use experiment by default on all channels (flutter/flutter#188887) 2026-07-02 [email protected] Roll Packages from e742106 to 420e135 (13 revisions) (flutter/flutter#188916) 2026-07-02 [email protected] Roll Fuchsia Linux SDK from I2h2eXk06RrA3pIG2... to sx_eN0J_f2BV6jqjW... (flutter/flutter#188915) 2026-07-02 [email protected] Clarify layout callback debug flag docs (flutter/flutter#186879) 2026-07-02 [email protected] Roll Skia from 0c4faca350cc to 5358ab75b840 (2 revisions) (flutter/flutter#188899) 2026-07-02 [email protected] Roll Dart SDK from e47361c7fe9a to 786212a2ce0d (2 revisions) (flutter/flutter#188898) 2026-07-01 [email protected] Roll Skia from 0fc8ba72e802 to 0c4faca350cc (2 revisions) (flutter/flutter#188886) 2026-07-01 [email protected] Hide draft PRs from the triage list (flutter/flutter#188885) 2026-07-01 [email protected] Stricten isSemantics and matchesSemantics children mismatch check (flutter/flutter#188827) 2026-07-01 [email protected] [AGP 9] Support Enabling Built-in Kotlin (flutter/flutter#188543) 2026-07-01 [email protected] Roll Skia from d19e557ac317 to 0fc8ba72e802 (4 revisions) (flutter/flutter#188879) 2026-07-01 [email protected] [ci] Increase test timeout for Mac_x64 build_tests shards (flutter/flutter#188804) 2026-07-01 [email protected] Roll Skia from bd4ae38ca3bb to d19e557ac317 (1 revision) (flutter/flutter#188865) 2026-07-01 [email protected] Roll Dart SDK from 26d723eb89af to e47361c7fe9a (5 revisions) (flutter/flutter#188864) 2026-07-01 [email protected] Update triage links for material_ui and cupertino_ui --> Design triage (flutter/flutter#188567) 2026-07-01 [email protected] [Impeller] Share a single ContextGLES among all PlaygroundImplGLES (flutter/flutter#188080) 2026-07-01 [email protected] Use null-aware elements in dev/devicelab/lib/integration_tests.dart (flutter/flutter#187852) 2026-07-01 [email protected] Roll Packages from 274ed3e to e742106 (18 revisions) (flutter/flutter#188863) 2026-07-01 [email protected] Add android 17 to embedding (flutter/flutter#187965) 2026-07-01 [email protected] Adds semantics role check to isSemantics and matchesSemantics (flutter/flutter#188825) 2026-07-01 [email protected] Roll Dart SDK from e1bdb9ce3327 to 26d723eb89af (3 revisions) (flutter/flutter#188795) 2026-07-01 [email protected] [web] Apply autocapitalize to text editing elements (flutter/flutter#188351) 2026-07-01 [email protected] Roll Fuchsia Linux SDK from RymJjIj7dd5vQ3Cnh... to I2h2eXk06RrA3pIG2... (flutter/flutter#188852) 2026-07-01 [email protected] [Impeller] Compute dispatch takes 3D workgroup counts and honors the shader local_size (flutter/flutter#188601) 2026-07-01 [email protected] Improve stylus support on linux (flutter/flutter#186831) 2026-07-01 [email protected] Resolve issue Catch am start failures with 'Error type' and prevent hang (flutter/flutter#187196) 2026-07-01 [email protected] Roll Skia from ef178c9898af to bd4ae38ca3bb (3 revisions) (flutter/flutter#188834) 2026-07-01 [email protected] Roll Skia from 3ac99be47229 to ef178c9898af (3 revisions) (flutter/flutter#188831) 2026-07-01 [email protected] Add a macosArm64Only feature flag (flutter/flutter#188598) 2026-07-01 [email protected] Roll Skia from 15302f1625b2 to 3ac99be47229 (1 revision) (flutter/flutter#188819) 2026-07-01 [email protected] [flutter_tools] Track asset transformer dependencies for hot reload (Reland #187947) (flutter/flutter#188808) 2026-06-30 [email protected] Add TapRegion samples (flutter/flutter#188685) 2026-06-30 [email protected] Print a warning in `flutter doctor` when running on Intel Macs (flutter/flutter#188760) 2026-06-30 [email protected] [framework] Keep scrollable semantics role stable (flutter/flutter#187963) 2026-06-30 [email protected] feat(skills): Add shepherd-prs skill for managing approved external contributor PRs (flutter/flutter#188534) 2026-06-30 [email protected] Roll Skia from 71947c4110b0 to 15302f1625b2 (17 revisions) (flutter/flutter#188815) 2026-06-30 [email protected] [Tool] Run re-entrant upgrade in original CWD (flutter/flutter#188794) ...

…lutter#12135) Manual roll Flutter from ca9f874f5284 to 6995038d96ef (44 revisions) Manual roll requested by [email protected] flutter/flutter@ca9f874...6995038 2026-07-03 [email protected] Roll Fuchsia Linux SDK from sx_eN0J_f2BV6jqjW... to Jr08vyeibMSv3Oxst... (flutter/flutter#188946) 2026-07-03 [email protected] Roll Packages from 420e135 to 2fbe873 (1 revision) (flutter/flutter#188945) 2026-07-03 [email protected] Roll Dart SDK from 786212a2ce0d to 1f9a08ce0638 (4 revisions) (flutter/flutter#188944) 2026-07-03 [email protected] Roll Skia from 5358ab75b840 to 919956953af6 (7 revisions) (flutter/flutter#188943) 2026-07-02 [email protected] [Windows] Keep regular windows in place when another is activated (flutter/flutter#188016) 2026-07-02 [email protected] [flutter_tools] Provide guided message when iOS/macOS build fails due to low minimum version (flutter/flutter#188812) 2026-07-02 [email protected] [Flutter GPU] Load a ShaderLibrary from shader bundle bytes (flutter/flutter#188596) 2026-07-02 [email protected] [tool] Enable record_use experiment by default on all channels (flutter/flutter#188887) 2026-07-02 [email protected] Roll Packages from e742106 to 420e135 (13 revisions) (flutter/flutter#188916) 2026-07-02 [email protected] Roll Fuchsia Linux SDK from I2h2eXk06RrA3pIG2... to sx_eN0J_f2BV6jqjW... (flutter/flutter#188915) 2026-07-02 [email protected] Clarify layout callback debug flag docs (flutter/flutter#186879) 2026-07-02 [email protected] Roll Skia from 0c4faca350cc to 5358ab75b840 (2 revisions) (flutter/flutter#188899) 2026-07-02 [email protected] Roll Dart SDK from e47361c7fe9a to 786212a2ce0d (2 revisions) (flutter/flutter#188898) 2026-07-01 [email protected] Roll Skia from 0fc8ba72e802 to 0c4faca350cc (2 revisions) (flutter/flutter#188886) 2026-07-01 [email protected] Hide draft PRs from the triage list (flutter/flutter#188885) 2026-07-01 [email protected] Stricten isSemantics and matchesSemantics children mismatch check (flutter/flutter#188827) 2026-07-01 [email protected] [AGP 9] Support Enabling Built-in Kotlin (flutter/flutter#188543) 2026-07-01 [email protected] Roll Skia from d19e557ac317 to 0fc8ba72e802 (4 revisions) (flutter/flutter#188879) 2026-07-01 [email protected] [ci] Increase test timeout for Mac_x64 build_tests shards (flutter/flutter#188804) 2026-07-01 [email protected] Roll Skia from bd4ae38ca3bb to d19e557ac317 (1 revision) (flutter/flutter#188865) 2026-07-01 [email protected] Roll Dart SDK from 26d723eb89af to e47361c7fe9a (5 revisions) (flutter/flutter#188864) 2026-07-01 [email protected] Update triage links for material_ui and cupertino_ui --> Design triage (flutter/flutter#188567) 2026-07-01 [email protected] [Impeller] Share a single ContextGLES among all PlaygroundImplGLES (flutter/flutter#188080) 2026-07-01 [email protected] Use null-aware elements in dev/devicelab/lib/integration_tests.dart (flutter/flutter#187852) 2026-07-01 [email protected] Roll Packages from 274ed3e to e742106 (18 revisions) (flutter/flutter#188863) 2026-07-01 [email protected] Add android 17 to embedding (flutter/flutter#187965) 2026-07-01 [email protected] Adds semantics role check to isSemantics and matchesSemantics (flutter/flutter#188825) 2026-07-01 [email protected] Roll Dart SDK from e1bdb9ce3327 to 26d723eb89af (3 revisions) (flutter/flutter#188795) 2026-07-01 [email protected] [web] Apply autocapitalize to text editing elements (flutter/flutter#188351) 2026-07-01 [email protected] Roll Fuchsia Linux SDK from RymJjIj7dd5vQ3Cnh... to I2h2eXk06RrA3pIG2... (flutter/flutter#188852) 2026-07-01 [email protected] [Impeller] Compute dispatch takes 3D workgroup counts and honors the shader local_size (flutter/flutter#188601) 2026-07-01 [email protected] Improve stylus support on linux (flutter/flutter#186831) 2026-07-01 [email protected] Resolve issue Catch am start failures with 'Error type' and prevent hang (flutter/flutter#187196) 2026-07-01 [email protected] Roll Skia from ef178c9898af to bd4ae38ca3bb (3 revisions) (flutter/flutter#188834) 2026-07-01 [email protected] Roll Skia from 3ac99be47229 to ef178c9898af (3 revisions) (flutter/flutter#188831) 2026-07-01 [email protected] Add a macosArm64Only feature flag (flutter/flutter#188598) 2026-07-01 [email protected] Roll Skia from 15302f1625b2 to 3ac99be47229 (1 revision) (flutter/flutter#188819) 2026-07-01 [email protected] [flutter_tools] Track asset transformer dependencies for hot reload (Reland #187947) (flutter/flutter#188808) 2026-06-30 [email protected] Add TapRegion samples (flutter/flutter#188685) 2026-06-30 [email protected] Print a warning in `flutter doctor` when running on Intel Macs (flutter/flutter#188760) 2026-06-30 [email protected] [framework] Keep scrollable semantics role stable (flutter/flutter#187963) 2026-06-30 [email protected] feat(skills): Add shepherd-prs skill for managing approved external contributor PRs (flutter/flutter#188534) 2026-06-30 [email protected] Roll Skia from 71947c4110b0 to 15302f1625b2 (17 revisions) (flutter/flutter#188815) 2026-06-30 [email protected] [Tool] Run re-entrant upgrade in original CWD (flutter/flutter#188794) ...

flutter-dashboard Bot added the CICD Run CI/CD label Jun 26, 2026

github-actions Bot added engine flutter/engine related. See also e: labels. e: impeller Impeller rendering backend issues and features requests labels Jun 26, 2026

bdero added the flutter-gpu label Jun 26, 2026

github-project-automation Bot added this to Flutter GPU Jun 26, 2026

github-project-automation Bot moved this to 🤔 Needs Triage in Flutter GPU Jun 26, 2026

bdero marked this pull request as ready for review June 26, 2026 03:45

gemini-code-assist Bot reviewed Jun 26, 2026

View reviewed changes

Comment thread engine/src/flutter/impeller/renderer/backend/metal/compute_pass_mtl.mm

Comment thread engine/src/flutter/impeller/renderer/compute_unittests.cc

github-actions Bot removed flutter-gpu CICD Run CI/CD labels Jun 26, 2026

bdero added CICD Run CI/CD flutter-gpu labels Jun 26, 2026

bdero requested a review from gaaclarke June 26, 2026 10:45

gaaclarke reviewed Jun 26, 2026

View reviewed changes

github-actions Bot removed the flutter-gpu label Jun 27, 2026

bdero added CICD Run CI/CD and removed CICD Run CI/CD labels Jun 27, 2026

bdero requested a review from gaaclarke June 29, 2026 07:36

gaaclarke previously approved these changes Jun 29, 2026

View reviewed changes

gaaclarke requested a review from walley892 June 29, 2026 16:32

bdero added CICD Run CI/CD and removed CICD Run CI/CD labels Jun 30, 2026

bdero dismissed gaaclarke’s stale review via 0fe0ba9 June 30, 2026 01:05

bdero force-pushed the bdero/compute-dispatch-contract branch from 5a60e7d to 0fe0ba9 Compare June 30, 2026 01:05

bdero added CICD Run CI/CD and removed CICD Run CI/CD labels Jun 30, 2026

walley892 reviewed Jun 30, 2026

View reviewed changes

Uh oh!

Conversation

bdero commented Jun 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Pre-launch Checklist

Uh oh!

gemini-code-assist Bot left a comment

Choose a reason for hiding this comment

Code Review

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gaaclarke left a comment

Choose a reason for hiding this comment

Uh oh!

gaaclarke commented Jun 29, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

bdero commented Jun 26, 2026 •

edited

Loading