| Flutter Impeller |

spirv_sksl.cc
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/compiler/spirv_sksl.h"
6 #include "impeller/compiler/uniform_sorter.h"
7
8 using namespace spv;
9 using namespace SPIRV_CROSS_NAMESPACE;
10
11 namespace impeller {
12 namespace compiler {
13
14 // This replaces the SPIRV_CROSS_THROW which aborts and drops the
15 // error message in non-debug modes.
16 void report_and_exit(const std::string& msg) {
17 fprintf(stderr, "There was a compiler error: %s\n", msg.c_str());
18 fflush(stderr);
19 exit(1);
20 }
21
22 #define FLUTTER_CROSS_THROW(x) report_and_exit(x)
23
24 std::string CompilerSkSL::compile() {
25 ir.fixup_reserved_names();
26
27 if (get_execution_model() != ExecutionModelFragment) {
28 FLUTTER_CROSS_THROW("Only fragment shaders are supported.'");
29 return "";
30 }
31
32 options.es = false;
33 options.version = 100;
34 options.vulkan_semantics = false;
35 options.enable_420pack_extension = false;
36 options.flatten_multidimensional_arrays = true;
37
38 backend.allow_precision_qualifiers = false;
39 backend.basic_int16_type = "short";
40 backend.basic_int_type = "int";
41 backend.basic_uint16_type = "ushort";
42 backend.basic_uint_type = "uint";
43 backend.double_literal_suffix = false;
44 backend.float_literal_suffix = false;
45 backend.long_long_literal_suffix = false;
46 backend.needs_row_major_load_workaround = true;
47 backend.nonuniform_qualifier = "";
48 backend.support_precise_qualifier = false;
49 backend.uint32_t_literal_suffix = false;
50 backend.use_array_constructor = true;
51 backend.workgroup_size_is_hidden = true;
52
53 fixup_user_functions();
54
55 fixup_anonymous_struct_names();
56 fixup_type_alias();
57 reorder_type_alias();
58 build_function_control_flow_graphs_and_analyze();
59 fixup_image_load_store_access();
60 update_active_builtins();
61 analyze_image_and_sampler_usage();
62 analyze_interlocked_resource_usage();
63
64 uint32_t pass_count = 0;
65 do {
66 reset(pass_count);
67
68 // Move constructor for this type is broken on GCC 4.9 ...
69 buffer.reset();
70
71 emit_header();
72 emit_resources();
73
74 emit_function(get<SPIRFunction>(ir.default_entry_point), Bitset());
75
76 pass_count++;
77 } while (is_forcing_recompilation());
78
79 statement("half4 main(float2 iFragCoord)");
80 begin_scope();
81 statement(" flutter_FragCoord = float4(iFragCoord, 0, 0);");
82 statement(" FLT_main();");
83 statement(" return " + output_name_ + ";");
84 end_scope();
85
86 return buffer.str();
87 }
88
89 void CompilerSkSL::fixup_user_functions() {
90 const std::string prefix = "FLT_flutter_local_";
91 ir.for_each_typed_id<SPIRFunction>([&](uint32_t, const SPIRFunction& func) {
92 const auto& original_name = get_name(func.self);
93 // Just in case. Don't add the prefix a second time.
94 if (original_name.rfind(prefix, 0) == 0) {
95 return;
96 }
97 std::string new_name = prefix + original_name;
98 set_name(func.self, new_name);
99 });
100
101 ir.for_each_typed_id<SPIRFunctionPrototype>(
102 [&](uint32_t, const SPIRFunctionPrototype& func) {
103 const auto& original_name = get_name(func.self);
104 // Just in case. Don't add the prefix a second time.
105 if (original_name.rfind(prefix, 0) == 0) {
106 return;
107 }
108 std::string new_name = prefix + original_name;
109 set_name(func.self, new_name);
110 });
111 }
112
113 void CompilerSkSL::emit_header() {
114 statement("// This SkSL shader is autogenerated by spirv-cross.");
115 statement("");
116 statement("float4 flutter_FragCoord;");
117 statement("");
118 }
119
120 void CompilerSkSL::emit_uniform(const SPIRVariable& var) {
121 auto& type = get<SPIRType>(var.basetype);
122 if (type.basetype == SPIRType::UInt && is_legacy()) {
123 FLUTTER_CROSS_THROW("SkSL does not support unsigned integers: '" +
124 get_name(var.self) + "'");
125 }
126
127 add_resource_name(var.self);
128 statement(variable_decl(var), ";");
129
130 // The Flutter FragmentProgram implementation passes additional uniforms along
131 // with shader uniforms that encode the shader width and height.
132 if (type.basetype == SPIRType::SampledImage) {
133 std::string name = to_name(var.self);
134 statement("uniform half2 " + name + "_size;");
135 }
136 }
137
138 bool CompilerSkSL::emit_constant_resources() {
139 bool emitted = false;
140
141 for (auto& id : ir.ids) {
142 if (id.get_type() == TypeConstant) {
143 auto& c = id.get<SPIRConstant>();
144 bool needs_declaration = c.specialization || c.is_used_as_lut;
145 if (needs_declaration) {
146 if (!options.vulkan_semantics && c.specialization) {
147 c.specialization_constant_macro_name = constant_value_macro_name(
148 get_decoration(c.self, DecorationSpecId));
149 }
150 emit_constant(c);
151 emitted = true;
152 }
153 } else if (id.get_type() == TypeConstantOp) {
154 emit_specialization_constant_op(id.get<SPIRConstantOp>());
155 emitted = true;
156 }
157 }
158
159 return emitted;
160 }
161
162 bool CompilerSkSL::emit_struct_resources() {
163 bool emitted = false;
164
165 // Output all basic struct types which are not Block or BufferBlock as these
166 // are declared inplace when such variables are instantiated.
167 for (auto& id : ir.ids) {
168 if (id.get_type() == TypeType) {
169 auto& type = id.get<SPIRType>();
170 if (type.basetype == SPIRType::Struct && type.array.empty() &&
171 !type.pointer &&
172 (!ir.meta[type.self].decoration.decoration_flags.get(
173 DecorationBlock) &&
174 !ir.meta[type.self].decoration.decoration_flags.get(
175 DecorationBufferBlock))) {
176 emit_struct(type);
177 emitted = true;
178 }
179 }
180 }
181
182 return emitted;
183 }
184
185 void CompilerSkSL::detect_unsupported_resources() {
186 for (auto& id : ir.ids) {
187 if (id.get_type() == TypeVariable) {
188 auto& var = id.get<SPIRVariable>();
189 auto& type = get<SPIRType>(var.basetype);
190
191 // UBOs and SSBOs are not supported.
192 if (var.storage != StorageClassFunction && type.pointer &&
193 type.storage == StorageClassUniform && !is_hidden_variable(var) &&
194 (ir.meta[type.self].decoration.decoration_flags.get(
195 DecorationBlock) ||
196 ir.meta[type.self].decoration.decoration_flags.get(
197 DecorationBufferBlock))) {
198 FLUTTER_CROSS_THROW("SkSL does not support UBOs or SSBOs: '" +
199 get_name(var.self) + "'");
200 }
201
202 // Push constant blocks are not supported.
203 if (!is_hidden_variable(var) && var.storage != StorageClassFunction &&
204 type.pointer && type.storage == StorageClassPushConstant) {
205 FLUTTER_CROSS_THROW("SkSL does not support push constant blocks: '" +
206 get_name(var.self) + "'");
207 }
208
209 // User specified inputs are not supported.
210 if (!is_hidden_variable(var) && var.storage != StorageClassFunction &&
211 type.pointer && type.storage == StorageClassInput) {
212 FLUTTER_CROSS_THROW("SkSL does not support inputs: '" +
213 get_name(var.self) + "'");
214 }
215 } else if (id.get_type() == TypeBlock) {
216 // Array initializers are not supported. Check TypeBlock IDs to detect
217 // this.
218 auto& block = id.get<SPIRBlock>();
219
220 // File and line information for use in the error message.
221 std::string file;
222 uint32_t line = 0;
223
224 for (auto instruction : block.ops) {
225 bool has_array_initializer = false;
226 if (instruction.op == OpLine) {
227 // OpLine information applies to all subsequent instructions until the
228 // next OpLine instruction.
229 file = get<SPIRString>(ir.spirv[instruction.offset]).str;
230 line = ir.spirv[instruction.offset + 1];
231 } else if (instruction.op == OpStore) {
232 // Check for OpStore instructions that store an array constant. This
233 // detects array initializations which use compile-time constants
234 // (e.g. `float[2] nums = float[](1.0, 2.0);`).
235 Variant& store_object_id = ir.ids[ir.spirv[instruction.offset + 1]];
236 if (store_object_id.get_type() == TypeConstant) {
237 auto& c = store_object_id.get<SPIRConstant>();
238 auto& type = get<SPIRType>(c.constant_type);
239 has_array_initializer = !type.array.empty();
240 }
241 } else if (instruction.op == OpCompositeConstruct) {
242 // Check for OpCompositeConstruct instructions with an array result.
243 // This detects array initializations which use variables
244 // (e.g. `float var = 2.0; float[2] nums = float[](1.0, var);`).
245 auto result_type_id = ir.spirv[instruction.offset];
246 auto& type = get<SPIRType>(result_type_id);
247 has_array_initializer = !type.array.empty();
248 }
249
250 if (has_array_initializer) {
251 FLUTTER_CROSS_THROW("SkSL does not support array initializers: " +
252 file + ":" + std::to_string(line));
253 }
254 }
255 }
256 }
257 }
258
259 bool CompilerSkSL::emit_uniform_resources() {
260 bool emitted = false;
261
262 // Output Uniform Constants (values, samplers, images, etc).
263 std::vector<ID> regular_uniforms =
264 SortUniforms(&ir, this, SPIRType::SampledImage, /*include=*/false);
265 std::vector<ID> shader_uniforms =
266 SortUniforms(&ir, this, SPIRType::SampledImage);
267 if (regular_uniforms.size() > 0 || shader_uniforms.size() > 0) {
268 emitted = true;
269 }
270
271 for (const auto& id : regular_uniforms) {
272 auto& var = get<SPIRVariable>(id);
273 emit_uniform(var);
274 }
275
276 for (const auto& id : shader_uniforms) {
277 auto& var = get<SPIRVariable>(id);
278 emit_uniform(var);
279 }
280
281 return emitted;
282 }
283
284 bool CompilerSkSL::emit_output_resources() {
285 bool emitted = false;
286
287 // Output 'out' variables. These are restricted to the cases handled by
288 // SkSL in 'emit_interface_block'.
289 for (auto& id : ir.ids) {
290 if (id.get_type() == TypeVariable) {
291 auto& var = id.get<SPIRVariable>();
292 auto& type = get<SPIRType>(var.basetype);
293 if (var.storage != StorageClassFunction && !is_hidden_variable(var) &&
294 type.pointer &&
295 (var.storage == StorageClassInput ||
296 var.storage == StorageClassOutput) &&
297 interface_variable_exists_in_entry_point(var.self)) {
298 emit_interface_block(var);
299 emitted = true;
300 }
301 }
302 }
303
304 return emitted;
305 }
306
307 bool CompilerSkSL::emit_global_variable_resources() {
308 bool emitted = false;
309
310 for (auto global : global_variables) {
311 auto& var = get<SPIRVariable>(global);
312 if (is_hidden_variable(var, true)) {
313 continue;
314 }
315 if (var.storage != StorageClassOutput) {
316 if (!variable_is_lut(var)) {
317 add_resource_name(var.self);
318 std::string initializer;
319 if (options.force_zero_initialized_variables &&
320 var.storage == StorageClassPrivate && !var.initializer &&
321 !var.static_expression &&
322 type_can_zero_initialize(get_variable_data_type(var))) {
323 initializer = join(" = ", to_zero_initialized_expression(
324 get_variable_data_type_id(var)));
325 }
326 statement(variable_decl(var), initializer, ";");
327 emitted = true;
328 }
329 } else if (var.initializer &&
330 maybe_get<SPIRConstant>(var.initializer) != nullptr) {
331 emit_output_variable_initializer(var);
332 }
333 }
334
335 return emitted;
336 }
337
338 bool CompilerSkSL::emit_undefined_values() {
339 bool emitted = false;
340
341 ir.for_each_typed_id<SPIRUndef>([&](uint32_t, const SPIRUndef& undef) {
342 auto& type = this->get<SPIRType>(undef.basetype);
343 // OpUndef can be void for some reason ...
344 if (type.basetype == SPIRType::Void) {
345 return;
346 }
347
348 std::string initializer;
349 if (options.force_zero_initialized_variables &&
350 type_can_zero_initialize(type)) {
351 initializer = join(" = ", to_zero_initialized_expression(undef.basetype));
352 }
353
354 statement(variable_decl(type, to_name(undef.self), undef.self), initializer,
355 ";");
356 emitted = true;
357 });
358
359 return emitted;
360 }
361
362 void CompilerSkSL::emit_resources() {
363 detect_unsupported_resources();
364
365 if (emit_constant_resources()) {
366 statement("");
367 }
368
369 if (emit_struct_resources()) {
370 statement("");
371 }
372
373 if (emit_uniform_resources()) {
374 statement("");
375 }
376
377 if (emit_output_resources()) {
378 statement("");
379 }
380
381 if (emit_global_variable_resources()) {
382 statement("");
383 }
384
385 if (emit_undefined_values()) {
386 statement("");
387 }
388 }
389
390 void CompilerSkSL::emit_interface_block(const SPIRVariable& var) {
391 auto& type = get<SPIRType>(var.basetype);
392 bool block =
393 ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock);
394 if (block) {
395 FLUTTER_CROSS_THROW("Interface blocks are not supported: '" +
396 to_name(var.self) + "'");
397 }
398
399 // The output is emitted as a global variable, which is returned from the
400 // wrapper around the 'main' function. Only one output variable is allowed.
401 add_resource_name(var.self);
402 statement(variable_decl(type, to_name(var.self), var.self), ";");
403 if (output_name_.empty()) {
404 output_name_ = to_name(var.self);
405 } else if (to_name(var.self) != output_name_) {
406 FLUTTER_CROSS_THROW("Only one output variable is supported: '" +
407 to_name(var.self) + "'");
408 }
409 }
410
411 void CompilerSkSL::emit_function_prototype(SPIRFunction& func,
412 const Bitset& return_flags) {
413 // If this is not the entrypoint, then no special processsing for SkSL is
414 // required.
415 if (func.self != ir.default_entry_point) {
416 CompilerGLSL::emit_function_prototype(func, return_flags);
417 return;
418 }
419
420 auto& type = get<SPIRType>(func.return_type);
421 if (type.basetype != SPIRType::Void) {
422 FLUTTER_CROSS_THROW(
423 "Return type of the entrypoint function must be 'void'");
424 }
425
426 if (func.arguments.size() != 0) {
427 FLUTTER_CROSS_THROW(
428 "The entry point function should not acept any parameters.");
429 }
430
431 processing_entry_point = true;
432
433 // If this is the entrypoint of a fragment shader, then GLSL requires the
434 // prototype to be "void main()", and so it is safe to rewrite as
435 // "void FLT_main()".
436 statement("void FLT_main()");
437 }
438
439 std::string CompilerSkSL::image_type_glsl(const SPIRType& type,
440 uint32_t id,
441 bool member) {
442 if (type.basetype != SPIRType::SampledImage || type.image.dim != Dim2D) {
443 FLUTTER_CROSS_THROW("Only sampler2D uniform image types are supported.");
444 return "???";
445 }
446 return "shader";
447 }
448
449 std::string CompilerSkSL::builtin_to_glsl(BuiltIn builtin,
450 StorageClass storage) {
451 std::string gl_builtin = CompilerGLSL::builtin_to_glsl(builtin, storage);
452 switch (builtin) {
453 case BuiltInFragCoord:
454 return "flutter_FragCoord";
455 default:
456 FLUTTER_CROSS_THROW("Builtin '" + gl_builtin + "' is not supported.");
457 break;
458 }
459
460 return "???";
461 }
462
463 std::string CompilerSkSL::to_texture_op(
464 const Instruction& i,
465 bool sparse,
466 bool* forward,
467 SmallVector<uint32_t>& inherited_expressions) {
468 auto op = static_cast<Op>(i.op);
469 if (op != OpImageSampleImplicitLod) {
470 FLUTTER_CROSS_THROW("Only simple shader sampling is supported.");
471 return "???";
472 }
473 return CompilerGLSL::to_texture_op(i, sparse, forward, inherited_expressions);
474 }
475
476 std::string CompilerSkSL::to_function_name(
477 const CompilerGLSL::TextureFunctionNameArguments& args) {
478 std::string name = to_expression(args.base.img);
479 return name + ".eval";
480 }
481
482 std::string CompilerSkSL::to_function_args(const TextureFunctionArguments& args,
483 bool* p_forward) {
484 std::string name = to_expression(args.base.img);
485
486 std::string glsl_args = CompilerGLSL::to_function_args(args, p_forward);
487 // SkSL only supports coordinates. All other arguments to texture are
488 // unsupported and will generate invalid SkSL.
489 if (args.grad_x || args.grad_y || args.lod || args.offset || args.sample ||
490 args.min_lod || args.sparse_texel || args.bias || args.component) {
491 FLUTTER_CROSS_THROW(
492 "Only sampler and position arguments are supported in texture() "
493 "calls.");
494 }
495
496 // GLSL puts the shader as the first argument, but in SkSL the shader is
497 // implicitly passed as the reciever of the 'eval' method. Therefore, the
498 // shader is removed from the GLSL argument list.
499 std::string no_shader;
500 auto npos = glsl_args.find(", "); // The first ','.
501 if (npos != std::string::npos) {
502 no_shader = glsl_args.substr(npos + 1); // The string after the first ','.
503 }
504
505 if (no_shader.empty()) {
506 FLUTTER_CROSS_THROW("Unexpected shader sampling arguments: '(" + glsl_args +
507 ")'");
508 return "()";
509 }
510
511 return name + "_size * (" + no_shader + ")";
512 }
513
514 } // namespace compiler
515 } // namespace impeller
type
GLenum type
Definition: blit_command_gles.cc:153
impeller::compiler::report_and_exit
void report_and_exit(const std::string &msg)
Definition: spirv_sksl.cc:16
impeller
Definition: allocation.cc:12
impeller::SortUniforms
std::vector< spirv_cross::ID > SortUniforms(const spirv_cross::ParsedIR *ir, const spirv_cross::Compiler *compiler, std::optional< spirv_cross::SPIRType::BaseType > type_filter, bool include)
Sorts uniform declarations in an IR according to decoration order.
Definition: uniform_sorter.cc:11
c
Point c
Definition: round_superellipse_param.cc:315
FLUTTER_CROSS_THROW
#define FLUTTER_CROSS_THROW(x)
Definition: spirv_sksl.cc:22
spirv_sksl.h
uniform_sorter.h