[Impeller] dont emulate command buffers for compute on Metal/Vulkan. (#49922)

Part of the work for: flutter/flutter#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.

Author: Jonah Williams

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

ci/licenses_golden/licenses_flutter

@@ -5387,6 +5387,8 @@ ORIGIN: ../../../flutter/impeller/renderer/backend/metal/blit_pass_mtl.h + ../..
 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_buffe
 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.cc + ../../
 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.cc + ../../../flutter/LICENSE
 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.h
 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_
 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/formats_vk.h
 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.cc
 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

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 @@ GeometryResult PointFieldGeometry::GetPositionBufferGPU(
   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 @@ GeometryResult PointFieldGeometry::GetPositionBufferGPU(
     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,29 +200,25 @@ GeometryResult PointFieldGeometry::GetPositionBufferGPU(
 
     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;
     frame_info.effect_transform = effect_transform.value();
     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 {};
   }

impeller/renderer/BUILD.gn

@@ -55,8 +55,6 @@ impeller_component("renderer") {
     "command.h",
     "command_buffer.cc",
     "command_buffer.h",
-    "compute_command.cc",
-    "compute_command.h",
     "compute_pass.cc",
     "compute_pass.h",
     "compute_pipeline_builder.cc",

impeller/renderer/backend/metal/BUILD.gn

@@ -14,6 +14,8 @@ impeller_component("metal") {
     "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",

impeller/renderer/backend/metal/command_buffer_mtl.mm

@@ -220,9 +220,13 @@ static bool LogMTLCommandBufferErrorIfPresent(id<MTLCommandBuffer> buffer) {
   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;
   }

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_

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

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 @@ class ComputePassMTL final : public ComputePass {
   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;