1 /*
2 * Copyright 2023 Google LLC
3 *
4 * Use of this source code is governed by a BSD-style license that can be
5 * found in the LICENSE file.
6 */
7
8 #include "src/gpu/graphite/compute/VelloComputeSteps.h"
9
10 namespace skgpu::graphite {
11
VelloStageName(vello_cpp::ShaderStage stage)12 std::string_view VelloStageName(vello_cpp::ShaderStage stage) {
13 auto name = vello_cpp::shader(stage).name();
14 return {name.data(), name.length()};
15 }
16
VelloStageLocalSize(vello_cpp::ShaderStage stage)17 WorkgroupSize VelloStageLocalSize(vello_cpp::ShaderStage stage) {
18 auto wgSize = vello_cpp::shader(stage).workgroup_size();
19 return WorkgroupSize(wgSize.x, wgSize.y, wgSize.z);
20 }
21
VelloWorkgroupBuffers(vello_cpp::ShaderStage stage)22 skia_private::TArray<ComputeStep::WorkgroupBufferDesc> VelloWorkgroupBuffers(
23 vello_cpp::ShaderStage stage) {
24 auto wgBuffers = vello_cpp::shader(stage).workgroup_buffers();
25 skia_private::TArray<ComputeStep::WorkgroupBufferDesc> result;
26 if (!wgBuffers.empty()) {
27 result.reserve(wgBuffers.size());
28 for (const auto& desc : wgBuffers) {
29 result.push_back({desc.size_in_bytes, desc.index});
30 }
31 }
32 return result;
33 }
34
VelloNativeShaderSource(vello_cpp::ShaderStage stage,ComputeStep::NativeShaderFormat format)35 ComputeStep::NativeShaderSource VelloNativeShaderSource(vello_cpp::ShaderStage stage,
36 ComputeStep::NativeShaderFormat format) {
37 using NativeFormat = ComputeStep::NativeShaderFormat;
38
39 const auto& shader = vello_cpp::shader(stage);
40 ::rust::Str source;
41 std::string entryPoint;
42 switch (format) {
43 #if VELLO_WGSL_SHADERS
44 case NativeFormat::kWGSL:
45 source = shader.wgsl();
46 entryPoint = "main";
47 break;
48 #endif
49 #if VELLO_MSL_SHADERS
50 case NativeFormat::kMSL:
51 source = shader.msl();
52 entryPoint = "main_";
53 break;
54 #endif
55 default:
56 return {std::string_view(), ""};
57 }
58
59 return {{source.data(), source.length()}, std::move(entryPoint)};
60 }
61
62 #define BUFFER_BINDING(slot, type, policy) \
63 { \
64 /*type=*/ComputeStep::ResourceType::k##type##Buffer, \
65 /*flow=*/ComputeStep::DataFlow::kShared, \
66 /*policy=*/ComputeStep::ResourcePolicy::k##policy, \
67 /*slot=*/kVelloSlot_##slot, \
68 }
69
70 #define TEXTURE_BINDING(slot, type, policy) \
71 { \
72 /*type=*/ComputeStep::ResourceType::k##type##Texture, \
73 /*flow=*/ComputeStep::DataFlow::kShared, \
74 /*policy=*/ComputeStep::ResourcePolicy::k##policy, \
75 /*slot=*/kVelloSlot_##slot, \
76 }
77
78 // PathtagReduce
VelloPathtagReduceStep()79 VelloPathtagReduceStep::VelloPathtagReduceStep()
80 : VelloStep({
81 BUFFER_BINDING(ConfigUniform, Uniform, Mapped),
82 BUFFER_BINDING(Scene, ReadOnlyStorage, Mapped),
83 BUFFER_BINDING(PathtagReduceOutput, Storage, None),
84 }) {}
85
86 // PathtagScanSmall
VelloPathtagScanSmallStep()87 VelloPathtagScanSmallStep::VelloPathtagScanSmallStep()
88 : VelloStep({
89 BUFFER_BINDING(ConfigUniform, Uniform, None),
90 BUFFER_BINDING(Scene, ReadOnlyStorage, None),
91 BUFFER_BINDING(PathtagReduceOutput, ReadOnlyStorage, None),
92 BUFFER_BINDING(TagMonoid, Storage, None),
93 }) {}
94
95 // PathtagReduce2
VelloPathtagReduce2Step()96 VelloPathtagReduce2Step::VelloPathtagReduce2Step()
97 : VelloStep({
98 BUFFER_BINDING(LargePathtagReduceFirstPassOutput, ReadOnlyStorage, None),
99 BUFFER_BINDING(LargePathtagReduceSecondPassOutput, Storage, None),
100 }) {}
101
102 // PathtagScan1
VelloPathtagScan1Step()103 VelloPathtagScan1Step::VelloPathtagScan1Step()
104 : VelloStep({
105 BUFFER_BINDING(LargePathtagReduceFirstPassOutput, ReadOnlyStorage, None),
106 BUFFER_BINDING(LargePathtagReduceSecondPassOutput, ReadOnlyStorage, None),
107 BUFFER_BINDING(LargePathtagScanFirstPassOutput, Storage, None),
108 }) {}
109
110 // PathtagScanLarge
VelloPathtagScanLargeStep()111 VelloPathtagScanLargeStep::VelloPathtagScanLargeStep()
112 : VelloStep({
113 BUFFER_BINDING(ConfigUniform, Uniform, None),
114 BUFFER_BINDING(Scene, ReadOnlyStorage, None),
115 BUFFER_BINDING(LargePathtagScanFirstPassOutput, ReadOnlyStorage, None),
116 BUFFER_BINDING(TagMonoid, Storage, None),
117 }) {}
118
119 // BboxClear
VelloBboxClearStep()120 VelloBboxClearStep::VelloBboxClearStep()
121 : VelloStep({
122 BUFFER_BINDING(ConfigUniform, Uniform, None),
123 BUFFER_BINDING(PathBBoxes, Storage, None),
124 }) {}
125
126 // Flatten
VelloFlattenStep()127 VelloFlattenStep::VelloFlattenStep()
128 : VelloStep({
129 BUFFER_BINDING(ConfigUniform, Uniform, None),
130 BUFFER_BINDING(Scene, ReadOnlyStorage, None),
131 BUFFER_BINDING(TagMonoid, ReadOnlyStorage, None),
132 BUFFER_BINDING(PathBBoxes, Storage, None),
133 BUFFER_BINDING(BumpAlloc, Storage, Clear),
134 BUFFER_BINDING(Lines, Storage, None),
135 }) {}
136
137 // DrawReduce
VelloDrawReduceStep()138 VelloDrawReduceStep::VelloDrawReduceStep()
139 : VelloStep({
140 BUFFER_BINDING(ConfigUniform, Uniform, None),
141 BUFFER_BINDING(Scene, ReadOnlyStorage, None),
142 BUFFER_BINDING(DrawReduceOutput, Storage, None),
143 }) {}
144
145 // DrawLeaf
VelloDrawLeafStep()146 VelloDrawLeafStep::VelloDrawLeafStep()
147 : VelloStep({
148 BUFFER_BINDING(ConfigUniform, Uniform, None),
149 BUFFER_BINDING(Scene, ReadOnlyStorage, None),
150 BUFFER_BINDING(DrawReduceOutput, ReadOnlyStorage, None),
151 BUFFER_BINDING(PathBBoxes, ReadOnlyStorage, None),
152 BUFFER_BINDING(DrawMonoid, Storage, None),
153 BUFFER_BINDING(InfoBinData, Storage, None),
154 BUFFER_BINDING(ClipInput, Storage, None),
155 }) {}
156
157 // ClipReduce
VelloClipReduceStep()158 VelloClipReduceStep::VelloClipReduceStep()
159 : VelloStep({
160 BUFFER_BINDING(ClipInput, ReadOnlyStorage, None),
161 BUFFER_BINDING(PathBBoxes, ReadOnlyStorage, None),
162 BUFFER_BINDING(ClipBicyclic, Storage, None),
163 BUFFER_BINDING(ClipElement, Storage, None),
164 }) {}
165
166 // ClipLeaf
VelloClipLeafStep()167 VelloClipLeafStep::VelloClipLeafStep()
168 : VelloStep({
169 BUFFER_BINDING(ConfigUniform, Uniform, None),
170 BUFFER_BINDING(ClipInput, ReadOnlyStorage, None),
171 BUFFER_BINDING(PathBBoxes, ReadOnlyStorage, None),
172 BUFFER_BINDING(ClipBicyclic, ReadOnlyStorage, None),
173 BUFFER_BINDING(ClipElement, ReadOnlyStorage, None),
174 BUFFER_BINDING(DrawMonoid, Storage, None),
175 BUFFER_BINDING(ClipBBoxes, Storage, None),
176 }) {}
177
178 // Binning
VelloBinningStep()179 VelloBinningStep::VelloBinningStep()
180 : VelloStep({
181 BUFFER_BINDING(ConfigUniform, Uniform, None),
182 BUFFER_BINDING(DrawMonoid, ReadOnlyStorage, None),
183 BUFFER_BINDING(PathBBoxes, ReadOnlyStorage, None),
184 BUFFER_BINDING(ClipBBoxes, ReadOnlyStorage, None),
185 BUFFER_BINDING(DrawBBoxes, Storage, None),
186 BUFFER_BINDING(BumpAlloc, Storage, None),
187 BUFFER_BINDING(InfoBinData, Storage, None),
188 BUFFER_BINDING(BinHeader, Storage, None),
189 }) {}
190
191 // TileAlloc
VelloTileAllocStep()192 VelloTileAllocStep::VelloTileAllocStep()
193 : VelloStep({
194 BUFFER_BINDING(ConfigUniform, Uniform, None),
195 BUFFER_BINDING(Scene, ReadOnlyStorage, None),
196 BUFFER_BINDING(DrawBBoxes, ReadOnlyStorage, None),
197 BUFFER_BINDING(BumpAlloc, Storage, None),
198 BUFFER_BINDING(Path, Storage, None),
199 BUFFER_BINDING(Tile, Storage, None),
200 }) {}
201
202 // PathCountSetup
VelloPathCountSetupStep()203 VelloPathCountSetupStep::VelloPathCountSetupStep()
204 : VelloStep({
205 BUFFER_BINDING(BumpAlloc, Storage, None),
206 BUFFER_BINDING(IndirectCount, Storage, None),
207 }) {}
208
209 // PathCount
VelloPathCountStep()210 VelloPathCountStep::VelloPathCountStep()
211 : VelloStep({
212 BUFFER_BINDING(ConfigUniform, Uniform, None),
213 BUFFER_BINDING(BumpAlloc, Storage, None),
214 BUFFER_BINDING(Lines, ReadOnlyStorage, None),
215 BUFFER_BINDING(Path, ReadOnlyStorage, None),
216 BUFFER_BINDING(Tile, Storage, None),
217 BUFFER_BINDING(SegmentCounts, Storage, None),
218 }) {}
219
220 // BackdropDyn
VelloBackdropDynStep()221 VelloBackdropDynStep::VelloBackdropDynStep()
222 : VelloStep({
223 BUFFER_BINDING(ConfigUniform, Uniform, None),
224 BUFFER_BINDING(BumpAlloc, Storage, None),
225 BUFFER_BINDING(Path, ReadOnlyStorage, None),
226 BUFFER_BINDING(Tile, Storage, None),
227 }) {}
228
229 // Coarse
VelloCoarseStep()230 VelloCoarseStep::VelloCoarseStep()
231 : VelloStep({
232 BUFFER_BINDING(ConfigUniform, Uniform, None),
233 BUFFER_BINDING(Scene, ReadOnlyStorage, None),
234 BUFFER_BINDING(DrawMonoid, ReadOnlyStorage, None),
235 BUFFER_BINDING(BinHeader, ReadOnlyStorage, None),
236 BUFFER_BINDING(InfoBinData, ReadOnlyStorage, None),
237 BUFFER_BINDING(Path, ReadOnlyStorage, None),
238 BUFFER_BINDING(Tile, Storage, None),
239 BUFFER_BINDING(BumpAlloc, Storage, None),
240 BUFFER_BINDING(PTCL, Storage, None),
241 }) {}
242
243 // PathTilingSetup
VelloPathTilingSetupStep()244 VelloPathTilingSetupStep::VelloPathTilingSetupStep()
245 : VelloStep({
246 BUFFER_BINDING(BumpAlloc, Storage, None),
247 BUFFER_BINDING(IndirectCount, Storage, None),
248 BUFFER_BINDING(PTCL, Storage, None),
249 }) {}
250
251 // PathTiling
VelloPathTilingStep()252 VelloPathTilingStep::VelloPathTilingStep()
253 : VelloStep({
254 BUFFER_BINDING(BumpAlloc, Storage, None),
255 BUFFER_BINDING(SegmentCounts, ReadOnlyStorage, None),
256 BUFFER_BINDING(Lines, ReadOnlyStorage, None),
257 BUFFER_BINDING(Path, ReadOnlyStorage, None),
258 BUFFER_BINDING(Tile, ReadOnlyStorage, None),
259 BUFFER_BINDING(Segments, Storage, None),
260 }) {}
261
262 // Fine
263 static constexpr ComputeStep::ResourceDesc kFineAreaResources[] = {
264 BUFFER_BINDING(ConfigUniform, Uniform, None),
265 BUFFER_BINDING(Segments, ReadOnlyStorage, None),
266 BUFFER_BINDING(PTCL, ReadOnlyStorage, None),
267 BUFFER_BINDING(InfoBinData, ReadOnlyStorage, None),
268 TEXTURE_BINDING(OutputImage, WriteOnlyStorage, None),
269 };
270
271 static constexpr ComputeStep::ResourceDesc kFineMsaaResources[] = {
272 BUFFER_BINDING(ConfigUniform, Uniform, None),
273 BUFFER_BINDING(Segments, ReadOnlyStorage, None),
274 BUFFER_BINDING(PTCL, ReadOnlyStorage, None),
275 BUFFER_BINDING(InfoBinData, ReadOnlyStorage, None),
276 TEXTURE_BINDING(OutputImage, WriteOnlyStorage, None),
277 BUFFER_BINDING(MaskLUT, ReadOnlyStorage, Mapped),
278 };
279
VelloFineAreaStep()280 VelloFineAreaStep::VelloFineAreaStep() : VelloFineStepBase(kFineAreaResources) {}
281
VelloFineMsaa16Step()282 VelloFineMsaa16Step::VelloFineMsaa16Step() : VelloFineMsaaStepBase(kFineMsaaResources) {}
283
VelloFineMsaa8Step()284 VelloFineMsaa8Step::VelloFineMsaa8Step() : VelloFineMsaaStepBase(kFineMsaaResources) {}
285
VelloFineAreaAlpha8Step()286 VelloFineAreaAlpha8Step::VelloFineAreaAlpha8Step() : VelloFineStepBase(kFineAreaResources) {}
287
VelloFineMsaa16Alpha8Step()288 VelloFineMsaa16Alpha8Step::VelloFineMsaa16Alpha8Step()
289 : VelloFineMsaaStepBase(kFineMsaaResources) {}
290
VelloFineMsaa8Alpha8Step()291 VelloFineMsaa8Alpha8Step::VelloFineMsaa8Alpha8Step() : VelloFineMsaaStepBase(kFineMsaaResources) {}
292
293 } // namespace skgpu::graphite
294