| Flutter Impeller |

compute_pass_mtl.mm
Go to the documentation of this file.
1 // Copyright 2013 The Flutter Authors. All rights reserved.
2 // Use of this source code is governed by a BSD-style license that can be
3 // found in the LICENSE file.
4
5 #include "impeller/renderer/backend/metal/compute_pass_mtl.h"
6
7 #include <Metal/Metal.h>
8 #include <memory>
9 #include <variant>
10
11 #include "flutter/fml/backtrace.h"
12 #include "flutter/fml/closure.h"
13 #include "flutter/fml/logging.h"
14 #include "flutter/fml/trace_event.h"
15 #include "impeller/base/backend_cast.h"
16 #include "impeller/core/formats.h"
17 #include "impeller/core/host_buffer.h"
18 #include "impeller/core/shader_types.h"
19 #include "impeller/renderer/backend/metal/compute_pipeline_mtl.h"
20 #include "impeller/renderer/backend/metal/device_buffer_mtl.h"
21 #include "impeller/renderer/backend/metal/formats_mtl.h"
22 #include "impeller/renderer/backend/metal/sampler_mtl.h"
23 #include "impeller/renderer/backend/metal/texture_mtl.h"
24 #include "impeller/renderer/compute_command.h"
25
26 namespace impeller {
27
28 ComputePassMTL::ComputePassMTL(std::weak_ptr<const Context> context,
29 id<MTLCommandBuffer> buffer)
30 : ComputePass(std::move(context)), buffer_(buffer) {
31 if (!buffer_) {
32 return;
33 }
34 is_valid_ = true;
35 }
36
37 ComputePassMTL::~ComputePassMTL() = default;
38
39 bool ComputePassMTL::IsValid() const {
40 return is_valid_;
41 }
42
43 void ComputePassMTL::OnSetLabel(const std::string& label) {
44 if (label.empty()) {
45 return;
46 }
47 label_ = label;
48 }
49
50 bool ComputePassMTL::OnEncodeCommands(const Context& context,
51 const ISize& grid_size,
52 const ISize& thread_group_size) const {
53 TRACE_EVENT0("impeller", "ComputePassMTL::EncodeCommands");
54 if (!IsValid()) {
55 return false;
56 }
57
58 FML_DCHECK(!grid_size.IsEmpty() && !thread_group_size.IsEmpty());
59
60 // TODO(dnfield): Support non-serial dispatch type on higher iOS versions.
61 auto compute_command_encoder = [buffer_ computeCommandEncoder];
62
63 if (!compute_command_encoder) {
64 return false;
65 }
66
67 if (!label_.empty()) {
68 [compute_command_encoder setLabel:@(label_.c_str())];
69 }
70
71 // Success or failure, the pass must end. The buffer can only process one pass
72 // at a time.
73 fml::ScopedCleanupClosure auto_end(
74 [compute_command_encoder]() { [compute_command_encoder endEncoding]; });
75
76 return EncodeCommands(context.GetResourceAllocator(), compute_command_encoder,
77 grid_size, thread_group_size);
78 }
79
80 //-----------------------------------------------------------------------------
81 /// @brief Ensures that bindings on the pass are not redundantly set or
82 /// updated. Avoids making the driver do additional checks and makes
83 /// the frame insights during profiling and instrumentation not
84 /// complain about the same.
85 ///
86 /// There should be no change to rendering if this caching was
87 /// absent.
88 ///
89 struct ComputePassBindingsCache {
90 explicit ComputePassBindingsCache(id<MTLComputeCommandEncoder> encoder)
91 : encoder_(encoder) {}
92
93 ComputePassBindingsCache(const ComputePassBindingsCache&) = delete;
94
95 ComputePassBindingsCache(ComputePassBindingsCache&&) = delete;
96
97 void SetComputePipelineState(id<MTLComputePipelineState> pipeline) {
98 if (pipeline == pipeline_) {
99 return;
100 }
101 pipeline_ = pipeline;
102 [encoder_ setComputePipelineState:pipeline_];
103 }
104
105 id<MTLComputePipelineState> GetPipeline() const { return pipeline_; }
106
107 void SetBuffer(uint64_t index, uint64_t offset, id<MTLBuffer> buffer) {
108 auto found = buffers_.find(index);
109 if (found != buffers_.end() && found->second.buffer == buffer) {
110 // The right buffer is bound. Check if its offset needs to be updated.
111 if (found->second.offset == offset) {
112 // Buffer and its offset is identical. Nothing to do.
113 return;
114 }
115
116 // Only the offset needs to be updated.
117 found->second.offset = offset;
118
119 [encoder_ setBufferOffset:offset atIndex:index];
120 return;
121 }
122
123 buffers_[index] = {buffer, static_cast<size_t>(offset)};
124 [encoder_ setBuffer:buffer offset:offset atIndex:index];
125 }
126
127 void SetTexture(uint64_t index, id<MTLTexture> texture) {
128 auto found = textures_.find(index);
129 if (found != textures_.end() && found->second == texture) {
130 // Already bound.
131 return;
132 }
133 textures_[index] = texture;
134 [encoder_ setTexture:texture atIndex:index];
135 return;
136 }
137
138 void SetSampler(uint64_t index, id<MTLSamplerState> sampler) {
139 auto found = samplers_.find(index);
140 if (found != samplers_.end() && found->second == sampler) {
141 // Already bound.
142 return;
143 }
144 samplers_[index] = sampler;
145 [encoder_ setSamplerState:sampler atIndex:index];
146 return;
147 }
148
149 private:
150 struct BufferOffsetPair {
151 id<MTLBuffer> buffer = nullptr;
152 size_t offset = 0u;
153 };
154 using BufferMap = std::map<uint64_t, BufferOffsetPair>;
155 using TextureMap = std::map<uint64_t, id<MTLTexture>>;
156 using SamplerMap = std::map<uint64_t, id<MTLSamplerState>>;
157
158 const id<MTLComputeCommandEncoder> encoder_;
159 id<MTLComputePipelineState> pipeline_ = nullptr;
160 BufferMap buffers_;
161 TextureMap textures_;
162 SamplerMap samplers_;
163 };
164
165 static bool Bind(ComputePassBindingsCache& pass,
166 Allocator& allocator,
167 size_t bind_index,
168 const BufferView& view) {
169 if (!view.buffer) {
170 return false;
171 }
172
173 auto device_buffer = view.buffer->GetDeviceBuffer(allocator);
174 if (!device_buffer) {
175 return false;
176 }
177
178 auto buffer = DeviceBufferMTL::Cast(*device_buffer).GetMTLBuffer();
179 // The Metal call is a void return and we don't want to make it on nil.
180 if (!buffer) {
181 return false;
182 }
183
184 pass.SetBuffer(bind_index, view.range.offset, buffer);
185 return true;
186 }
187
188 static bool Bind(ComputePassBindingsCache& pass,
189 size_t bind_index,
190 const Sampler& sampler,
191 const Texture& texture) {
192 if (!sampler.IsValid() || !texture.IsValid()) {
193 return false;
194 }
195
196 pass.SetTexture(bind_index, TextureMTL::Cast(texture).GetMTLTexture());
197 pass.SetSampler(bind_index, SamplerMTL::Cast(sampler).GetMTLSamplerState());
198 return true;
199 }
200
201 bool ComputePassMTL::EncodeCommands(const std::shared_ptr<Allocator>& allocator,
202 id<MTLComputeCommandEncoder> encoder,
203 const ISize& grid_size,
204 const ISize& thread_group_size) const {
205 if (grid_size.width == 0 || grid_size.height == 0) {
206 return true;
207 }
208
209 ComputePassBindingsCache pass_bindings(encoder);
210
211 fml::closure pop_debug_marker = [encoder]() { [encoder popDebugGroup]; };
212 for (const auto& command : commands_) {
213 #ifdef IMPELLER_DEBUG
214 fml::ScopedCleanupClosure auto_pop_debug_marker(pop_debug_marker);
215 if (!command.label.empty()) {
216 [encoder pushDebugGroup:@(command.label.c_str())];
217 } else {
218 auto_pop_debug_marker.Release();
219 }
220 #endif
221
222 pass_bindings.SetComputePipelineState(
223 ComputePipelineMTL::Cast(*command.pipeline)
224 .GetMTLComputePipelineState());
225
226 for (const auto& buffer : command.bindings.buffers) {
227 if (!Bind(pass_bindings, *allocator, buffer.first,
228 buffer.second.view.resource)) {
229 return false;
230 }
231 }
232
233 for (const auto& data : command.bindings.sampled_images) {
234 if (!Bind(pass_bindings, data.first, *data.second.sampler.resource,
235 *data.second.texture.resource)) {
236 return false;
237 }
238 }
239
240 // TODO(dnfield): use feature detection to support non-uniform threadgroup
241 // sizes.
242 // https://github.com/flutter/flutter/issues/110619
243 auto width = grid_size.width;
244 auto height = grid_size.height;
245
246 auto maxTotalThreadsPerThreadgroup = static_cast<int64_t>(
247 pass_bindings.GetPipeline().maxTotalThreadsPerThreadgroup);
248
249 // Special case for linear processing.
250 if (height == 1) {
251 int64_t threadGroups = std::max(
252 static_cast<int64_t>(
253 std::ceil(width * 1.0 / maxTotalThreadsPerThreadgroup * 1.0)),
254 1LL);
255 [encoder dispatchThreadgroups:MTLSizeMake(threadGroups, 1, 1)
256 threadsPerThreadgroup:MTLSizeMake(maxTotalThreadsPerThreadgroup,
257 1, 1)];
258 } else {
259 while (width * height > maxTotalThreadsPerThreadgroup) {
260 width = std::max(1LL, width / 2);
261 height = std::max(1LL, height / 2);
262 }
263
264 auto size = MTLSizeMake(width, height, 1);
265 [encoder dispatchThreadgroups:size threadsPerThreadgroup:size];
266 }
267 }
268
269 return true;
270 }
271
272 } // namespace impeller
host_buffer.h
impeller::Range::offset
size_t offset
Definition: range.h:14
compute_pipeline_mtl.h
formats.h
formats_mtl.h
impeller::ComputePass::commands_
std::vector< ComputeCommand > commands_
Definition: compute_pass.h:62
impeller::BufferView::range
Range range
Definition: buffer_view.h:16
impeller::ComputePassBindingsCache::SetTexture
void SetTexture(uint64_t index, id< MTLTexture > texture)
Definition: compute_pass_mtl.mm:127
impeller::Texture::IsValid
virtual bool IsValid() const =0
impeller::ComputePassBindingsCache::SetComputePipelineState
void SetComputePipelineState(id< MTLComputePipelineState > pipeline)
Definition: compute_pass_mtl.mm:97
impeller::ComputePass::EncodeCommands
bool EncodeCommands() const
Encode the recorded commands to the underlying command buffer.
Definition: compute_pass.cc:50
impeller::Texture
Definition: texture.h:17
impeller::Sampler
Definition: sampler.h:12
impeller::ComputePassBindingsCache::ComputePassBindingsCache
ComputePassBindingsCache(id< MTLComputeCommandEncoder > encoder)
Definition: compute_pass_mtl.mm:90
impeller::ComputePassMTL::~ComputePassMTL
~ComputePassMTL() override
impeller::ComputePipelineMTL::GetMTLComputePipelineState
id< MTLComputePipelineState > GetMTLComputePipelineState() const
Definition: compute_pipeline_mtl.mm:25
impeller::BufferView::buffer
std::shared_ptr< const Buffer > buffer
Definition: buffer_view.h:14
backend_cast.h
impeller::Allocator
An object that allocates device memory.
Definition: allocator.h:25
impeller::ComputePassBindingsCache
Ensures that bindings on the pass are not redundantly set or updated. Avoids making the driver do add...
Definition: compute_pass_mtl.mm:89
impeller::ComputePassBindingsCache::SetSampler
void SetSampler(uint64_t index, id< MTLSamplerState > sampler)
Definition: compute_pass_mtl.mm:138
impeller::DeviceBufferMTL::GetMTLBuffer
id< MTLBuffer > GetMTLBuffer() const
Definition: device_buffer_mtl.mm:21
impeller::ISize
TSize< int64_t > ISize
Definition: size.h:136
impeller::BufferView
Definition: buffer_view.h:13
impeller::ComputePassBindingsCache::SetBuffer
void SetBuffer(uint64_t index, uint64_t offset, id< MTLBuffer > buffer)
Definition: compute_pass_mtl.mm:107
std
Definition: comparable.h:98
impeller::Sampler::IsValid
virtual bool IsValid() const =0
texture_mtl.h
impeller::BackendCast< DeviceBufferMTL, Buffer >::Cast
static DeviceBufferMTL & Cast(Buffer &base)
Definition: backend_cast.h:14
impeller::ComputePassBindingsCache::GetPipeline
id< MTLComputePipelineState > GetPipeline() const
Definition: compute_pass_mtl.mm:105
compute_command.h
compute_pass_mtl.h
impeller::Bind
static bool Bind(ComputePassBindingsCache &pass, Allocator &allocator, size_t bind_index, const BufferView &view)
Definition: compute_pass_mtl.mm:165
sampler_mtl.h
shader_types.h
impeller
Definition: aiks_context.cc:10
device_buffer_mtl.h