Skip to content

[Flutter GPU] Compute shaders on Metal and Vulkan - #193805

Merged
auto-submit[bot] merged 16 commits into
flutter:masterfrom
bdero:bdero/flutter-gpu-compute-metadata
Oct 26, 2026
Merged

auto-submit[bot] merged 16 commits into
flutter:masterfrom
bdero:bdero/flutter-gpu-compute-metadata

Conversation

@bdero

@bdero bdero commented Oct 4, 2026 •

Copy link
Copy Markdown
Member

Fixes #188477, fixes #188479, fixes #188480, fixes #193868, fixes #193869, fixes #193872. Part of #188474 and #188481.

Adds compute to Flutter GPU on Metal and Vulkan. Stacked on #194167.

  • Shader bundles (format version 3) carry each shader's storage buffers and compute workgroup size.
  • gpuContext.createComputePipeline, CommandBuffer.createComputePass and ComputePass, with storage, uniform and texture bindings, binding sets and dispatch. gpuContext.supportsCompute and the compute limits are exposed; on OpenGL ES compute is unsupported and the compute calls throw.
  • Impeller: workgroup sizes are checked against device limits when a pipeline is created, bindings persist across dispatches on Vulkan, dispatches and passes synchronize automatically, and reloading a compute shader no longer crashes.
  • End-to-end tests on Metal and Vulkan, including a compute pass writing vertices that a render pass draws.

Pre-launch Checklist

@github-actions github-actions Bot added engine flutter/engine related. See also e: labels. e: impeller Impeller rendering backend issues and features requests flutter-gpu team-fluttergpu Owned by Flutter GPU team labels Oct 4, 2026
@bdero
bdero force-pushed the bdero/flutter-gpu-compute-metadata branch from c17a112 to a344fee Compare October 5, 2026 22:26
@bdero bdero added the CICD Run CI/CD label Oct 5, 2026
@bdero
bdero marked this pull request as ready for review October 6, 2026 00:13
@bdero
bdero requested a review from gaaclarke October 6, 2026 00:14

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces support for compute shaders in the Impeller shader compiler and Flutter GPU runtime by reflecting, serializing, and loading workgroup sizes and storage buffer metadata. It updates the shader bundle FlatBuffer schema to version 3 and adds comprehensive unit tests. Feedback on the changes includes addressing a potential null pointer dereference when parsing optional FlatBuffer string fields for storage buffer names, and simplifying a workgroup size validation check by directly inspecting array elements instead of using std::find.

@bdero
bdero force-pushed the bdero/flutter-gpu-compute-metadata branch from 3b120fe to 7de2074 Compare October 6, 2026 00:28

@gaaclarke gaaclarke left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is looking good to me. Since we are incrementing the shader bundle version I do have a concern that we don't have a full integration test. It would be nice to see what didn't work before this PR now working holistically to justify the addition of the new data. Thoughts?

// resulting flat buffer format.
enum ShaderBundleFormatVersion:uint32 {
kVersion = 2
kVersion = 3

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

👍 because it's not a breaking change, but previously created bundles will lack information that is needed to exist.

Factcheck me =)

@bdero bdero Oct 11, 2026 •

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Mostly right. Old v2 bundles won't load until they're rebuilt, since the version has to match exactly, but flutter_gpu_shaders rebuilds them automatically when impellerc changes. Without the bump they'd load with the compute metadata missing, so the bump gives a clear error instead.

@bdero

bdero commented Oct 11, 2026 •

Copy link
Copy Markdown
Member Author

Agreed. I'm pulling the compute pass API (#188480) and what it depends on into this PR, so it will come with Dart tests on Metal and Vulkan that dispatch a compute shader and check the output. I'll ping you when it's up.

@bdero bdero mentioned this pull request Oct 11, 2026
2 of 22 tasks
@bdero
bdero force-pushed the bdero/flutter-gpu-compute-metadata branch from 7de2074 to 53c7bfe Compare October 11, 2026 03:25
@bdero
bdero changed the base branch from master to bdero/flutter-gpu-pass-lifecycle October 11, 2026 03:25
@bdero bdero changed the title [Flutter GPU] Shader bundles carry compute workgroup size and storage buffers [Flutter GPU] Compute shaders on Metal and Vulkan Oct 11, 2026
@bdero

bdero commented Oct 11, 2026

Copy link
Copy Markdown
Member Author

This now includes the compute pass API, with end-to-end tests on Metal and Vulkan that dispatch a compute shader and draw its output. It's stacked on #194167, so the diff shows only the compute changes until that lands.

@bdero
bdero requested a review from gaaclarke October 11, 2026 03:37
@bdero

bdero commented Oct 11, 2026

Copy link
Copy Markdown
Member Author

/gemini review

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces compute shader support to Flutter GPU, adding ComputePass and ComputePipeline APIs to both the C++ engine and Dart bindings, alongside backend implementations for Metal and Vulkan. The review feedback suggests preventing write-after-write hazards in Vulkan blit passes by adding write access flags to barriers, skipping global memory barriers in Vulkan compute passes when no dispatches are recorded, and refactoring buffer bounds checks in ComputePass and BindingSet to use subtraction to avoid potential integer overflow vulnerabilities.

@bdero
bdero force-pushed the bdero/flutter-gpu-compute-metadata branch from 53c7bfe to 4f9ffb8 Compare October 11, 2026 07:07

@gaaclarke gaaclarke left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks brandon, this is looking really exciting. It's mostly looking good. High level here's my feedback:

  1. I'm concerned about editing the blit barriers for graphics, can we make those barriers more contextual to using compute shaders
  2. Use absl error handling and data structures
  3. Make sure we got docstrings on datatypes

size_t binding = 0u;
};

struct ShaderStorageBuffer {

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Add docstring please.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 24838e2.

Comment on lines +475 to +479
TEST(ShaderBundleTest, RejectsComputeShaderWithSpecializedWorkgroupSize) {
EXPECT_FALSE(
GenerateComputeBundle("flutter_gpu_compute_specialized_size.comp")
.has_value());
}

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

no action required: I really wish this function used absl::Status instead of std::optional to denote error states. Not introduced here though so no need to address.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Agreed.

ASSERT_TRUE(bundle.has_value());
// NOLINTNEXTLINE(bugprone-unchecked-optional-access)
for (const auto& shader : bundle->shaders) {
EXPECT_EQ(shader->vulkan->workgroup_size, nullptr);

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think you have a test that has a workgroup size specified?

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This one checks render shaders have none. GenerateShaderBundleFlatbufferReflectsComputeShader covers the 8x4x2 size.

workgroup_count[2] == 0u) {
return fml::Status(fml::StatusCode::kCancelled,
"Invalid workgroup count for compute command.");
PopCommandLabel();

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: use ScopedCleanupClosure to remove the label

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 24838e2.

Comment on lines 39 to 41
if (!encoder_) {
return;
}

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

no action required: This should be a factory function that returns absl::StatusOr instead of creating invalid objects. If encoder fails to be created the destructor will be fine because objc eats calls to nil.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Agreed.

if (auto found = buffer_bindings_.find(&binding->metadata);
found != buffer_bindings_.end()) {
view = &found->second;
} else if (const auto* set_binding = FindInSets(

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

remove auto please

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 24838e2.

Comment on lines +227 to +231
for (int i = 0;; i++) {
const Shader::TextureBinding* binding = shader.GetUniformTextureAt(i);
if (!binding) {
break;
}

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same here

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 24838e2.

found != texture_bindings_.end()) {
texture = found->second.texture;
sampler = found->second.sampler;
} else if (const auto* set_binding =

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

remove auto

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 24838e2.

context_->GetCapabilities()->GetMaximumComputeWorkgroupCount();
if (workgroup_count[0] > max_count[0] || workgroup_count[1] > max_count[1] ||
workgroup_count[2] > max_count[2]) {
return "The workgroup count " + std::to_string(workgroup_count[0]) + "x" +

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

absl::StrCat

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 1f36261.

Comment on lines +41 to +42
std::shared_ptr<impeller::Pipeline<impeller::ComputePipelineDescriptor>>
GetOrCreatePipeline(impeller::Context& context, std::string& out_error);

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

return

absl::StatusOr<std::shared_ptr<impeller::Pipeline<impeller::ComputePipelineDescriptor>>>

(probably could use some typedefs)

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 1f36261.

@gaaclarke

Copy link
Copy Markdown
Member

Also, @bdero can you refrain from rebasing this PR after I've done this monster code review, please =)

@bdero
bdero force-pushed the bdero/flutter-gpu-pass-lifecycle branch from 83e96f5 to cb39a29 Compare October 19, 2026 23:05
@auto-submit
auto-submit Bot deleted the branch flutter:bdero/flutter-gpu-pass-lifecycle October 20, 2026 22:06
@auto-submit auto-submit Bot closed this Oct 20, 2026
@github-project-automation github-project-automation Bot moved this from 🤔 Needs Triage to ✅ Done in Flutter GPU Oct 20, 2026
@bdero bdero reopened this Oct 20, 2026
bdero added 12 commits October 24, 2026 17:31
RemovePipelinesWithEntryPoint dereferenced the entry point for the removed
function's stage on every cached render pipeline, which is null when a
compute function is removed. Skip pipelines without an entry point for
that stage on Metal, Vulkan and OpenGL ES, and evict the cached compute
pipelines that use the function on Metal and Vulkan.

Fixes flutter#193872.
… sets

Shader.getStorageBufferSlot looks up a compute shader's storage buffer by
name, caching its reflection index per hot reload like UniformSlot.
Binding sets take storage buffers. A bound view is checked in Dart: it
covers the block, holds whole elements of a trailing runtime-sized array,
and starts at a multiple of the new
GpuContext.minimumStorageBufferByteAlignment. HostBuffer aligns each
emplacement to both the uniform and the storage alignment.

Fixes flutter#188479.
GpuContext.createComputePipeline builds a pipeline from one compute
shader, with the shader's workgroup size and descriptor set layouts. A
workgroup size over the device limits throws with the reason. A pipeline
whose shader hot reloads rebuilds at its next dispatch. GpuContext
exposes the compute limits: workgroup invocations, workgroup size,
workgroup count and shared memory size.

CommandBuffer.createComputePass returns a ComputePass that follows the
pass lifecycle. It binds pipelines, storage buffers, uniforms, sampled
textures and binding sets, which persist across dispatches, and
dispatch takes workgroup counts. A dispatch throws if a binding the
shader declares is unbound. A compute shader keeps the resources
Metal's compiler removes, so it takes the same bindings on every
backend; the pass skips the backend bind for them on Metal.

GpuContext.supportsCompute is true on Metal and Vulkan and false on
OpenGL ES, where compute pipeline and pass creation throw. A Metal
compute pass now ends its encoder when it is released unended, like
render and blit passes.

Fixes flutter#188480.
Adds compute shaders to the test shader bundles and gpu_test.dart
coverage for a compute pass writing vertices that a render pass draws,
dependent dispatches without a barrier, binding sets, sampled textures,
storage buffer bind validation, unused resources that must be bound on
every backend, the pass lifecycle, the compute limits, the workgroup
size limit, and the supportsCompute gate on OpenGL ES. A reload test
checks that an existing ComputePipeline runs reloaded shader code.

Part of flutter#188481.
The prefix sum test shader now uses a literal workgroup size, and the
dispatch contract and synchronization tests add six shaders.
Instead of widening the blit barriers and the render pass outgoing
dependency to include the compute stage, the first dispatch of a Vulkan
compute pass records one barrier that waits on earlier graphics, transfer
and host writes, so only command buffers that use compute pay for it.
The blit pass and render pass builder are back to their state on master.
Adds tests for compute passes that read a texture or a buffer written by
a blit pass.
ComputePipelineDescriptor::ValidateWorkgroupSize returns an absl::Status
and builds its message only on error. In Flutter GPU, ComputePass,
ComputePipeline, Shader::RegisterSync and CommandBuffer::EndComputePass
return absl::Status or absl::StatusOr, with messages built by
absl::StrCat. The FFI entry points convert the status to the error
string or bool Dart already receives. Binding maps use
absl::flat_hash_map, and a null pipeline in ComputePassVK::SetPipeline
or a null pass in EndComputePass is a DCHECK.
Adds docstrings to the storage buffer reflection types, the shader
bundle schema, ComputePassVK (in doxygen style) and the other types the
change adds, and documents the lifetime of the binding returned by
Shader::GetStorageBufferAt. Groups the compute limits in Capabilities
into a ComputeLimits struct, pops the debug label with a
ScopedCleanupClosure, adds argument comments to the Vulkan calls,
removes the performance-move-const-arg suppressions and simplifies the
binding set lookup in the Flutter GPU compute pass, and builds the test
shader bundle descriptions in BUILD.gn from one entry per line.
Shader.getStorageBufferSlot throws an ArgumentError for a name the shader
does not declare, so StorageBufferSlot.sizeInBytes and
runtimeArrayStrideInBytes are no longer nullable. If a hot reload removes
the storage buffer, using the slot throws a StateError. The reload test
bundle now replaces ComputeScaleWithUnused with a shader that drops its
Unused storage buffer, to test that.
Adds a test that draws a computed triangle covering part of the viewport
and checks pixels inside and outside each of its edges, rather than only
that the corners are one color. Names the lime and red pixel values the
compute tests read back, and notes that lime is the full-intensity green.
…buffer

The storage buffer slot and the Vulkan bound resource type now have
default values, for cppcoreguidelines-pro-type-member-init. The blit
readback test writes its device-private texture through a staging buffer,
since Metal disallows CPU writes to private textures.
Master replaced fml::Status with absl (flutter#194161).
@bdero
bdero force-pushed the bdero/flutter-gpu-compute-metadata branch from cc69dab to 43a04fe Compare October 25, 2026 02:14
@gaaclarke gaaclarke added the autosubmit Merge PR when tree becomes green via auto submit App label Oct 26, 2026
@auto-submit
auto-submit Bot added this pull request to the merge queue Oct 26, 2026
@auto-submit auto-submit Bot removed the autosubmit Merge PR when tree becomes green via auto submit App label Oct 26, 2026
@auto-submit

auto-submit Bot commented Oct 26, 2026

Copy link
Copy Markdown
Contributor

auto label is removed for flutter/flutter/193805, Failed to enqueue flutter/flutter/193805 with HTTP 400: UnknownException(TimeoutException after 0:00:05.000000: No stream event, stack:

).

Merged via the queue into flutter:master with commit a3a21c8 Oct 26, 2026
34 of 35 checks passed
@github-project-automation github-project-automation Bot moved this from 🤔 Needs Triage to ✅ Done in Flutter GPU Oct 26, 2026
suryaKommana2662 pushed a commit to suryaKommana2662/flutter that referenced this pull request Oct 29, 2026
Fixes flutter#188477, fixes flutter#188479, fixes
flutter#188480, fixes flutter#193868, fixes
flutter#193869, fixes flutter#193872. Part of
flutter#188474 and flutter#188481.

Adds compute to Flutter GPU on Metal and Vulkan. Stacked on flutter#194167.

- Shader bundles (format version 3) carry each shader's storage buffers
and compute workgroup size.
- `gpuContext.createComputePipeline`, `CommandBuffer.createComputePass`
and `ComputePass`, with storage, uniform and texture bindings, binding
sets and `dispatch`. `gpuContext.supportsCompute` and the compute limits
are exposed; on OpenGL ES compute is unsupported and the compute calls
throw.
- Impeller: workgroup sizes are checked against device limits when a
pipeline is created, bindings persist across dispatches on Vulkan,
dispatches and passes synchronize automatically, and reloading a compute
shader no longer crashes.
- End-to-end tests on Metal and Vulkan, including a compute pass writing
vertices that a render pass draws.

## Pre-launch Checklist

- [x] I read the [Contributor Guide] and followed the process outlined
there for submitting PRs.
- [x] I read the [AI contribution guidelines] and understand my
responsibilities, or I am not using AI tools.
- [x] I read the [Tree Hygiene] wiki page, which explains my
responsibilities.
- [x] I read and followed the [Flutter Style Guide], including [Features
we expect every widget to implement].
- [x] I signed the [CLA].
- [x] I listed at least one issue that this PR fixes in the description
above.
- [x] I updated/added relevant documentation (doc comments with `///`).
- [x] I added new tests to check the change I am making, or this PR is
[test-exempt].
- [x] I followed the [breaking change policy] and added [Data Driven
Fixes] where supported.
- [ ] All existing and new tests are passing.

<!-- Links -->
[Contributor Guide]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Tree-hygiene.md#overview
[AI contribution guidelines]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Tree-hygiene.md#ai-contribution-guidelines
[Tree Hygiene]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Tree-hygiene.md
[test-exempt]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Tree-hygiene.md#tests
[Flutter Style Guide]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Style-guide-for-Flutter-repo.md
[Features we expect every widget to implement]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Style-guide-for-Flutter-repo.md#features-we-expect-every-widget-to-implement
[CLA]: https://cla.developers.google.com/
[breaking change policy]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Tree-hygiene.md#handling-breaking-changes
[Data Driven Fixes]:
https://github.com/flutter/flutter/blob/main/docs/contributing/Data-driven-Fixes.md
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CICD Run CI/CD e: impeller Impeller rendering backend issues and features requests engine flutter/engine related. See also e: labels. flutter-gpu team-fluttergpu Owned by Flutter GPU team

Projects

Status: ✅ Done

3 participants