Repository navigation
[Flutter GPU] Compute shaders on Metal and Vulkan - #193805
Conversation
c17a112 to
a344fee
Compare
There was a problem hiding this comment.
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.
3b120fe to
7de2074
Compare
gaaclarke
left a comment
There was a problem hiding this comment.
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 |
There was a problem hiding this comment.
👍 because it's not a breaking change, but previously created bundles will lack information that is needed to exist.
Factcheck me =)
There was a problem hiding this comment.
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.
|
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. |
7de2074 to
53c7bfe
Compare
|
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. |
|
/gemini review |
There was a problem hiding this comment.
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.
53c7bfe to
4f9ffb8
Compare
gaaclarke
left a comment
There was a problem hiding this comment.
Thanks brandon, this is looking really exciting. It's mostly looking good. High level here's my feedback:
- I'm concerned about editing the blit barriers for graphics, can we make those barriers more contextual to using compute shaders
- Use absl error handling and data structures
- Make sure we got docstrings on datatypes
| size_t binding = 0u; | ||
| }; | ||
|
|
||
| struct ShaderStorageBuffer { |
| TEST(ShaderBundleTest, RejectsComputeShaderWithSpecializedWorkgroupSize) { | ||
| EXPECT_FALSE( | ||
| GenerateComputeBundle("flutter_gpu_compute_specialized_size.comp") | ||
| .has_value()); | ||
| } |
There was a problem hiding this comment.
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.
| ASSERT_TRUE(bundle.has_value()); | ||
| // NOLINTNEXTLINE(bugprone-unchecked-optional-access) | ||
| for (const auto& shader : bundle->shaders) { | ||
| EXPECT_EQ(shader->vulkan->workgroup_size, nullptr); |
There was a problem hiding this comment.
I don't think you have a test that has a workgroup size specified?
There was a problem hiding this comment.
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(); |
There was a problem hiding this comment.
nit: use ScopedCleanupClosure to remove the label
| if (!encoder_) { | ||
| return; | ||
| } |
There was a problem hiding this comment.
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.
| if (auto found = buffer_bindings_.find(&binding->metadata); | ||
| found != buffer_bindings_.end()) { | ||
| view = &found->second; | ||
| } else if (const auto* set_binding = FindInSets( |
| for (int i = 0;; i++) { | ||
| const Shader::TextureBinding* binding = shader.GetUniformTextureAt(i); | ||
| if (!binding) { | ||
| break; | ||
| } |
| found != texture_bindings_.end()) { | ||
| texture = found->second.texture; | ||
| sampler = found->second.sampler; | ||
| } else if (const auto* set_binding = |
| 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" + |
| std::shared_ptr<impeller::Pipeline<impeller::ComputePipelineDescriptor>> | ||
| GetOrCreatePipeline(impeller::Context& context, std::string& out_error); |
There was a problem hiding this comment.
return
absl::StatusOr<std::shared_ptr<impeller::Pipeline<impeller::ComputePipelineDescriptor>>>(probably could use some typedefs)
|
Also, @bdero can you refrain from rebasing this PR after I've done this monster code review, please =) |
83e96f5 to
cb39a29
Compare
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).
cc69dab to
43a04fe
Compare
|
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: ). |
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
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.
gpuContext.createComputePipeline,CommandBuffer.createComputePassandComputePass, with storage, uniform and texture bindings, binding sets anddispatch.gpuContext.supportsComputeand the compute limits are exposed; on OpenGL ES compute is unsupported and the compute calls throw.Pre-launch Checklist
///).