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