[Impeller] dont emulate command buffers for compute on Metal/Vulkan. (#49922)
Part of the work for: https://github.com/flutter/flutter/issues/140804
Removes ComputeCommand and has both ComputePassMTL and ComputePassVK immediately record to their respective comand buffers. This combines the commits for Metal/Vulkan as they are both fairly small. After this, I can add barriers to fix compute for Vulkan.
diff --git a/ci/licenses_golden/excluded_files b/ci/licenses_golden/excluded_files
index 98b6d25..e4445af 100644
--- a/ci/licenses_golden/excluded_files
+++ b/ci/licenses_golden/excluded_files
@@ -182,7 +182,6 @@
../../../flutter/impeller/renderer/backend/vulkan/context_vk_unittests.cc
../../../flutter/impeller/renderer/backend/vulkan/descriptor_pool_vk_unittests.cc
../../../flutter/impeller/renderer/backend/vulkan/fence_waiter_vk_unittests.cc
-../../../flutter/impeller/renderer/backend/vulkan/pass_bindings_cache_unittests.cc
../../../flutter/impeller/renderer/backend/vulkan/resource_manager_vk_unittests.cc
../../../flutter/impeller/renderer/backend/vulkan/test
../../../flutter/impeller/renderer/blit_pass_unittests.cc
diff --git a/ci/licenses_golden/licenses_flutter b/ci/licenses_golden/licenses_flutter
index 6bee18a..58fd408 100644
--- a/ci/licenses_golden/licenses_flutter
+++ b/ci/licenses_golden/licenses_flutter
@@ -5387,6 +5387,8 @@
ORIGIN: ../../../flutter/impeller/renderer/backend/metal/blit_pass_mtl.mm + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/metal/command_buffer_mtl.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/metal/command_buffer_mtl.mm + ../../../flutter/LICENSE
+ORIGIN: ../../../flutter/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.h + ../../../flutter/LICENSE
+ORIGIN: ../../../flutter/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.mm + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/metal/compute_pass_mtl.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/metal/compute_pass_mtl.mm + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/metal/compute_pipeline_mtl.h + ../../../flutter/LICENSE
@@ -5431,8 +5433,6 @@
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/android_hardware_buffer_texture_source_vk.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/barrier_vk.cc + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/barrier_vk.h + ../../../flutter/LICENSE
-ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/binding_helpers_vk.cc + ../../../flutter/LICENSE
-ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/binding_helpers_vk.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/blit_command_vk.cc + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/blit_command_vk.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/blit_pass_vk.cc + ../../../flutter/LICENSE
@@ -5464,8 +5464,6 @@
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/formats_vk.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/gpu_tracer_vk.cc + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/gpu_tracer_vk.h + ../../../flutter/LICENSE
-ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/pass_bindings_cache.cc + ../../../flutter/LICENSE
-ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/pass_bindings_cache.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/pipeline_cache_vk.cc + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/pipeline_cache_vk.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/backend/vulkan/pipeline_library_vk.cc + ../../../flutter/LICENSE
@@ -5519,8 +5517,6 @@
ORIGIN: ../../../flutter/impeller/renderer/command.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/command_buffer.cc + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/command_buffer.h + ../../../flutter/LICENSE
-ORIGIN: ../../../flutter/impeller/renderer/compute_command.cc + ../../../flutter/LICENSE
-ORIGIN: ../../../flutter/impeller/renderer/compute_command.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/compute_pass.cc + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/compute_pass.h + ../../../flutter/LICENSE
ORIGIN: ../../../flutter/impeller/renderer/compute_pipeline_builder.cc + ../../../flutter/LICENSE
@@ -8231,6 +8227,8 @@
FILE: ../../../flutter/impeller/renderer/backend/metal/blit_pass_mtl.mm
FILE: ../../../flutter/impeller/renderer/backend/metal/command_buffer_mtl.h
FILE: ../../../flutter/impeller/renderer/backend/metal/command_buffer_mtl.mm
+FILE: ../../../flutter/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.h
+FILE: ../../../flutter/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.mm
FILE: ../../../flutter/impeller/renderer/backend/metal/compute_pass_mtl.h
FILE: ../../../flutter/impeller/renderer/backend/metal/compute_pass_mtl.mm
FILE: ../../../flutter/impeller/renderer/backend/metal/compute_pipeline_mtl.h
@@ -8275,8 +8273,6 @@
FILE: ../../../flutter/impeller/renderer/backend/vulkan/android_hardware_buffer_texture_source_vk.h
FILE: ../../../flutter/impeller/renderer/backend/vulkan/barrier_vk.cc
FILE: ../../../flutter/impeller/renderer/backend/vulkan/barrier_vk.h
-FILE: ../../../flutter/impeller/renderer/backend/vulkan/binding_helpers_vk.cc
-FILE: ../../../flutter/impeller/renderer/backend/vulkan/binding_helpers_vk.h
FILE: ../../../flutter/impeller/renderer/backend/vulkan/blit_command_vk.cc
FILE: ../../../flutter/impeller/renderer/backend/vulkan/blit_command_vk.h
FILE: ../../../flutter/impeller/renderer/backend/vulkan/blit_pass_vk.cc
@@ -8309,8 +8305,6 @@
FILE: ../../../flutter/impeller/renderer/backend/vulkan/gpu_tracer_vk.cc
FILE: ../../../flutter/impeller/renderer/backend/vulkan/gpu_tracer_vk.h
FILE: ../../../flutter/impeller/renderer/backend/vulkan/limits_vk.h
-FILE: ../../../flutter/impeller/renderer/backend/vulkan/pass_bindings_cache.cc
-FILE: ../../../flutter/impeller/renderer/backend/vulkan/pass_bindings_cache.h
FILE: ../../../flutter/impeller/renderer/backend/vulkan/pipeline_cache_vk.cc
FILE: ../../../flutter/impeller/renderer/backend/vulkan/pipeline_cache_vk.h
FILE: ../../../flutter/impeller/renderer/backend/vulkan/pipeline_library_vk.cc
@@ -8364,8 +8358,6 @@
FILE: ../../../flutter/impeller/renderer/command.h
FILE: ../../../flutter/impeller/renderer/command_buffer.cc
FILE: ../../../flutter/impeller/renderer/command_buffer.h
-FILE: ../../../flutter/impeller/renderer/compute_command.cc
-FILE: ../../../flutter/impeller/renderer/compute_command.h
FILE: ../../../flutter/impeller/renderer/compute_pass.cc
FILE: ../../../flutter/impeller/renderer/compute_pass.h
FILE: ../../../flutter/impeller/renderer/compute_pipeline_builder.cc
diff --git a/impeller/entity/geometry/point_field_geometry.cc b/impeller/entity/geometry/point_field_geometry.cc
index 27369da..0e7bf69 100644
--- a/impeller/entity/geometry/point_field_geometry.cc
+++ b/impeller/entity/geometry/point_field_geometry.cc
@@ -5,7 +5,6 @@
#include "impeller/entity/geometry/point_field_geometry.h"
#include "impeller/renderer/command_buffer.h"
-#include "impeller/renderer/compute_command.h"
namespace impeller {
@@ -168,9 +167,9 @@
BufferView output;
{
using PS = PointsComputeShader;
- ComputeCommand cmd;
- DEBUG_COMMAND_INFO(cmd, "Points Geometry");
- cmd.pipeline = renderer.GetPointComputePipeline();
+
+ compute_pass->SetPipeline(renderer.GetPointComputePipeline());
+ compute_pass->SetCommandLabel("Points Geometry");
PS::FrameInfo frame_info;
frame_info.count = points_.size();
@@ -180,11 +179,11 @@
frame_info.points_per_circle = points_per_circle;
frame_info.divisions_per_circle = vertices_per_geom;
- PS::BindFrameInfo(cmd, host_buffer.EmplaceUniform(frame_info));
- PS::BindGeometryData(cmd, geometry_buffer);
- PS::BindPointData(cmd, points_data);
+ PS::BindFrameInfo(*compute_pass, host_buffer.EmplaceUniform(frame_info));
+ PS::BindGeometryData(*compute_pass, geometry_buffer);
+ PS::BindPointData(*compute_pass, points_data);
- if (!compute_pass->AddCommand(std::move(cmd))) {
+ if (!compute_pass->Compute(ISize(total, 1)).ok()) {
return {};
}
output = geometry_buffer;
@@ -201,9 +200,8 @@
using UV = UvComputeShader;
- ComputeCommand cmd;
- DEBUG_COMMAND_INFO(cmd, "UV Geometry");
- cmd.pipeline = renderer.GetUvComputePipeline();
+ compute_pass->SetCommandLabel("UV Geometry");
+ compute_pass->SetPipeline(renderer.GetUvComputePipeline());
UV::FrameInfo frame_info;
frame_info.count = total;
@@ -211,19 +209,16 @@
frame_info.texture_origin = {0, 0};
frame_info.texture_size = Vector2(texture_coverage.value().GetSize());
- UV::BindFrameInfo(cmd, host_buffer.EmplaceUniform(frame_info));
- UV::BindGeometryData(cmd, geometry_buffer);
- UV::BindGeometryUVData(cmd, geometry_uv_buffer);
+ UV::BindFrameInfo(*compute_pass, host_buffer.EmplaceUniform(frame_info));
+ UV::BindGeometryData(*compute_pass, geometry_buffer);
+ UV::BindGeometryUVData(*compute_pass, geometry_uv_buffer);
- if (!compute_pass->AddCommand(std::move(cmd))) {
+ if (!compute_pass->Compute(ISize(total, 1)).ok()) {
return {};
}
output = geometry_uv_buffer;
}
- compute_pass->SetGridSize(ISize(total, 1));
- compute_pass->SetThreadGroupSize(ISize(total, 1));
-
if (!compute_pass->EncodeCommands() || !cmd_buffer->SubmitCommands()) {
return {};
}
diff --git a/impeller/renderer/BUILD.gn b/impeller/renderer/BUILD.gn
index 39e98fd..1456b65 100644
--- a/impeller/renderer/BUILD.gn
+++ b/impeller/renderer/BUILD.gn
@@ -55,8 +55,6 @@
"command.h",
"command_buffer.cc",
"command_buffer.h",
- "compute_command.cc",
- "compute_command.h",
"compute_pass.cc",
"compute_pass.h",
"compute_pipeline_builder.cc",
diff --git a/impeller/renderer/backend/metal/BUILD.gn b/impeller/renderer/backend/metal/BUILD.gn
index 4e64976..2c9e0f2 100644
--- a/impeller/renderer/backend/metal/BUILD.gn
+++ b/impeller/renderer/backend/metal/BUILD.gn
@@ -14,6 +14,8 @@
"blit_pass_mtl.mm",
"command_buffer_mtl.h",
"command_buffer_mtl.mm",
+ "compute_pass_bindings_cache_mtl.h",
+ "compute_pass_bindings_cache_mtl.mm",
"compute_pass_mtl.h",
"compute_pass_mtl.mm",
"compute_pipeline_mtl.h",
diff --git a/impeller/renderer/backend/metal/command_buffer_mtl.mm b/impeller/renderer/backend/metal/command_buffer_mtl.mm
index d9c43d1..5909426 100644
--- a/impeller/renderer/backend/metal/command_buffer_mtl.mm
+++ b/impeller/renderer/backend/metal/command_buffer_mtl.mm
@@ -220,9 +220,13 @@
if (!buffer_) {
return nullptr;
}
+ auto context = context_.lock();
+ if (!context) {
+ return nullptr;
+ }
auto pass =
- std::shared_ptr<ComputePassMTL>(new ComputePassMTL(context_, buffer_));
+ std::shared_ptr<ComputePassMTL>(new ComputePassMTL(context, buffer_));
if (!pass->IsValid()) {
return nullptr;
}
diff --git a/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.h b/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.h
new file mode 100644
index 0000000..eff7557
--- /dev/null
+++ b/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.h
@@ -0,0 +1,62 @@
+// Copyright 2013 The Flutter Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+
+#ifndef FLUTTER_IMPELLER_RENDERER_BACKEND_METAL_COMPUTE_PASS_BINDINGS_CACHE_MTL_H_
+#define FLUTTER_IMPELLER_RENDERER_BACKEND_METAL_COMPUTE_PASS_BINDINGS_CACHE_MTL_H_
+
+#include <Metal/Metal.h>
+
+#include "flutter/fml/macros.h"
+#include "impeller/renderer/compute_pass.h"
+#include "impeller/renderer/pipeline_descriptor.h"
+
+namespace impeller {
+
+//-----------------------------------------------------------------------------
+/// @brief Ensures that bindings on the pass are not redundantly set or
+/// updated. Avoids making the driver do additional checks and makes
+/// the frame insights during profiling and instrumentation not
+/// complain about the same.
+///
+/// There should be no change to rendering if this caching was
+/// absent.
+///
+struct ComputePassBindingsCacheMTL {
+ explicit ComputePassBindingsCacheMTL() {}
+
+ ComputePassBindingsCacheMTL(const ComputePassBindingsCacheMTL&) = delete;
+
+ ComputePassBindingsCacheMTL(ComputePassBindingsCacheMTL&&) = delete;
+
+ void SetComputePipelineState(id<MTLComputePipelineState> pipeline);
+
+ id<MTLComputePipelineState> GetPipeline() const;
+
+ void SetEncoder(id<MTLComputeCommandEncoder> encoder);
+
+ void SetBuffer(uint64_t index, uint64_t offset, id<MTLBuffer> buffer);
+
+ void SetTexture(uint64_t index, id<MTLTexture> texture);
+
+ void SetSampler(uint64_t index, id<MTLSamplerState> sampler);
+
+ private:
+ struct BufferOffsetPair {
+ id<MTLBuffer> buffer = nullptr;
+ size_t offset = 0u;
+ };
+ using BufferMap = std::map<uint64_t, BufferOffsetPair>;
+ using TextureMap = std::map<uint64_t, id<MTLTexture>>;
+ using SamplerMap = std::map<uint64_t, id<MTLSamplerState>>;
+
+ id<MTLComputeCommandEncoder> encoder_;
+ id<MTLComputePipelineState> pipeline_ = nullptr;
+ BufferMap buffers_;
+ TextureMap textures_;
+ SamplerMap samplers_;
+};
+
+} // namespace impeller
+
+#endif // FLUTTER_IMPELLER_RENDERER_BACKEND_METAL_COMPUTE_PASS_BINDINGS_CACHE_MTL_H_
diff --git a/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.mm b/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.mm
new file mode 100644
index 0000000..50e91e7
--- /dev/null
+++ b/impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.mm
@@ -0,0 +1,73 @@
+// Copyright 2013 The Flutter Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+
+#include "impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.h"
+
+namespace impeller {
+
+void ComputePassBindingsCacheMTL::SetComputePipelineState(
+ id<MTLComputePipelineState> pipeline) {
+ if (pipeline == pipeline_) {
+ return;
+ }
+ pipeline_ = pipeline;
+ [encoder_ setComputePipelineState:pipeline_];
+}
+
+id<MTLComputePipelineState> ComputePassBindingsCacheMTL::GetPipeline() const {
+ return pipeline_;
+}
+
+void ComputePassBindingsCacheMTL::SetEncoder(
+ id<MTLComputeCommandEncoder> encoder) {
+ encoder_ = encoder;
+}
+
+void ComputePassBindingsCacheMTL::SetBuffer(uint64_t index,
+ uint64_t offset,
+ id<MTLBuffer> buffer) {
+ auto found = buffers_.find(index);
+ if (found != buffers_.end() && found->second.buffer == buffer) {
+ // The right buffer is bound. Check if its offset needs to be updated.
+ if (found->second.offset == offset) {
+ // Buffer and its offset is identical. Nothing to do.
+ return;
+ }
+
+ // Only the offset needs to be updated.
+ found->second.offset = offset;
+
+ [encoder_ setBufferOffset:offset atIndex:index];
+ return;
+ }
+
+ buffers_[index] = {buffer, static_cast<size_t>(offset)};
+ [encoder_ setBuffer:buffer offset:offset atIndex:index];
+}
+
+void ComputePassBindingsCacheMTL::SetTexture(uint64_t index,
+ id<MTLTexture> texture) {
+ auto found = textures_.find(index);
+ if (found != textures_.end() && found->second == texture) {
+ // Already bound.
+ return;
+ }
+ textures_[index] = texture;
+ [encoder_ setTexture:texture atIndex:index];
+ return;
+}
+
+void ComputePassBindingsCacheMTL::SetSampler(uint64_t index,
+ id<MTLSamplerState> sampler) {
+ auto found = samplers_.find(index);
+ if (found != samplers_.end() && found->second == sampler) {
+ // Already bound.
+ return;
+ }
+ samplers_[index] = sampler;
+ [encoder_ setSamplerState:sampler atIndex:index];
+ return;
+}
+
+} // namespace impeller
diff --git a/impeller/renderer/backend/metal/compute_pass_mtl.h b/impeller/renderer/backend/metal/compute_pass_mtl.h
index 8127496..91a45ea 100644
--- a/impeller/renderer/backend/metal/compute_pass_mtl.h
+++ b/impeller/renderer/backend/metal/compute_pass_mtl.h
@@ -8,7 +8,9 @@
#include <Metal/Metal.h>
#include "flutter/fml/macros.h"
+#include "impeller/renderer/backend/metal/compute_pass_bindings_cache_mtl.h"
#include "impeller/renderer/compute_pass.h"
+#include "impeller/renderer/pipeline_descriptor.h"
namespace impeller {
@@ -21,27 +23,48 @@
friend class CommandBufferMTL;
id<MTLCommandBuffer> buffer_ = nil;
- std::string label_;
+ id<MTLComputeCommandEncoder> encoder_ = nil;
+ ComputePassBindingsCacheMTL pass_bindings_cache_ =
+ ComputePassBindingsCacheMTL();
bool is_valid_ = false;
+ bool has_label_ = false;
- ComputePassMTL(std::weak_ptr<const Context> context,
+ ComputePassMTL(std::shared_ptr<const Context> context,
id<MTLCommandBuffer> buffer);
// |ComputePass|
bool IsValid() const override;
// |ComputePass|
+ fml::Status Compute(const ISize& grid_size) override;
+
+ // |ComputePass|
+ void SetCommandLabel(std::string_view label) override;
+
+ // |ComputePass|
void OnSetLabel(const std::string& label) override;
// |ComputePass|
- bool OnEncodeCommands(const Context& context,
- const ISize& grid_size,
- const ISize& thread_group_size) const override;
+ void SetPipeline(const std::shared_ptr<Pipeline<ComputePipelineDescriptor>>&
+ pipeline) override;
- bool EncodeCommands(const std::shared_ptr<Allocator>& allocator,
- id<MTLComputeCommandEncoder> pass,
- const ISize& grid_size,
- const ISize& thread_group_size) const;
+ // |ComputePass|
+ bool BindResource(ShaderStage stage,
+ DescriptorType type,
+ const ShaderUniformSlot& slot,
+ const ShaderMetadata& metadata,
+ BufferView view) override;
+
+ // |ComputePass|
+ bool BindResource(ShaderStage stage,
+ DescriptorType type,
+ const SampledImageSlot& slot,
+ const ShaderMetadata& metadata,
+ std::shared_ptr<const Texture> texture,
+ std::shared_ptr<const Sampler> sampler) override;
+
+ // |ComputePass|
+ bool EncodeCommands() const override;
ComputePassMTL(const ComputePassMTL&) = delete;
diff --git a/impeller/renderer/backend/metal/compute_pass_mtl.mm b/impeller/renderer/backend/metal/compute_pass_mtl.mm
index 5abab34..e322819 100644
--- a/impeller/renderer/backend/metal/compute_pass_mtl.mm
+++ b/impeller/renderer/backend/metal/compute_pass_mtl.mm
@@ -6,13 +6,14 @@
#include <Metal/Metal.h>
#include <memory>
-#include <variant>
#include "flutter/fml/backtrace.h"
#include "flutter/fml/closure.h"
#include "flutter/fml/logging.h"
#include "flutter/fml/trace_event.h"
+#include "fml/status.h"
#include "impeller/base/backend_cast.h"
+#include "impeller/core/device_buffer.h"
#include "impeller/core/formats.h"
#include "impeller/core/host_buffer.h"
#include "impeller/core/shader_types.h"
@@ -22,16 +23,20 @@
#include "impeller/renderer/backend/metal/sampler_mtl.h"
#include "impeller/renderer/backend/metal/texture_mtl.h"
#include "impeller/renderer/command.h"
-#include "impeller/renderer/compute_command.h"
namespace impeller {
-ComputePassMTL::ComputePassMTL(std::weak_ptr<const Context> context,
+ComputePassMTL::ComputePassMTL(std::shared_ptr<const Context> context,
id<MTLCommandBuffer> buffer)
: ComputePass(std::move(context)), buffer_(buffer) {
if (!buffer_) {
return;
}
+ encoder_ = [buffer_ computeCommandEncoder];
+ if (!encoder_) {
+ return;
+ }
+ pass_bindings_cache_.SetEncoder(encoder_);
is_valid_ = true;
}
@@ -42,231 +47,114 @@
}
void ComputePassMTL::OnSetLabel(const std::string& label) {
+#ifdef IMPELLER_DEBUG
if (label.empty()) {
return;
}
- label_ = label;
+ [encoder_ setLabel:@(label.c_str())];
+#endif // IMPELLER_DEBUG
}
-bool ComputePassMTL::OnEncodeCommands(const Context& context,
- const ISize& grid_size,
- const ISize& thread_group_size) const {
- TRACE_EVENT0("impeller", "ComputePassMTL::EncodeCommands");
- if (!IsValid()) {
- return false;
- }
-
- FML_DCHECK(!grid_size.IsEmpty() && !thread_group_size.IsEmpty());
-
- // TODO(dnfield): Support non-serial dispatch type on higher iOS versions.
- auto compute_command_encoder = [buffer_ computeCommandEncoder];
-
- if (!compute_command_encoder) {
- return false;
- }
-
- if (!label_.empty()) {
- [compute_command_encoder setLabel:@(label_.c_str())];
- }
-
- // Success or failure, the pass must end. The buffer can only process one pass
- // at a time.
- fml::ScopedCleanupClosure auto_end(
- [compute_command_encoder]() { [compute_command_encoder endEncoding]; });
-
- return EncodeCommands(context.GetResourceAllocator(), compute_command_encoder,
- grid_size, thread_group_size);
+void ComputePassMTL::SetCommandLabel(std::string_view label) {
+ has_label_ = true;
+ [encoder_ pushDebugGroup:@(label.data())];
}
-//-----------------------------------------------------------------------------
-/// @brief Ensures that bindings on the pass are not redundantly set or
-/// updated. Avoids making the driver do additional checks and makes
-/// the frame insights during profiling and instrumentation not
-/// complain about the same.
-///
-/// There should be no change to rendering if this caching was
-/// absent.
-///
-struct ComputePassBindingsCache {
- explicit ComputePassBindingsCache(id<MTLComputeCommandEncoder> encoder)
- : encoder_(encoder) {}
+// |ComputePass|
+void ComputePassMTL::SetPipeline(
+ const std::shared_ptr<Pipeline<ComputePipelineDescriptor>>& pipeline) {
+ pass_bindings_cache_.SetComputePipelineState(
+ ComputePipelineMTL::Cast(*pipeline).GetMTLComputePipelineState());
+}
- ComputePassBindingsCache(const ComputePassBindingsCache&) = delete;
-
- ComputePassBindingsCache(ComputePassBindingsCache&&) = delete;
-
- void SetComputePipelineState(id<MTLComputePipelineState> pipeline) {
- if (pipeline == pipeline_) {
- return;
- }
- pipeline_ = pipeline;
- [encoder_ setComputePipelineState:pipeline_];
- }
-
- id<MTLComputePipelineState> GetPipeline() const { return pipeline_; }
-
- void SetBuffer(uint64_t index, uint64_t offset, id<MTLBuffer> buffer) {
- auto found = buffers_.find(index);
- if (found != buffers_.end() && found->second.buffer == buffer) {
- // The right buffer is bound. Check if its offset needs to be updated.
- if (found->second.offset == offset) {
- // Buffer and its offset is identical. Nothing to do.
- return;
- }
-
- // Only the offset needs to be updated.
- found->second.offset = offset;
-
- [encoder_ setBufferOffset:offset atIndex:index];
- return;
- }
-
- buffers_[index] = {buffer, static_cast<size_t>(offset)};
- [encoder_ setBuffer:buffer offset:offset atIndex:index];
- }
-
- void SetTexture(uint64_t index, id<MTLTexture> texture) {
- auto found = textures_.find(index);
- if (found != textures_.end() && found->second == texture) {
- // Already bound.
- return;
- }
- textures_[index] = texture;
- [encoder_ setTexture:texture atIndex:index];
- return;
- }
-
- void SetSampler(uint64_t index, id<MTLSamplerState> sampler) {
- auto found = samplers_.find(index);
- if (found != samplers_.end() && found->second == sampler) {
- // Already bound.
- return;
- }
- samplers_[index] = sampler;
- [encoder_ setSamplerState:sampler atIndex:index];
- return;
- }
-
- private:
- struct BufferOffsetPair {
- id<MTLBuffer> buffer = nullptr;
- size_t offset = 0u;
- };
- using BufferMap = std::map<uint64_t, BufferOffsetPair>;
- using TextureMap = std::map<uint64_t, id<MTLTexture>>;
- using SamplerMap = std::map<uint64_t, id<MTLSamplerState>>;
-
- const id<MTLComputeCommandEncoder> encoder_;
- id<MTLComputePipelineState> pipeline_ = nullptr;
- BufferMap buffers_;
- TextureMap textures_;
- SamplerMap samplers_;
-};
-
-static bool Bind(ComputePassBindingsCache& pass,
- Allocator& allocator,
- size_t bind_index,
- const BufferView& view) {
+// |ComputePass|
+bool ComputePassMTL::BindResource(ShaderStage stage,
+ DescriptorType type,
+ const ShaderUniformSlot& slot,
+ const ShaderMetadata& metadata,
+ BufferView view) {
if (!view.buffer) {
return false;
}
- auto device_buffer = view.buffer;
+ const std::shared_ptr<const DeviceBuffer>& device_buffer = view.buffer;
if (!device_buffer) {
return false;
}
- auto buffer = DeviceBufferMTL::Cast(*device_buffer).GetMTLBuffer();
+ id<MTLBuffer> buffer = DeviceBufferMTL::Cast(*device_buffer).GetMTLBuffer();
// The Metal call is a void return and we don't want to make it on nil.
if (!buffer) {
return false;
}
- pass.SetBuffer(bind_index, view.range.offset, buffer);
+ pass_bindings_cache_.SetBuffer(slot.ext_res_0, view.range.offset, buffer);
return true;
}
-static bool Bind(ComputePassBindingsCache& pass,
- size_t bind_index,
- const Sampler& sampler,
- const Texture& texture) {
- if (!sampler.IsValid() || !texture.IsValid()) {
+// |ComputePass|
+bool ComputePassMTL::BindResource(ShaderStage stage,
+ DescriptorType type,
+ const SampledImageSlot& slot,
+ const ShaderMetadata& metadata,
+ std::shared_ptr<const Texture> texture,
+ std::shared_ptr<const Sampler> sampler) {
+ if (!sampler->IsValid() || !texture->IsValid()) {
return false;
}
- pass.SetTexture(bind_index, TextureMTL::Cast(texture).GetMTLTexture());
- pass.SetSampler(bind_index, SamplerMTL::Cast(sampler).GetMTLSamplerState());
+ pass_bindings_cache_.SetTexture(slot.texture_index,
+ TextureMTL::Cast(*texture).GetMTLTexture());
+ pass_bindings_cache_.SetSampler(
+ slot.texture_index, SamplerMTL::Cast(*sampler).GetMTLSamplerState());
return true;
}
-bool ComputePassMTL::EncodeCommands(const std::shared_ptr<Allocator>& allocator,
- id<MTLComputeCommandEncoder> encoder,
- const ISize& grid_size,
- const ISize& thread_group_size) const {
- if (grid_size.width == 0 || grid_size.height == 0) {
- return true;
+fml::Status ComputePassMTL::Compute(const ISize& grid_size) {
+ if (grid_size.IsEmpty()) {
+ return fml::Status(fml::StatusCode::kUnknown,
+ "Invalid grid size for compute command.");
+ }
+ // TODO(dnfield): use feature detection to support non-uniform threadgroup
+ // sizes.
+ // https://github.com/flutter/flutter/issues/110619
+ auto width = grid_size.width;
+ auto height = grid_size.height;
+
+ auto max_total_threads_per_threadgroup = static_cast<int64_t>(
+ pass_bindings_cache_.GetPipeline().maxTotalThreadsPerThreadgroup);
+
+ // Special case for linear processing.
+ if (height == 1) {
+ int64_t thread_groups = std::max(
+ static_cast<int64_t>(
+ std::ceil(width * 1.0 / max_total_threads_per_threadgroup * 1.0)),
+ 1LL);
+ [encoder_
+ dispatchThreadgroups:MTLSizeMake(thread_groups, 1, 1)
+ threadsPerThreadgroup:MTLSizeMake(max_total_threads_per_threadgroup, 1,
+ 1)];
+ } else {
+ while (width * height > max_total_threads_per_threadgroup) {
+ width = std::max(1LL, width / 2);
+ height = std::max(1LL, height / 2);
+ }
+
+ auto size = MTLSizeMake(width, height, 1);
+ [encoder_ dispatchThreadgroups:size threadsPerThreadgroup:size];
}
- ComputePassBindingsCache pass_bindings(encoder);
-
- fml::closure pop_debug_marker = [encoder]() { [encoder popDebugGroup]; };
- for (const ComputeCommand& command : commands_) {
#ifdef IMPELLER_DEBUG
- fml::ScopedCleanupClosure auto_pop_debug_marker(pop_debug_marker);
- if (!command.label.empty()) {
- [encoder pushDebugGroup:@(command.label.c_str())];
- } else {
- auto_pop_debug_marker.Release();
- }
-#endif
-
- pass_bindings.SetComputePipelineState(
- ComputePipelineMTL::Cast(*command.pipeline)
- .GetMTLComputePipelineState());
-
- for (const BufferAndUniformSlot& buffer : command.bindings.buffers) {
- if (!Bind(pass_bindings, *allocator, buffer.slot.ext_res_0,
- buffer.view.resource)) {
- return false;
- }
- }
-
- for (const TextureAndSampler& data : command.bindings.sampled_images) {
- if (!Bind(pass_bindings, data.slot.texture_index, *data.sampler,
- *data.texture.resource)) {
- return false;
- }
- }
-
- // TODO(dnfield): use feature detection to support non-uniform threadgroup
- // sizes.
- // https://github.com/flutter/flutter/issues/110619
- auto width = grid_size.width;
- auto height = grid_size.height;
-
- auto maxTotalThreadsPerThreadgroup = static_cast<int64_t>(
- pass_bindings.GetPipeline().maxTotalThreadsPerThreadgroup);
-
- // Special case for linear processing.
- if (height == 1) {
- int64_t threadGroups = std::max(
- static_cast<int64_t>(
- std::ceil(width * 1.0 / maxTotalThreadsPerThreadgroup * 1.0)),
- 1LL);
- [encoder dispatchThreadgroups:MTLSizeMake(threadGroups, 1, 1)
- threadsPerThreadgroup:MTLSizeMake(maxTotalThreadsPerThreadgroup,
- 1, 1)];
- } else {
- while (width * height > maxTotalThreadsPerThreadgroup) {
- width = std::max(1LL, width / 2);
- height = std::max(1LL, height / 2);
- }
-
- auto size = MTLSizeMake(width, height, 1);
- [encoder dispatchThreadgroups:size threadsPerThreadgroup:size];
- }
+ if (has_label_) {
+ [encoder_ popDebugGroup];
}
+ has_label_ = false;
+#endif // IMPELLER_DEBUG
+ return fml::Status();
+}
+bool ComputePassMTL::EncodeCommands() const {
+ [encoder_ endEncoding];
return true;
}
diff --git a/impeller/renderer/backend/vulkan/BUILD.gn b/impeller/renderer/backend/vulkan/BUILD.gn
index 7a7822c..63cf211 100644
--- a/impeller/renderer/backend/vulkan/BUILD.gn
+++ b/impeller/renderer/backend/vulkan/BUILD.gn
@@ -14,7 +14,6 @@
"context_vk_unittests.cc",
"descriptor_pool_vk_unittests.cc",
"fence_waiter_vk_unittests.cc",
- "pass_bindings_cache_unittests.cc",
"resource_manager_vk_unittests.cc",
"test/gpu_tracer_unittests.cc",
"test/mock_vulkan.cc",
@@ -35,8 +34,6 @@
"android_hardware_buffer_texture_source_vk.h",
"barrier_vk.cc",
"barrier_vk.h",
- "binding_helpers_vk.cc",
- "binding_helpers_vk.h",
"blit_command_vk.cc",
"blit_command_vk.h",
"blit_pass_vk.cc",
@@ -68,8 +65,6 @@
"gpu_tracer_vk.cc",
"gpu_tracer_vk.h",
"limits_vk.h",
- "pass_bindings_cache.cc",
- "pass_bindings_cache.h",
"pipeline_cache_vk.cc",
"pipeline_cache_vk.h",
"pipeline_library_vk.cc",
diff --git a/impeller/renderer/backend/vulkan/binding_helpers_vk.cc b/impeller/renderer/backend/vulkan/binding_helpers_vk.cc
deleted file mode 100644
index 92b3031..0000000
--- a/impeller/renderer/backend/vulkan/binding_helpers_vk.cc
+++ /dev/null
@@ -1,162 +0,0 @@
-// Copyright 2013 The Flutter Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-
-#include "impeller/renderer/backend/vulkan/binding_helpers_vk.h"
-#include "fml/status.h"
-#include "impeller/core/allocator.h"
-#include "impeller/core/device_buffer.h"
-#include "impeller/core/shader_types.h"
-#include "impeller/renderer/backend/vulkan/command_buffer_vk.h"
-#include "impeller/renderer/backend/vulkan/command_encoder_vk.h"
-#include "impeller/renderer/backend/vulkan/command_pool_vk.h"
-#include "impeller/renderer/backend/vulkan/compute_pipeline_vk.h"
-#include "impeller/renderer/backend/vulkan/context_vk.h"
-#include "impeller/renderer/backend/vulkan/sampler_vk.h"
-#include "impeller/renderer/backend/vulkan/texture_vk.h"
-#include "impeller/renderer/command.h"
-#include "impeller/renderer/compute_command.h"
-#include "vulkan/vulkan_core.h"
-
-namespace impeller {
-
-static bool BindImages(
- const Bindings& bindings,
- Allocator& allocator,
- const std::shared_ptr<CommandEncoderVK>& encoder,
- vk::DescriptorSet& vk_desc_set,
- std::array<vk::DescriptorImageInfo, kMaxBindings>& image_workspace,
- size_t& image_offset,
- std::array<vk::WriteDescriptorSet, kMaxBindings + kMaxBindings>&
- write_workspace,
- size_t& write_offset) {
- for (const TextureAndSampler& data : bindings.sampled_images) {
- const std::shared_ptr<const Texture>& texture = data.texture.resource;
- const TextureVK& texture_vk = TextureVK::Cast(*texture);
- const SamplerVK& sampler = SamplerVK::Cast(*data.sampler);
-
- if (!encoder->Track(texture) ||
- !encoder->Track(sampler.GetSharedSampler())) {
- return false;
- }
-
- const SampledImageSlot& slot = data.slot;
-
- vk::DescriptorImageInfo image_info;
- image_info.imageLayout = vk::ImageLayout::eShaderReadOnlyOptimal;
- image_info.sampler = sampler.GetSampler();
- image_info.imageView = texture_vk.GetImageView();
- image_workspace[image_offset++] = image_info;
-
- vk::WriteDescriptorSet write_set;
- write_set.dstSet = vk_desc_set;
- write_set.dstBinding = slot.binding;
- write_set.descriptorCount = 1u;
- write_set.descriptorType = vk::DescriptorType::eCombinedImageSampler;
- write_set.pImageInfo = &image_workspace[image_offset - 1];
-
- write_workspace[write_offset++] = write_set;
- }
-
- return true;
-};
-
-static bool BindBuffers(
- const Bindings& bindings,
- Allocator& allocator,
- const std::shared_ptr<CommandEncoderVK>& encoder,
- vk::DescriptorSet& vk_desc_set,
- const std::vector<DescriptorSetLayout>& desc_set,
- std::array<vk::DescriptorBufferInfo, kMaxBindings>& buffer_workspace,
- size_t& buffer_offset,
- std::array<vk::WriteDescriptorSet, kMaxBindings + kMaxBindings>&
- write_workspace,
- size_t& write_offset) {
- for (const BufferAndUniformSlot& data : bindings.buffers) {
- const std::shared_ptr<const DeviceBuffer>& device_buffer =
- data.view.resource.buffer;
-
- auto buffer = DeviceBufferVK::Cast(*device_buffer).GetBuffer();
- if (!buffer) {
- return false;
- }
-
- if (!encoder->Track(device_buffer)) {
- return false;
- }
-
- uint32_t offset = data.view.resource.range.offset;
-
- vk::DescriptorBufferInfo buffer_info;
- buffer_info.buffer = buffer;
- buffer_info.offset = offset;
- buffer_info.range = data.view.resource.range.length;
- buffer_workspace[buffer_offset++] = buffer_info;
-
- // TODO(jonahwilliams): remove this part by storing more data in
- // ShaderUniformSlot.
- const ShaderUniformSlot& uniform = data.slot;
- auto layout_it =
- std::find_if(desc_set.begin(), desc_set.end(),
- [&uniform](const DescriptorSetLayout& layout) {
- return layout.binding == uniform.binding;
- });
- if (layout_it == desc_set.end()) {
- VALIDATION_LOG << "Failed to get descriptor set layout for binding "
- << uniform.binding;
- return false;
- }
- auto layout = *layout_it;
-
- vk::WriteDescriptorSet write_set;
- write_set.dstSet = vk_desc_set;
- write_set.dstBinding = uniform.binding;
- write_set.descriptorCount = 1u;
- write_set.descriptorType = ToVKDescriptorType(layout.descriptor_type);
- write_set.pBufferInfo = &buffer_workspace[buffer_offset - 1];
-
- write_workspace[write_offset++] = write_set;
- }
- return true;
-}
-
-fml::StatusOr<vk::DescriptorSet> AllocateAndBindDescriptorSets(
- const ContextVK& context,
- const std::shared_ptr<CommandEncoderVK>& encoder,
- Allocator& allocator,
- const ComputeCommand& command,
- std::array<vk::DescriptorImageInfo, kMaxBindings>& image_workspace,
- std::array<vk::DescriptorBufferInfo, kMaxBindings>& buffer_workspace,
- std::array<vk::WriteDescriptorSet, kMaxBindings + kMaxBindings>&
- write_workspace) {
- auto descriptor_result = encoder->AllocateDescriptorSets(
- ComputePipelineVK::Cast(*command.pipeline).GetDescriptorSetLayout(),
- context);
- if (!descriptor_result.ok()) {
- return descriptor_result.status();
- }
- auto descriptor_set = descriptor_result.value();
-
- size_t buffer_offset = 0u;
- size_t image_offset = 0u;
- size_t write_offset = 0u;
-
- auto& pipeline_descriptor = command.pipeline->GetDescriptor();
- auto& desc_set = pipeline_descriptor.GetDescriptorSetLayouts();
-
- if (!BindBuffers(command.bindings, allocator, encoder, descriptor_set,
- desc_set, buffer_workspace, buffer_offset, write_workspace,
- write_offset) ||
- !BindImages(command.bindings, allocator, encoder, descriptor_set,
- image_workspace, image_offset, write_workspace,
- write_offset)) {
- return fml::Status(fml::StatusCode::kUnknown,
- "Failed to bind texture or buffer.");
- }
- context.GetDevice().updateDescriptorSets(write_offset, write_workspace.data(),
- 0u, {});
-
- return descriptor_set;
-}
-
-} // namespace impeller
diff --git a/impeller/renderer/backend/vulkan/binding_helpers_vk.h b/impeller/renderer/backend/vulkan/binding_helpers_vk.h
deleted file mode 100644
index be35567..0000000
--- a/impeller/renderer/backend/vulkan/binding_helpers_vk.h
+++ /dev/null
@@ -1,32 +0,0 @@
-// Copyright 2013 The Flutter Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-
-#ifndef FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_BINDING_HELPERS_VK_H_
-#define FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_BINDING_HELPERS_VK_H_
-
-#include "fml/status_or.h"
-#include "impeller/renderer/backend/vulkan/context_vk.h"
-#include "impeller/renderer/backend/vulkan/texture_vk.h"
-#include "impeller/renderer/command.h"
-#include "impeller/renderer/compute_command.h"
-
-namespace impeller {
-
-// Limit on the total number of buffer and image bindings that allow the Vulkan
-// backend to avoid dynamic heap allocations.
-static constexpr size_t kMaxBindings = 32;
-
-fml::StatusOr<vk::DescriptorSet> AllocateAndBindDescriptorSets(
- const ContextVK& context,
- const std::shared_ptr<CommandEncoderVK>& encoder,
- Allocator& allocator,
- const ComputeCommand& command,
- std::array<vk::DescriptorImageInfo, kMaxBindings>& image_workspace,
- std::array<vk::DescriptorBufferInfo, kMaxBindings>& buffer_workspace,
- std::array<vk::WriteDescriptorSet, kMaxBindings + kMaxBindings>&
- write_workspace);
-
-} // namespace impeller
-
-#endif // FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_BINDING_HELPERS_VK_H_
diff --git a/impeller/renderer/backend/vulkan/command_buffer_vk.cc b/impeller/renderer/backend/vulkan/command_buffer_vk.cc
index 6b7b14f..c63dbec 100644
--- a/impeller/renderer/backend/vulkan/command_buffer_vk.cc
+++ b/impeller/renderer/backend/vulkan/command_buffer_vk.cc
@@ -100,8 +100,8 @@
return nullptr;
}
auto pass =
- std::shared_ptr<ComputePassVK>(new ComputePassVK(context, //
- weak_from_this() //
+ std::shared_ptr<ComputePassVK>(new ComputePassVK(context, //
+ shared_from_this() //
));
if (!pass->IsValid()) {
return nullptr;
diff --git a/impeller/renderer/backend/vulkan/compute_pass_vk.cc b/impeller/renderer/backend/vulkan/compute_pass_vk.cc
index 4788ec1..8324879 100644
--- a/impeller/renderer/backend/vulkan/compute_pass_vk.cc
+++ b/impeller/renderer/backend/vulkan/compute_pass_vk.cc
@@ -4,19 +4,25 @@
#include "impeller/renderer/backend/vulkan/compute_pass_vk.h"
-#include "flutter/fml/trace_event.h"
-#include "impeller/renderer/backend/vulkan/binding_helpers_vk.h"
#include "impeller/renderer/backend/vulkan/command_buffer_vk.h"
#include "impeller/renderer/backend/vulkan/compute_pipeline_vk.h"
+#include "impeller/renderer/backend/vulkan/formats_vk.h"
+#include "impeller/renderer/backend/vulkan/sampler_vk.h"
#include "impeller/renderer/backend/vulkan/texture_vk.h"
-#include "impeller/renderer/command.h"
+#include "vulkan/vulkan_structs.hpp"
namespace impeller {
-ComputePassVK::ComputePassVK(std::weak_ptr<const Context> context,
- std::weak_ptr<CommandBufferVK> command_buffer)
+ComputePassVK::ComputePassVK(std::shared_ptr<const Context> context,
+ std::shared_ptr<CommandBufferVK> command_buffer)
: ComputePass(std::move(context)),
command_buffer_(std::move(command_buffer)) {
+ // TOOD(dnfield): This should be moved to caps. But for now keeping this
+ // in parallel with Metal.
+ max_wg_size_ = ContextVK::Cast(*context_)
+ .GetPhysicalDevice()
+ .getProperties()
+ .limits.maxComputeWorkGroupSize;
is_valid_ = true;
}
@@ -33,125 +39,182 @@
label_ = label;
}
-static bool UpdateBindingLayouts(const Bindings& bindings,
- const vk::CommandBuffer& buffer) {
- BarrierVK barrier;
- barrier.cmd_buffer = buffer;
- barrier.src_access = vk::AccessFlagBits::eTransferWrite;
- barrier.src_stage = vk::PipelineStageFlagBits::eTransfer;
- barrier.dst_access = vk::AccessFlagBits::eShaderRead;
- barrier.dst_stage = vk::PipelineStageFlagBits::eComputeShader;
-
- barrier.new_layout = vk::ImageLayout::eShaderReadOnlyOptimal;
-
- for (const TextureAndSampler& data : bindings.sampled_images) {
- if (!TextureVK::Cast(*data.texture.resource).SetLayout(barrier)) {
- return false;
- }
- }
- return true;
+// |RenderPass|
+void ComputePassVK::SetCommandLabel(std::string_view label) {
+#ifdef IMPELLER_DEBUG
+ command_buffer_->GetEncoder()->PushDebugGroup(label);
+ has_label_ = true;
+#endif // IMPELLER_DEBUG
}
-static bool UpdateBindingLayouts(const ComputeCommand& command,
- const vk::CommandBuffer& buffer) {
- return UpdateBindingLayouts(command.bindings, buffer);
+// |ComputePass|
+void ComputePassVK::SetPipeline(
+ const std::shared_ptr<Pipeline<ComputePipelineDescriptor>>& pipeline) {
+ const auto& pipeline_vk = ComputePipelineVK::Cast(*pipeline);
+ const vk::CommandBuffer& command_buffer_vk =
+ command_buffer_->GetEncoder()->GetCommandBuffer();
+ command_buffer_vk.bindPipeline(vk::PipelineBindPoint::eCompute,
+ pipeline_vk.GetPipeline());
+ pipeline_layout_ = pipeline_vk.GetPipelineLayout();
+
+ auto descriptor_result =
+ command_buffer_->GetEncoder()->AllocateDescriptorSets(
+ pipeline_vk.GetDescriptorSetLayout(), ContextVK::Cast(*context_));
+ if (!descriptor_result.ok()) {
+ return;
+ }
+ descriptor_set_ = descriptor_result.value();
+ pipeline_valid_ = true;
}
-static bool UpdateBindingLayouts(const std::vector<ComputeCommand>& commands,
- const vk::CommandBuffer& buffer) {
- for (const ComputeCommand& command : commands) {
- if (!UpdateBindingLayouts(command, buffer)) {
- return false;
- }
- }
- return true;
-}
-
-bool ComputePassVK::OnEncodeCommands(const Context& context,
- const ISize& grid_size,
- const ISize& thread_group_size) const {
- TRACE_EVENT0("impeller", "ComputePassVK::EncodeCommands");
- if (!IsValid()) {
- return false;
+// |ComputePass|
+fml::Status ComputePassVK::Compute(const ISize& grid_size) {
+ if (grid_size.IsEmpty() || !pipeline_valid_) {
+ bound_image_offset_ = 0u;
+ bound_buffer_offset_ = 0u;
+ descriptor_write_offset_ = 0u;
+ has_label_ = false;
+ pipeline_valid_ = false;
+ return fml::Status(fml::StatusCode::kCancelled,
+ "Invalid pipeline or empty grid.");
}
- FML_DCHECK(!grid_size.IsEmpty() && !thread_group_size.IsEmpty());
-
- const auto& vk_context = ContextVK::Cast(context);
- auto command_buffer = command_buffer_.lock();
- if (!command_buffer) {
- VALIDATION_LOG << "Command buffer died before commands could be encoded.";
- return false;
- }
- auto encoder = command_buffer->GetEncoder();
- if (!encoder) {
- return false;
+ const ContextVK& context_vk = ContextVK::Cast(*context_);
+ for (auto i = 0u; i < descriptor_write_offset_; i++) {
+ write_workspace_[i].dstSet = descriptor_set_;
}
- fml::ScopedCleanupClosure pop_marker(
- [&encoder]() { encoder->PopDebugGroup(); });
- if (!label_.empty()) {
- encoder->PushDebugGroup(label_.c_str());
+ context_vk.GetDevice().updateDescriptorSets(descriptor_write_offset_,
+ write_workspace_.data(), 0u, {});
+ const vk::CommandBuffer& command_buffer_vk =
+ command_buffer_->GetEncoder()->GetCommandBuffer();
+
+ command_buffer_vk.bindDescriptorSets(
+ vk::PipelineBindPoint::eCompute, // bind point
+ pipeline_layout_, // layout
+ 0, // first set
+ 1, // set count
+ &descriptor_set_, // sets
+ 0, // offset count
+ nullptr // offsets
+ );
+
+ int64_t width = grid_size.width;
+ int64_t height = grid_size.height;
+
+ // Special case for linear processing.
+ if (height == 1) {
+ int64_t minimum = 1;
+ int64_t threadGroups = std::max(
+ static_cast<int64_t>(std::ceil(width * 1.0 / max_wg_size_[0] * 1.0)),
+ minimum);
+ command_buffer_vk.dispatch(threadGroups, 1, 1);
} else {
- pop_marker.Release();
+ while (width > max_wg_size_[0]) {
+ width = std::max(static_cast<int64_t>(1), width / 2);
+ }
+ while (height > max_wg_size_[1]) {
+ height = std::max(static_cast<int64_t>(1), height / 2);
+ }
+ command_buffer_vk.dispatch(width, height, 1);
}
- auto cmd_buffer = encoder->GetCommandBuffer();
- if (!UpdateBindingLayouts(commands_, cmd_buffer)) {
- VALIDATION_LOG << "Could not update binding layouts for compute pass.";
+#ifdef IMPELLER_DEBUG
+ if (has_label_) {
+ command_buffer_->GetEncoder()->PopDebugGroup();
+ }
+ has_label_ = false;
+#endif // IMPELLER_DEBUG
+
+ bound_image_offset_ = 0u;
+ bound_buffer_offset_ = 0u;
+ descriptor_write_offset_ = 0u;
+ has_label_ = false;
+ pipeline_valid_ = false;
+
+ return fml::Status();
+}
+
+// |ResourceBinder|
+bool ComputePassVK::BindResource(ShaderStage stage,
+ DescriptorType type,
+ const ShaderUniformSlot& slot,
+ const ShaderMetadata& metadata,
+ BufferView view) {
+ return BindResource(slot.binding, type, view);
+}
+
+// |ResourceBinder|
+bool ComputePassVK::BindResource(ShaderStage stage,
+ DescriptorType type,
+ const SampledImageSlot& slot,
+ const ShaderMetadata& metadata,
+ std::shared_ptr<const Texture> texture,
+ std::shared_ptr<const Sampler> sampler) {
+ if (bound_image_offset_ >= kMaxBindings) {
return false;
}
- auto& allocator = *context.GetResourceAllocator();
+ const TextureVK& texture_vk = TextureVK::Cast(*texture);
+ const SamplerVK& sampler_vk = SamplerVK::Cast(*sampler);
- TRACE_EVENT0("impeller", "EncodeComputePassCommands");
- for (const auto& command : commands_) {
- auto desc_set_result = AllocateAndBindDescriptorSets(
- vk_context, encoder, allocator, command, image_workspace_,
- buffer_workspace_, write_workspace_);
- if (!desc_set_result.ok()) {
- return false;
- }
- auto desc_set = desc_set_result.value();
-
- const auto& pipeline_vk = ComputePipelineVK::Cast(*command.pipeline);
-
- cmd_buffer.bindPipeline(vk::PipelineBindPoint::eCompute,
- pipeline_vk.GetPipeline());
- cmd_buffer.bindDescriptorSets(
- vk::PipelineBindPoint::eCompute, // bind point
- pipeline_vk.GetPipelineLayout(), // layout
- 0, // first set
- {vk::DescriptorSet{desc_set}}, // sets
- nullptr // offsets
- );
-
- // TOOD(dnfield): This should be moved to caps. But for now keeping this
- // in parallel with Metal.
- auto device_properties = vk_context.GetPhysicalDevice().getProperties();
-
- auto max_wg_size = device_properties.limits.maxComputeWorkGroupSize;
-
- int64_t width = grid_size.width;
- int64_t height = grid_size.height;
-
- // Special case for linear processing.
- if (height == 1) {
- int64_t minimum = 1;
- int64_t threadGroups = std::max(
- static_cast<int64_t>(std::ceil(width * 1.0 / max_wg_size[0] * 1.0)),
- minimum);
- cmd_buffer.dispatch(threadGroups, 1, 1);
- } else {
- while (width > max_wg_size[0]) {
- width = std::max(static_cast<int64_t>(1), width / 2);
- }
- while (height > max_wg_size[1]) {
- height = std::max(static_cast<int64_t>(1), height / 2);
- }
- cmd_buffer.dispatch(width, height, 1);
- }
+ if (!command_buffer_->GetEncoder()->Track(texture) ||
+ !command_buffer_->GetEncoder()->Track(sampler_vk.GetSharedSampler())) {
+ return false;
}
+ vk::DescriptorImageInfo image_info;
+ image_info.imageLayout = vk::ImageLayout::eShaderReadOnlyOptimal;
+ image_info.sampler = sampler_vk.GetSampler();
+ image_info.imageView = texture_vk.GetImageView();
+ image_workspace_[bound_image_offset_++] = image_info;
+
+ vk::WriteDescriptorSet write_set;
+ write_set.dstBinding = slot.binding;
+ write_set.descriptorCount = 1u;
+ write_set.descriptorType = ToVKDescriptorType(type);
+ write_set.pImageInfo = &image_workspace_[bound_image_offset_ - 1];
+
+ write_workspace_[descriptor_write_offset_++] = write_set;
+ return true;
+}
+
+bool ComputePassVK::BindResource(size_t binding,
+ DescriptorType type,
+ const BufferView& view) {
+ if (bound_buffer_offset_ >= kMaxBindings) {
+ return false;
+ }
+
+ const std::shared_ptr<const DeviceBuffer>& device_buffer = view.buffer;
+ auto buffer = DeviceBufferVK::Cast(*device_buffer).GetBuffer();
+ if (!buffer) {
+ return false;
+ }
+
+ if (!command_buffer_->GetEncoder()->Track(device_buffer)) {
+ return false;
+ }
+
+ uint32_t offset = view.range.offset;
+
+ vk::DescriptorBufferInfo buffer_info;
+ buffer_info.buffer = buffer;
+ buffer_info.offset = offset;
+ buffer_info.range = view.range.length;
+ buffer_workspace_[bound_buffer_offset_++] = buffer_info;
+
+ vk::WriteDescriptorSet write_set;
+ write_set.dstBinding = binding;
+ write_set.descriptorCount = 1u;
+ write_set.descriptorType = ToVKDescriptorType(type);
+ write_set.pBufferInfo = &buffer_workspace_[bound_buffer_offset_ - 1];
+
+ write_workspace_[descriptor_write_offset_++] = write_set;
+ return true;
+}
+
+// |ComputePass|
+bool ComputePassVK::EncodeCommands() const {
return true;
}
diff --git a/impeller/renderer/backend/vulkan/compute_pass_vk.h b/impeller/renderer/backend/vulkan/compute_pass_vk.h
index 2e2259c..7b6fad8 100644
--- a/impeller/renderer/backend/vulkan/compute_pass_vk.h
+++ b/impeller/renderer/backend/vulkan/compute_pass_vk.h
@@ -5,8 +5,6 @@
#ifndef FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_COMPUTE_PASS_VK_H_
#define FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_COMPUTE_PASS_VK_H_
-#include "flutter/fml/macros.h"
-#include "impeller/renderer/backend/vulkan/binding_helpers_vk.h"
#include "impeller/renderer/backend/vulkan/command_encoder_vk.h"
#include "impeller/renderer/compute_pass.h"
@@ -22,16 +20,26 @@
private:
friend class CommandBufferVK;
- std::weak_ptr<CommandBufferVK> command_buffer_;
+ std::shared_ptr<CommandBufferVK> command_buffer_;
std::string label_;
+ std::array<uint32_t, 3> max_wg_size_ = {};
bool is_valid_ = false;
- mutable std::array<vk::DescriptorImageInfo, kMaxBindings> image_workspace_;
- mutable std::array<vk::DescriptorBufferInfo, kMaxBindings> buffer_workspace_;
- mutable std::array<vk::WriteDescriptorSet, kMaxBindings + kMaxBindings>
- write_workspace_;
- ComputePassVK(std::weak_ptr<const Context> context,
- std::weak_ptr<CommandBufferVK> command_buffer);
+ // Per-command state.
+ std::array<vk::DescriptorImageInfo, kMaxBindings> image_workspace_;
+ std::array<vk::DescriptorBufferInfo, kMaxBindings> buffer_workspace_;
+ std::array<vk::WriteDescriptorSet, kMaxBindings + kMaxBindings>
+ write_workspace_;
+ size_t bound_image_offset_ = 0u;
+ size_t bound_buffer_offset_ = 0u;
+ size_t descriptor_write_offset_ = 0u;
+ bool has_label_ = false;
+ bool pipeline_valid_ = false;
+ vk::DescriptorSet descriptor_set_ = {};
+ vk::PipelineLayout pipeline_layout_ = {};
+
+ ComputePassVK(std::shared_ptr<const Context> context,
+ std::shared_ptr<CommandBufferVK> command_buffer);
// |ComputePass|
bool IsValid() const override;
@@ -40,9 +48,36 @@
void OnSetLabel(const std::string& label) override;
// |ComputePass|
- bool OnEncodeCommands(const Context& context,
- const ISize& grid_size,
- const ISize& thread_group_size) const override;
+ bool EncodeCommands() const override;
+
+ // |ComputePass|
+ void SetCommandLabel(std::string_view label) override;
+
+ // |ComputePass|
+ void SetPipeline(const std::shared_ptr<Pipeline<ComputePipelineDescriptor>>&
+ pipeline) override;
+
+ // |ComputePass|
+ fml::Status Compute(const ISize& grid_size) override;
+
+ // |ResourceBinder|
+ bool BindResource(ShaderStage stage,
+ DescriptorType type,
+ const ShaderUniformSlot& slot,
+ const ShaderMetadata& metadata,
+ BufferView view) override;
+
+ // |ResourceBinder|
+ bool BindResource(ShaderStage stage,
+ DescriptorType type,
+ const SampledImageSlot& slot,
+ const ShaderMetadata& metadata,
+ std::shared_ptr<const Texture> texture,
+ std::shared_ptr<const Sampler> sampler) override;
+
+ bool BindResource(size_t binding,
+ DescriptorType type,
+ const BufferView& view);
};
} // namespace impeller
diff --git a/impeller/renderer/backend/vulkan/pass_bindings_cache.cc b/impeller/renderer/backend/vulkan/pass_bindings_cache.cc
deleted file mode 100644
index 00d8774..0000000
--- a/impeller/renderer/backend/vulkan/pass_bindings_cache.cc
+++ /dev/null
@@ -1,72 +0,0 @@
-// Copyright 2013 The Flutter Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-
-#include "impeller/renderer/backend/vulkan/pass_bindings_cache.h"
-
-namespace impeller {
-void PassBindingsCache::BindPipeline(vk::CommandBuffer command_buffer,
- vk::PipelineBindPoint pipeline_bind_point,
- vk::Pipeline pipeline) {
- switch (pipeline_bind_point) {
- case vk::PipelineBindPoint::eGraphics:
- if (graphics_pipeline_.has_value() &&
- graphics_pipeline_.value() == pipeline) {
- return;
- }
- graphics_pipeline_ = pipeline;
- break;
- case vk::PipelineBindPoint::eCompute:
- if (compute_pipeline_.has_value() &&
- compute_pipeline_.value() == pipeline) {
- return;
- }
- compute_pipeline_ = pipeline;
- break;
- default:
- break;
- }
- command_buffer.bindPipeline(pipeline_bind_point, pipeline);
-}
-
-void PassBindingsCache::SetStencilReference(vk::CommandBuffer command_buffer,
- vk::StencilFaceFlags face_mask,
- uint32_t reference) {
- if (stencil_face_flags_.has_value() &&
- face_mask == stencil_face_flags_.value() &&
- reference == stencil_reference_) {
- return;
- }
- stencil_face_flags_ = face_mask;
- stencil_reference_ = reference;
- command_buffer.setStencilReference(face_mask, reference);
-}
-
-void PassBindingsCache::SetScissor(vk::CommandBuffer command_buffer,
- uint32_t first_scissor,
- uint32_t scissor_count,
- const vk::Rect2D* scissors) {
- if (first_scissor == 0 && scissor_count == 1) {
- if (scissors_.has_value() && scissors_.value() == scissors[0]) {
- return;
- }
- scissors_ = scissors[0];
- }
- command_buffer.setScissor(first_scissor, scissor_count, scissors);
-}
-
-void PassBindingsCache::SetViewport(vk::CommandBuffer command_buffer,
- uint32_t first_viewport,
- uint32_t viewport_count,
- const vk::Viewport* viewports) {
- if (first_viewport == 0 && viewport_count == 1) {
- // Note that this is doing equality checks on floating point numbers.
- if (viewport_.has_value() && viewport_.value() == viewports[0]) {
- return;
- }
- viewport_ = viewports[0];
- }
- command_buffer.setViewport(first_viewport, viewport_count, viewports);
-}
-
-} // namespace impeller
diff --git a/impeller/renderer/backend/vulkan/pass_bindings_cache.h b/impeller/renderer/backend/vulkan/pass_bindings_cache.h
deleted file mode 100644
index 9eccee0..0000000
--- a/impeller/renderer/backend/vulkan/pass_bindings_cache.h
+++ /dev/null
@@ -1,50 +0,0 @@
-// Copyright 2013 The Flutter Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-
-#ifndef FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_PASS_BINDINGS_CACHE_H_
-#define FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_PASS_BINDINGS_CACHE_H_
-
-#include <cstdint>
-#include <optional>
-
-#include "flutter/impeller/renderer/backend/vulkan/vk.h"
-
-namespace impeller {
-
-class PassBindingsCache {
- public:
- void BindPipeline(vk::CommandBuffer command_buffer,
- vk::PipelineBindPoint pipeline_bind_point,
- vk::Pipeline pipeline);
-
- void SetStencilReference(vk::CommandBuffer command_buffer,
- vk::StencilFaceFlags face_mask,
- uint32_t reference);
-
- void SetScissor(vk::CommandBuffer command_buffer,
- uint32_t first_scissor,
- uint32_t scissor_count,
- const vk::Rect2D* scissors);
-
- void SetViewport(vk::CommandBuffer command_buffer,
- uint32_t first_viewport,
- uint32_t viewport_count,
- const vk::Viewport* viewports);
-
- private:
- // bindPipeline
- std::optional<vk::Pipeline> graphics_pipeline_;
- std::optional<vk::Pipeline> compute_pipeline_;
- // setStencilReference
- std::optional<vk::StencilFaceFlags> stencil_face_flags_;
- uint32_t stencil_reference_ = 0;
- // setScissor
- std::optional<vk::Rect2D> scissors_;
- // setViewport
- std::optional<vk::Viewport> viewport_;
-};
-
-} // namespace impeller
-
-#endif // FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_PASS_BINDINGS_CACHE_H_
diff --git a/impeller/renderer/backend/vulkan/pass_bindings_cache_unittests.cc b/impeller/renderer/backend/vulkan/pass_bindings_cache_unittests.cc
deleted file mode 100644
index e76213c..0000000
--- a/impeller/renderer/backend/vulkan/pass_bindings_cache_unittests.cc
+++ /dev/null
@@ -1,82 +0,0 @@
-// Copyright 2013 The Flutter Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-
-#include "flutter/testing/testing.h"
-#include "impeller/renderer/backend/vulkan/command_encoder_vk.h"
-#include "impeller/renderer/backend/vulkan/pass_bindings_cache.h"
-#include "impeller/renderer/backend/vulkan/test/mock_vulkan.h"
-
-namespace impeller {
-namespace testing {
-
-namespace {
-int32_t CountStringViewInstances(const std::vector<std::string>& strings,
- std::string_view target) {
- int32_t count = 0;
- for (const std::string& str : strings) {
- if (str == target) {
- count++;
- }
- }
- return count;
-}
-} // namespace
-
-TEST(PassBindingsCacheTest, bindPipeline) {
- auto context = MockVulkanContextBuilder().Build();
- PassBindingsCache cache;
- auto encoder = std::make_unique<CommandEncoderFactoryVK>(context)->Create();
- auto buffer = encoder->GetCommandBuffer();
- VkPipeline vk_pipeline = reinterpret_cast<VkPipeline>(0xfeedface);
- vk::Pipeline pipeline(vk_pipeline);
- cache.BindPipeline(buffer, vk::PipelineBindPoint::eGraphics, pipeline);
- cache.BindPipeline(buffer, vk::PipelineBindPoint::eGraphics, pipeline);
- std::shared_ptr<std::vector<std::string>> functions =
- GetMockVulkanFunctions(context->GetDevice());
- EXPECT_EQ(CountStringViewInstances(*functions, "vkCmdBindPipeline"), 1);
-}
-
-TEST(PassBindingsCacheTest, setStencilReference) {
- auto context = MockVulkanContextBuilder().Build();
- PassBindingsCache cache;
- auto encoder = std::make_unique<CommandEncoderFactoryVK>(context)->Create();
- auto buffer = encoder->GetCommandBuffer();
- cache.SetStencilReference(
- buffer, vk::StencilFaceFlagBits::eVkStencilFrontAndBack, 123);
- cache.SetStencilReference(
- buffer, vk::StencilFaceFlagBits::eVkStencilFrontAndBack, 123);
- std::shared_ptr<std::vector<std::string>> functions =
- GetMockVulkanFunctions(context->GetDevice());
- EXPECT_EQ(CountStringViewInstances(*functions, "vkCmdSetStencilReference"),
- 1);
-}
-
-TEST(PassBindingsCacheTest, setScissor) {
- auto context = MockVulkanContextBuilder().Build();
- PassBindingsCache cache;
- auto encoder = std::make_unique<CommandEncoderFactoryVK>(context)->Create();
- auto buffer = encoder->GetCommandBuffer();
- vk::Rect2D scissors;
- cache.SetScissor(buffer, 0, 1, &scissors);
- cache.SetScissor(buffer, 0, 1, &scissors);
- std::shared_ptr<std::vector<std::string>> functions =
- GetMockVulkanFunctions(context->GetDevice());
- EXPECT_EQ(CountStringViewInstances(*functions, "vkCmdSetScissor"), 1);
-}
-
-TEST(PassBindingsCacheTest, setViewport) {
- auto context = MockVulkanContextBuilder().Build();
- PassBindingsCache cache;
- auto encoder = std::make_unique<CommandEncoderFactoryVK>(context)->Create();
- auto buffer = encoder->GetCommandBuffer();
- vk::Viewport viewports;
- cache.SetViewport(buffer, 0, 1, &viewports);
- cache.SetViewport(buffer, 0, 1, &viewports);
- std::shared_ptr<std::vector<std::string>> functions =
- GetMockVulkanFunctions(context->GetDevice());
- EXPECT_EQ(CountStringViewInstances(*functions, "vkCmdSetViewport"), 1);
-}
-
-} // namespace testing
-} // namespace impeller
diff --git a/impeller/renderer/backend/vulkan/pipeline_vk.h b/impeller/renderer/backend/vulkan/pipeline_vk.h
index ffe2451..c27d186 100644
--- a/impeller/renderer/backend/vulkan/pipeline_vk.h
+++ b/impeller/renderer/backend/vulkan/pipeline_vk.h
@@ -14,6 +14,10 @@
namespace impeller {
+// Limit on the total number of buffer and image bindings that allow the Vulkan
+// backend to avoid dynamic heap allocations.
+static constexpr size_t kMaxBindings = 32;
+
class PipelineVK final
: public Pipeline<PipelineDescriptor>,
public BackendCast<PipelineVK, Pipeline<PipelineDescriptor>> {
diff --git a/impeller/renderer/backend/vulkan/render_pass_vk.cc b/impeller/renderer/backend/vulkan/render_pass_vk.cc
index 0489bfc..5bbe7b4 100644
--- a/impeller/renderer/backend/vulkan/render_pass_vk.cc
+++ b/impeller/renderer/backend/vulkan/render_pass_vk.cc
@@ -14,7 +14,6 @@
#include "impeller/core/formats.h"
#include "impeller/core/texture.h"
#include "impeller/renderer/backend/vulkan/barrier_vk.h"
-#include "impeller/renderer/backend/vulkan/binding_helpers_vk.h"
#include "impeller/renderer/backend/vulkan/command_buffer_vk.h"
#include "impeller/renderer/backend/vulkan/command_encoder_vk.h"
#include "impeller/renderer/backend/vulkan/context_vk.h"
diff --git a/impeller/renderer/backend/vulkan/render_pass_vk.h b/impeller/renderer/backend/vulkan/render_pass_vk.h
index c7da1ae..ad9e89d 100644
--- a/impeller/renderer/backend/vulkan/render_pass_vk.h
+++ b/impeller/renderer/backend/vulkan/render_pass_vk.h
@@ -6,7 +6,6 @@
#define FLUTTER_IMPELLER_RENDERER_BACKEND_VULKAN_RENDER_PASS_VK_H_
#include "impeller/core/buffer_view.h"
-#include "impeller/renderer/backend/vulkan/binding_helpers_vk.h"
#include "impeller/renderer/backend/vulkan/context_vk.h"
#include "impeller/renderer/backend/vulkan/pipeline_vk.h"
#include "impeller/renderer/backend/vulkan/shared_object_vk.h"
diff --git a/impeller/renderer/compute_command.cc b/impeller/renderer/compute_command.cc
deleted file mode 100644
index e1a0601..0000000
--- a/impeller/renderer/compute_command.cc
+++ /dev/null
@@ -1,59 +0,0 @@
-// Copyright 2013 The Flutter Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-
-#include "impeller/renderer/compute_command.h"
-
-#include <utility>
-
-#include "impeller/base/validation.h"
-#include "impeller/core/formats.h"
-#include "impeller/core/shader_types.h"
-
-namespace impeller {
-
-bool ComputeCommand::BindResource(ShaderStage stage,
- DescriptorType type,
- const ShaderUniformSlot& slot,
- const ShaderMetadata& metadata,
- BufferView view) {
- if (stage != ShaderStage::kCompute) {
- VALIDATION_LOG << "Use Command for non-compute shader stages.";
- return false;
- }
- if (!view) {
- return false;
- }
-
- bindings.buffers.emplace_back(
- BufferAndUniformSlot{.slot = slot, .view = {&metadata, std::move(view)}});
- return true;
-}
-
-bool ComputeCommand::BindResource(ShaderStage stage,
- DescriptorType type,
- const SampledImageSlot& slot,
- const ShaderMetadata& metadata,
- std::shared_ptr<const Texture> texture,
- std::shared_ptr<const Sampler> sampler) {
- if (stage != ShaderStage::kCompute) {
- VALIDATION_LOG << "Use Command for non-compute shader stages.";
- return false;
- }
- if (!sampler || !sampler->IsValid()) {
- return false;
- }
- if (!texture || !texture->IsValid()) {
- return false;
- }
-
- bindings.sampled_images.emplace_back(TextureAndSampler{
- .slot = slot,
- .texture = {&metadata, std::move(texture)},
- .sampler = std::move(sampler),
- });
-
- return false;
-}
-
-} // namespace impeller
diff --git a/impeller/renderer/compute_command.h b/impeller/renderer/compute_command.h
deleted file mode 100644
index c0c56e9..0000000
--- a/impeller/renderer/compute_command.h
+++ /dev/null
@@ -1,76 +0,0 @@
-// Copyright 2013 The Flutter Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-
-#ifndef FLUTTER_IMPELLER_RENDERER_COMPUTE_COMMAND_H_
-#define FLUTTER_IMPELLER_RENDERER_COMPUTE_COMMAND_H_
-
-#include <memory>
-#include <string>
-
-#include "impeller/core/buffer_view.h"
-#include "impeller/core/formats.h"
-#include "impeller/core/resource_binder.h"
-#include "impeller/core/sampler.h"
-#include "impeller/core/shader_types.h"
-#include "impeller/core/texture.h"
-#include "impeller/renderer/command.h"
-#include "impeller/renderer/compute_pipeline_descriptor.h"
-#include "impeller/renderer/pipeline.h"
-
-namespace impeller {
-
-//------------------------------------------------------------------------------
-/// @brief An object used to specify compute work to the GPU along with
-/// references to resources the GPU will used when doing said work.
-///
-/// To construct a valid command, follow these steps:
-/// * Specify a valid pipeline.
-/// * (Optional) Specify a debug label.
-///
-/// Command are very lightweight objects and can be created
-/// frequently and on demand. The resources referenced in commands
-/// views into buffers managed by other allocators and resource
-/// managers.
-///
-struct ComputeCommand : public ResourceBinder {
- //----------------------------------------------------------------------------
- /// The pipeline to use for this command.
- ///
- std::shared_ptr<Pipeline<ComputePipelineDescriptor>> pipeline;
- //----------------------------------------------------------------------------
- /// The buffer, texture, and sampler bindings used by the compute pipeline
- /// stage.
- ///
- Bindings bindings;
-
-#ifdef IMPELLER_DEBUG
- //----------------------------------------------------------------------------
- /// The debugging label to use for the command.
- ///
- std::string label;
-#endif // IMPELLER_DEBUG
-
- // |ResourceBinder|
- bool BindResource(ShaderStage stage,
- DescriptorType type,
- const ShaderUniformSlot& slot,
- const ShaderMetadata& metadata,
- BufferView view) override;
-
- // |ResourceBinder|
- bool BindResource(ShaderStage stage,
- DescriptorType type,
- const SampledImageSlot& slot,
- const ShaderMetadata& metadata,
- std::shared_ptr<const Texture> texture,
- std::shared_ptr<const Sampler> sampler) override;
-
- constexpr explicit operator bool() const {
- return pipeline && pipeline->IsValid();
- }
-};
-
-} // namespace impeller
-
-#endif // FLUTTER_IMPELLER_RENDERER_COMPUTE_COMMAND_H_
diff --git a/impeller/renderer/compute_pass.cc b/impeller/renderer/compute_pass.cc
index 8446fff..e069aca 100644
--- a/impeller/renderer/compute_pass.cc
+++ b/impeller/renderer/compute_pass.cc
@@ -3,14 +3,10 @@
// found in the LICENSE file.
#include "impeller/renderer/compute_pass.h"
-#include <memory>
-
-#include "fml/logging.h"
-#include "impeller/base/validation.h"
namespace impeller {
-ComputePass::ComputePass(std::weak_ptr<const Context> context)
+ComputePass::ComputePass(std::shared_ptr<const Context> context)
: context_(std::move(context)) {}
ComputePass::~ComputePass() = default;
@@ -22,37 +18,4 @@
OnSetLabel(label);
}
-void ComputePass::SetGridSize(const ISize& size) {
- grid_size_ = size;
-}
-
-void ComputePass::SetThreadGroupSize(const ISize& size) {
- thread_group_size_ = size;
-}
-
-bool ComputePass::AddCommand(ComputeCommand command) {
- if (!command) {
- VALIDATION_LOG
- << "Attempted to add an invalid command to the compute pass.";
- return false;
- }
-
- commands_.emplace_back(std::move(command));
- return true;
-}
-
-bool ComputePass::EncodeCommands() const {
- if (grid_size_.IsEmpty() || thread_group_size_.IsEmpty()) {
- FML_DLOG(WARNING) << "Attempted to encode a compute pass with an empty "
- "grid or thread group size.";
- return false;
- }
- auto context = context_.lock();
- // The context could have been collected in the meantime.
- if (!context) {
- return false;
- }
- return OnEncodeCommands(*context, grid_size_, thread_group_size_);
-}
-
} // namespace impeller
diff --git a/impeller/renderer/compute_pass.h b/impeller/renderer/compute_pass.h
index 4ba6e80..40c313e 100644
--- a/impeller/renderer/compute_pass.h
+++ b/impeller/renderer/compute_pass.h
@@ -7,20 +7,20 @@
#include <string>
-#include "impeller/renderer/compute_command.h"
+#include "fml/status.h"
+#include "impeller/core/resource_binder.h"
+#include "impeller/renderer/compute_pipeline_descriptor.h"
+#include "impeller/renderer/pipeline_descriptor.h"
namespace impeller {
-class HostBuffer;
-class Allocator;
-
//------------------------------------------------------------------------------
/// @brief Compute passes encode compute shader into the underlying command
/// buffer.
///
/// @see `CommandBuffer`
///
-class ComputePass {
+class ComputePass : public ResourceBinder {
public:
virtual ~ComputePass();
@@ -28,48 +28,29 @@
void SetLabel(const std::string& label);
- void SetGridSize(const ISize& size);
+ virtual void SetCommandLabel(std::string_view label) = 0;
- void SetThreadGroupSize(const ISize& size);
+ virtual void SetPipeline(
+ const std::shared_ptr<Pipeline<ComputePipelineDescriptor>>& pipeline) = 0;
- //----------------------------------------------------------------------------
- /// @brief Record a command for subsequent encoding to the underlying
- /// command buffer. No work is encoded into the command buffer at
- /// this time.
- ///
- /// @param[in] command The command
- ///
- /// @return If the command was valid for subsequent commitment.
- ///
- bool AddCommand(ComputeCommand command);
+ virtual fml::Status Compute(const ISize& grid_size) = 0;
//----------------------------------------------------------------------------
/// @brief Encode the recorded commands to the underlying command buffer.
///
- /// @param transients_allocator The transients allocator.
- ///
/// @return If the commands were encoded to the underlying command
/// buffer.
///
- bool EncodeCommands() const;
+ virtual bool EncodeCommands() const = 0;
protected:
- const std::weak_ptr<const Context> context_;
- std::vector<ComputeCommand> commands_;
+ const std::shared_ptr<const Context> context_;
- explicit ComputePass(std::weak_ptr<const Context> context);
+ explicit ComputePass(std::shared_ptr<const Context> context);
virtual void OnSetLabel(const std::string& label) = 0;
- virtual bool OnEncodeCommands(const Context& context,
- const ISize& grid_size,
- const ISize& thread_group_size) const = 0;
-
private:
- std::shared_ptr<HostBuffer> transients_buffer_;
- ISize grid_size_ = ISize(32, 32);
- ISize thread_group_size_ = ISize(32, 32);
-
ComputePass(const ComputePass&) = delete;
ComputePass& operator=(const ComputePass&) = delete;
diff --git a/impeller/renderer/compute_subgroup_unittests.cc b/impeller/renderer/compute_subgroup_unittests.cc
index ee26381..144b635 100644
--- a/impeller/renderer/compute_subgroup_unittests.cc
+++ b/impeller/renderer/compute_subgroup_unittests.cc
@@ -24,7 +24,6 @@
#include "impeller/geometry/path_component.h"
#include "impeller/playground/compute_playground_test.h"
#include "impeller/renderer/command_buffer.h"
-#include "impeller/renderer/compute_command.h"
#include "impeller/renderer/compute_pipeline_builder.h"
#include "impeller/renderer/compute_tessellator.h"
#include "impeller/renderer/path_polyline.comp.h"
diff --git a/impeller/renderer/compute_tessellator.cc b/impeller/renderer/compute_tessellator.cc
index ebcf76d..5cd9877 100644
--- a/impeller/renderer/compute_tessellator.cc
+++ b/impeller/renderer/compute_tessellator.cc
@@ -121,21 +121,17 @@
context->GetPipelineLibrary()->GetPipeline(pipeline_desc).Get();
FML_DCHECK(compute_pipeline);
- pass->SetGridSize(ISize(line_count, 1));
- pass->SetThreadGroupSize(ISize(line_count, 1));
+ pass->SetPipeline(compute_pipeline);
+ pass->SetCommandLabel("Generate Polyline");
- ComputeCommand cmd;
- DEBUG_COMMAND_INFO(cmd, "Generate Polyline");
- cmd.pipeline = compute_pipeline;
+ PS::BindConfig(*pass, host_buffer.EmplaceUniform(config));
+ PS::BindCubics(*pass, host_buffer.EmplaceStorageBuffer(cubics));
+ PS::BindQuads(*pass, host_buffer.EmplaceStorageBuffer(quads));
+ PS::BindLines(*pass, host_buffer.EmplaceStorageBuffer(lines));
+ PS::BindComponents(*pass, host_buffer.EmplaceStorageBuffer(components));
+ PS::BindPolyline(*pass, DeviceBuffer::AsBufferView(polyline_buffer));
- PS::BindConfig(cmd, host_buffer.EmplaceUniform(config));
- PS::BindCubics(cmd, host_buffer.EmplaceStorageBuffer(cubics));
- PS::BindQuads(cmd, host_buffer.EmplaceStorageBuffer(quads));
- PS::BindLines(cmd, host_buffer.EmplaceStorageBuffer(lines));
- PS::BindComponents(cmd, host_buffer.EmplaceStorageBuffer(components));
- PS::BindPolyline(cmd, DeviceBuffer::AsBufferView(polyline_buffer));
-
- if (!pass->AddCommand(std::move(cmd))) {
+ if (!pass->Compute(ISize(line_count, 1)).ok()) {
return Status::kCommandInvalid;
}
}
@@ -149,12 +145,8 @@
context->GetPipelineLibrary()->GetPipeline(pipeline_desc).Get();
FML_DCHECK(compute_pipeline);
- pass->SetGridSize(ISize(line_count, 1));
- pass->SetThreadGroupSize(ISize(line_count, 1));
-
- ComputeCommand cmd;
- DEBUG_COMMAND_INFO(cmd, "Compute Stroke");
- cmd.pipeline = compute_pipeline;
+ pass->SetPipeline(compute_pipeline);
+ pass->SetCommandLabel("Compute Stroke");
SS::Config config{
.width = stroke_width_,
@@ -162,13 +154,13 @@
.join = static_cast<uint32_t>(stroke_join_),
.miter_limit = miter_limit_,
};
- SS::BindConfig(cmd, host_buffer.EmplaceUniform(config));
+ SS::BindConfig(*pass, host_buffer.EmplaceUniform(config));
- SS::BindPolyline(cmd, DeviceBuffer::AsBufferView(polyline_buffer));
- SS::BindVertexBufferCount(cmd, std::move(vertex_buffer_count));
- SS::BindVertexBuffer(cmd, std::move(vertex_buffer));
+ SS::BindPolyline(*pass, DeviceBuffer::AsBufferView(polyline_buffer));
+ SS::BindVertexBufferCount(*pass, std::move(vertex_buffer_count));
+ SS::BindVertexBuffer(*pass, std::move(vertex_buffer));
- if (!pass->AddCommand(std::move(cmd))) {
+ if (!pass->Compute(ISize(line_count, 1)).ok()) {
return Status::kCommandInvalid;
}
}
diff --git a/impeller/renderer/compute_unittests.cc b/impeller/renderer/compute_unittests.cc
index ce416c6..d8fd55d 100644
--- a/impeller/renderer/compute_unittests.cc
+++ b/impeller/renderer/compute_unittests.cc
@@ -3,20 +3,14 @@
// found in the LICENSE file.
#include "flutter/fml/synchronization/waitable_event.h"
-#include "flutter/fml/time/time_point.h"
#include "flutter/testing/testing.h"
#include "gmock/gmock.h"
-#include "impeller/base/strings.h"
-#include "impeller/core/formats.h"
#include "impeller/core/host_buffer.h"
#include "impeller/fixtures/sample.comp.h"
#include "impeller/fixtures/stage1.comp.h"
#include "impeller/fixtures/stage2.comp.h"
-#include "impeller/geometry/path.h"
-#include "impeller/geometry/path_component.h"
#include "impeller/playground/compute_playground_test.h"
#include "impeller/renderer/command_buffer.h"
-#include "impeller/renderer/compute_command.h"
#include "impeller/renderer/compute_pipeline_builder.h"
#include "impeller/renderer/pipeline_library.h"
#include "impeller/renderer/prefix_sum_test.comp.h"
@@ -54,11 +48,7 @@
static constexpr size_t kCount = 5;
- pass->SetGridSize(ISize(kCount, 1));
- pass->SetThreadGroupSize(ISize(kCount, 1));
-
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline;
+ pass->SetPipeline(compute_pipeline);
CS::Info info{.count = kCount};
CS::Input0<kCount> input_0;
@@ -76,12 +66,12 @@
auto output_buffer = CreateHostVisibleDeviceBuffer<CS::Output<kCount>>(
context, "Output Buffer");
- CS::BindInfo(cmd, host_buffer->EmplaceUniform(info));
- CS::BindInput0(cmd, host_buffer->EmplaceStorageBuffer(input_0));
- CS::BindInput1(cmd, host_buffer->EmplaceStorageBuffer(input_1));
- CS::BindOutput(cmd, DeviceBuffer::AsBufferView(output_buffer));
+ CS::BindInfo(*pass, host_buffer->EmplaceUniform(info));
+ CS::BindInput0(*pass, host_buffer->EmplaceStorageBuffer(input_0));
+ CS::BindInput1(*pass, host_buffer->EmplaceStorageBuffer(input_1));
+ CS::BindOutput(*pass, DeviceBuffer::AsBufferView(output_buffer));
- ASSERT_TRUE(pass->AddCommand(std::move(cmd)));
+ ASSERT_TRUE(pass->Compute(ISize(kCount, 1)).ok());
ASSERT_TRUE(pass->EncodeCommands());
fml::AutoResetWaitableEvent latch;
@@ -131,11 +121,7 @@
static constexpr size_t kCount = 5;
- pass->SetGridSize(ISize(kCount, 1));
- pass->SetThreadGroupSize(ISize(kCount, 1));
-
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline;
+ pass->SetPipeline(compute_pipeline);
CS::InputData<kCount> input_data;
input_data.count = kCount;
@@ -146,10 +132,10 @@
auto output_buffer = CreateHostVisibleDeviceBuffer<CS::OutputData<kCount>>(
context, "Output Buffer");
- CS::BindInputData(cmd, host_buffer->EmplaceStorageBuffer(input_data));
- CS::BindOutputData(cmd, DeviceBuffer::AsBufferView(output_buffer));
+ CS::BindInputData(*pass, host_buffer->EmplaceStorageBuffer(input_data));
+ CS::BindOutputData(*pass, DeviceBuffer::AsBufferView(output_buffer));
- ASSERT_TRUE(pass->AddCommand(std::move(cmd)));
+ ASSERT_TRUE(pass->Compute(ISize(kCount, 1)).ok());
ASSERT_TRUE(pass->EncodeCommands());
fml::AutoResetWaitableEvent latch;
@@ -195,18 +181,14 @@
static constexpr size_t kCount = 2048;
- pass->SetGridSize(ISize(kCount, 1));
- pass->SetThreadGroupSize(ISize(kCount, 1));
-
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline;
+ pass->SetPipeline(compute_pipeline);
auto output_buffer = CreateHostVisibleDeviceBuffer<CS::OutputData<kCount>>(
context, "Output Buffer");
- CS::BindOutputData(cmd, DeviceBuffer::AsBufferView(output_buffer));
+ CS::BindOutputData(*pass, DeviceBuffer::AsBufferView(output_buffer));
- ASSERT_TRUE(pass->AddCommand(std::move(cmd)));
+ ASSERT_TRUE(pass->Compute(ISize(kCount, 1)).ok());
ASSERT_TRUE(pass->EncodeCommands());
fml::AutoResetWaitableEvent latch;
@@ -248,10 +230,7 @@
static constexpr size_t kCount = 1023;
- pass->SetGridSize(ISize(kCount, 1));
-
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline;
+ pass->SetPipeline(compute_pipeline);
CS::InputData<kCount> input_data;
input_data.count = kCount;
@@ -262,10 +241,10 @@
auto output_buffer = CreateHostVisibleDeviceBuffer<CS::OutputData<kCount>>(
context, "Output Buffer");
- CS::BindInputData(cmd, host_buffer->EmplaceStorageBuffer(input_data));
- CS::BindOutputData(cmd, DeviceBuffer::AsBufferView(output_buffer));
+ CS::BindInputData(*pass, host_buffer->EmplaceStorageBuffer(input_data));
+ CS::BindOutputData(*pass, DeviceBuffer::AsBufferView(output_buffer));
- pass->AddCommand(std::move(cmd));
+ pass->Compute(ISize(kCount, 1));
pass->EncodeCommands();
host_buffer->Reset();
return cmd_buffer->SubmitCommands();
@@ -305,9 +284,6 @@
static constexpr size_t kCount1 = 5;
static constexpr size_t kCount2 = kCount1 * 2;
- pass->SetGridSize(ISize(512, 1));
- pass->SetThreadGroupSize(ISize(512, 1));
-
CS1::Input<kCount1> input_1;
input_1.count = kCount1;
for (size_t i = 0; i < kCount1; i++) {
@@ -326,22 +302,20 @@
context, "Output Buffer Stage 2");
{
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline_1;
+ pass->SetPipeline(compute_pipeline_1);
- CS1::BindInput(cmd, host_buffer->EmplaceStorageBuffer(input_1));
- CS1::BindOutput(cmd, DeviceBuffer::AsBufferView(output_buffer_1));
+ CS1::BindInput(*pass, host_buffer->EmplaceStorageBuffer(input_1));
+ CS1::BindOutput(*pass, DeviceBuffer::AsBufferView(output_buffer_1));
- ASSERT_TRUE(pass->AddCommand(std::move(cmd)));
+ ASSERT_TRUE(pass->Compute(ISize(512, 1)).ok());
}
{
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline_2;
+ pass->SetPipeline(compute_pipeline_2);
- CS1::BindInput(cmd, DeviceBuffer::AsBufferView(output_buffer_1));
- CS2::BindOutput(cmd, DeviceBuffer::AsBufferView(output_buffer_2));
- ASSERT_TRUE(pass->AddCommand(std::move(cmd)));
+ CS1::BindInput(*pass, DeviceBuffer::AsBufferView(output_buffer_1));
+ CS2::BindOutput(*pass, DeviceBuffer::AsBufferView(output_buffer_2));
+ ASSERT_TRUE(pass->Compute(ISize(512, 1)).ok());
}
ASSERT_TRUE(pass->EncodeCommands());
@@ -393,10 +367,7 @@
static constexpr size_t kCount = 5;
- pass->SetGridSize(ISize(kCount, 1));
-
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline;
+ pass->SetPipeline(compute_pipeline);
CS::Info info{.count = kCount};
CS::Input0<kCount> input_0;
@@ -414,12 +385,12 @@
auto output_buffer = CreateHostVisibleDeviceBuffer<CS::Output<kCount>>(
context, "Output Buffer");
- CS::BindInfo(cmd, host_buffer->EmplaceUniform(info));
- CS::BindInput0(cmd, host_buffer->EmplaceStorageBuffer(input_0));
- CS::BindInput1(cmd, host_buffer->EmplaceStorageBuffer(input_1));
- CS::BindOutput(cmd, DeviceBuffer::AsBufferView(output_buffer));
+ CS::BindInfo(*pass, host_buffer->EmplaceUniform(info));
+ CS::BindInput0(*pass, host_buffer->EmplaceStorageBuffer(input_0));
+ CS::BindInput1(*pass, host_buffer->EmplaceStorageBuffer(input_1));
+ CS::BindOutput(*pass, DeviceBuffer::AsBufferView(output_buffer));
- ASSERT_TRUE(pass->AddCommand(std::move(cmd)));
+ ASSERT_TRUE(pass->Compute(ISize(kCount, 1)).ok());
ASSERT_TRUE(pass->EncodeCommands());
fml::AutoResetWaitableEvent latch;
@@ -469,13 +440,7 @@
static constexpr size_t kCount = 5;
- // Intentionally making the grid size zero in one dimension. No GPU will
- // tolerate this.
- pass->SetGridSize(ISize(0, 1));
- pass->SetThreadGroupSize(ISize(0, 1));
-
- ComputeCommand cmd;
- cmd.pipeline = compute_pipeline;
+ pass->SetPipeline(compute_pipeline);
CS::Info info{.count = kCount};
CS::Input0<kCount> input_0;
@@ -493,13 +458,14 @@
auto output_buffer = CreateHostVisibleDeviceBuffer<CS::Output<kCount>>(
context, "Output Buffer");
- CS::BindInfo(cmd, host_buffer->EmplaceUniform(info));
- CS::BindInput0(cmd, host_buffer->EmplaceStorageBuffer(input_0));
- CS::BindInput1(cmd, host_buffer->EmplaceStorageBuffer(input_1));
- CS::BindOutput(cmd, DeviceBuffer::AsBufferView(output_buffer));
+ CS::BindInfo(*pass, host_buffer->EmplaceUniform(info));
+ CS::BindInput0(*pass, host_buffer->EmplaceStorageBuffer(input_0));
+ CS::BindInput1(*pass, host_buffer->EmplaceStorageBuffer(input_1));
+ CS::BindOutput(*pass, DeviceBuffer::AsBufferView(output_buffer));
- ASSERT_TRUE(pass->AddCommand(std::move(cmd)));
- ASSERT_FALSE(pass->EncodeCommands());
+ // Intentionally making the grid size zero in one dimension. No GPU will
+ // tolerate this.
+ ASSERT_FALSE(pass->Compute(ISize(0, 1)).ok());
}
} // namespace testing
