spirv_cpp.cpp

   1 /*
   2  * Copyright 2015-2021 Arm Limited
   3  * SPDX-License-Identifier: Apache-2.0 OR MIT
   4  *
   5  * Licensed under the Apache License, Version 2.0 (the "License");
   6  * you may not use this file except in compliance with the License.
   7  * You may obtain a copy of the License at
   8  *
   9  *     http://www.apache.org/licenses/LICENSE-2.0
  10  *
  11  * Unless required by applicable law or agreed to in writing, software
  12  * distributed under the License is distributed on an "AS IS" BASIS,
  13  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  14  * See the License for the specific language governing permissions and
  15  * limitations under the License.
  16  */
  17
  18 /*
  19  * At your option, you may choose to accept this material under either:
  20  *  1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or
  21  *  2. The MIT License, found at <http://opensource.org/licenses/MIT>.
  22  */
  23
  24 #include "spirv_cpp.hpp"
  25
  26 using namespace spv;
  27 using namespace SPIRV_CROSS_NAMESPACE;
  28 using namespace std;
  29
  30 void CompilerCPP::emit_buffer_block(const SPIRVariable &var)
  31 {
  32         add_resource_name(var.self);
  33
  34         auto &type = get<SPIRType>(var.basetype);
  35         auto instance_name = to_name(var.self);
  36
  37         uint32_t descriptor_set = ir.meta[var.self].decoration.set;
  38         uint32_t binding = ir.meta[var.self].decoration.binding;
  39
  40         emit_block_struct(type);
  41         auto buffer_name = to_name(type.self);
  42
  43         statement("internal::Resource<", buffer_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
  44         statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
  45         resource_registrations.push_back(
  46             join("s.register_resource(", instance_name, "__", ", ", descriptor_set, ", ", binding, ");"));
  47         statement("");
  48 }
  49
  50 void CompilerCPP::emit_interface_block(const SPIRVariable &var)
  51 {
  52         add_resource_name(var.self);
  53
  54         auto &type = get<SPIRType>(var.basetype);
  55
  56         const char *qual = var.storage == StorageClassInput ? "StageInput" : "StageOutput";
  57         const char *lowerqual = var.storage == StorageClassInput ? "stage_input" : "stage_output";
  58         auto instance_name = to_name(var.self);
  59         uint32_t location = ir.meta[var.self].decoration.location;
  60
  61         string buffer_name;
  62         auto flags = ir.meta[type.self].decoration.decoration_flags;
  63         if (flags.get(DecorationBlock))
  64         {
  65                 emit_block_struct(type);
  66                 buffer_name = to_name(type.self);
  67         }
  68         else
  69                 buffer_name = type_to_glsl(type);
  70
  71         statement("internal::", qual, "<", buffer_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
  72         statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
  73         resource_registrations.push_back(join("s.register_", lowerqual, "(", instance_name, "__", ", ", location, ");"));
  74         statement("");
  75 }
  76
  77 void CompilerCPP::emit_shared(const SPIRVariable &var)
  78 {
  79         add_resource_name(var.self);
  80
  81         auto instance_name = to_name(var.self);
  82         statement(CompilerGLSL::variable_decl(var), ";");
  83         statement_no_indent("#define ", instance_name, " __res->", instance_name);
  84 }
  85
  86 void CompilerCPP::emit_uniform(const SPIRVariable &var)
  87 {
  88         add_resource_name(var.self);
  89
  90         auto &type = get<SPIRType>(var.basetype);
  91         auto instance_name = to_name(var.self);
  92
  93         uint32_t descriptor_set = ir.meta[var.self].decoration.set;
  94         uint32_t binding = ir.meta[var.self].decoration.binding;
  95         uint32_t location = ir.meta[var.self].decoration.location;
  96
  97         string type_name = type_to_glsl(type);
  98         remap_variable_type_name(type, instance_name, type_name);
  99
 100         if (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage ||
 101             type.basetype == SPIRType::AtomicCounter)
 102         {
 103                 statement("internal::Resource<", type_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
 104                 statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
 105                 resource_registrations.push_back(
 106                     join("s.register_resource(", instance_name, "__", ", ", descriptor_set, ", ", binding, ");"));
 107         }
 108         else
 109         {
 110                 statement("internal::UniformConstant<", type_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
 111                 statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
 112                 resource_registrations.push_back(
 113                     join("s.register_uniform_constant(", instance_name, "__", ", ", location, ");"));
 114         }
 115
 116         statement("");
 117 }
 118
 119 void CompilerCPP::emit_push_constant_block(const SPIRVariable &var)
 120 {
 121         add_resource_name(var.self);
 122
 123         auto &type = get<SPIRType>(var.basetype);
 124         auto &flags = ir.meta[var.self].decoration.decoration_flags;
 125         if (flags.get(DecorationBinding) || flags.get(DecorationDescriptorSet))
 126                 SPIRV_CROSS_THROW("Push constant blocks cannot be compiled to GLSL with Binding or Set syntax. "
 127                                   "Remap to location with reflection API first or disable these decorations.");
 128
 129         emit_block_struct(type);
 130         auto buffer_name = to_name(type.self);
 131         auto instance_name = to_name(var.self);
 132
 133         statement("internal::PushConstant<", buffer_name, type_to_array_glsl(type, var.self), "> ", instance_name, ";");
 134         statement_no_indent("#define ", instance_name, " __res->", instance_name, ".get()");
 135         resource_registrations.push_back(join("s.register_push_constant(", instance_name, "__", ");"));
 136         statement("");
 137 }
 138
 139 void CompilerCPP::emit_block_struct(SPIRType &type)
 140 {
 141         // C++ can't do interface blocks, so we fake it by emitting a separate struct.
 142         // However, these structs are not allowed to alias anything, so remove it before
 143         // emitting the struct.
 144         //
 145         // The type we have here needs to be resolved to the non-pointer type so we can remove aliases.
 146         auto &self = get<SPIRType>(type.self);
 147         self.type_alias = 0;
 148         emit_struct(self);
 149 }
 150
 151 void CompilerCPP::emit_resources()
 152 {
 153         for (auto &id : ir.ids)
 154         {
 155                 if (id.get_type() == TypeConstant)
 156                 {
 157                         auto &c = id.get<SPIRConstant>();
 158
 159                         bool needs_declaration = c.specialization || c.is_used_as_lut;
 160
 161                         if (needs_declaration)
 162                         {
 163                                 if (!options.vulkan_semantics && c.specialization)
 164                                 {
 165                                         c.specialization_constant_macro_name =
 166                                             constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
 167                                 }
 168                                 emit_constant(c);
 169                         }
 170                 }
 171                 else if (id.get_type() == TypeConstantOp)
 172                 {
 173                         emit_specialization_constant_op(id.get<SPIRConstantOp>());
 174                 }
 175         }
 176
 177         // Output all basic struct types which are not Block or BufferBlock as these are declared inplace
 178         // when such variables are instantiated.
 179         for (auto &id : ir.ids)
 180         {
 181                 if (id.get_type() == TypeType)
 182                 {
 183                         auto &type = id.get<SPIRType>();
 184                         if (type.basetype == SPIRType::Struct && type.array.empty() && !type.pointer &&
 185                             (!ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) &&
 186                              !ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock)))
 187                         {
 188                                 emit_struct(type);
 189                         }
 190                 }
 191         }
 192
 193         statement("struct Resources : ", resource_type);
 194         begin_scope();
 195
 196         // Output UBOs and SSBOs
 197         for (auto &id : ir.ids)
 198         {
 199                 if (id.get_type() == TypeVariable)
 200                 {
 201                         auto &var = id.get<SPIRVariable>();
 202                         auto &type = get<SPIRType>(var.basetype);
 203
 204                         if (var.storage != StorageClassFunction && type.pointer && type.storage == StorageClassUniform &&
 205                             !is_hidden_variable(var) &&
 206                             (ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) ||
 207                              ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock)))
 208                         {
 209                                 emit_buffer_block(var);
 210                         }
 211                 }
 212         }
 213
 214         // Output push constant blocks
 215         for (auto &id : ir.ids)
 216         {
 217                 if (id.get_type() == TypeVariable)
 218                 {
 219                         auto &var = id.get<SPIRVariable>();
 220                         auto &type = get<SPIRType>(var.basetype);
 221                         if (!is_hidden_variable(var) && var.storage != StorageClassFunction && type.pointer &&
 222                             type.storage == StorageClassPushConstant)
 223                         {
 224                                 emit_push_constant_block(var);
 225                         }
 226                 }
 227         }
 228
 229         // Output in/out interfaces.
 230         for (auto &id : ir.ids)
 231         {
 232                 if (id.get_type() == TypeVariable)
 233                 {
 234                         auto &var = id.get<SPIRVariable>();
 235                         auto &type = get<SPIRType>(var.basetype);
 236
 237                         if (var.storage != StorageClassFunction && !is_hidden_variable(var) && type.pointer &&
 238                             (var.storage == StorageClassInput || var.storage == StorageClassOutput) &&
 239                             interface_variable_exists_in_entry_point(var.self))
 240                         {
 241                                 emit_interface_block(var);
 242                         }
 243                 }
 244         }
 245
 246         // Output Uniform Constants (values, samplers, images, etc).
 247         for (auto &id : ir.ids)
 248         {
 249                 if (id.get_type() == TypeVariable)
 250                 {
 251                         auto &var = id.get<SPIRVariable>();
 252                         auto &type = get<SPIRType>(var.basetype);
 253
 254                         if (var.storage != StorageClassFunction && !is_hidden_variable(var) && type.pointer &&
 255                             (type.storage == StorageClassUniformConstant || type.storage == StorageClassAtomicCounter))
 256                         {
 257                                 emit_uniform(var);
 258                         }
 259                 }
 260         }
 261
 262         // Global variables.
 263         bool emitted = false;
 264         for (auto global : global_variables)
 265         {
 266                 auto &var = get<SPIRVariable>(global);
 267                 if (var.storage == StorageClassWorkgroup)
 268                 {
 269                         emit_shared(var);
 270                         emitted = true;
 271                 }
 272         }
 273
 274         if (emitted)
 275                 statement("");
 276
 277         statement("inline void init(spirv_cross_shader& s)");
 278         begin_scope();
 279         statement(resource_type, "::init(s);");
 280         for (auto &reg : resource_registrations)
 281                 statement(reg);
 282         end_scope();
 283         resource_registrations.clear();
 284
 285         end_scope_decl();
 286
 287         statement("");
 288         statement("Resources* __res;");
 289         if (get_entry_point().model == ExecutionModelGLCompute)
 290                 statement("ComputePrivateResources __priv_res;");
 291         statement("");
 292
 293         // Emit regular globals which are allocated per invocation.
 294         emitted = false;
 295         for (auto global : global_variables)
 296         {
 297                 auto &var = get<SPIRVariable>(global);
 298                 if (var.storage == StorageClassPrivate)
 299                 {
 300                         if (var.storage == StorageClassWorkgroup)
 301                                 emit_shared(var);
 302                         else
 303                                 statement(CompilerGLSL::variable_decl(var), ";");
 304                         emitted = true;
 305                 }
 306         }
 307
 308         if (emitted)
 309                 statement("");
 310 }
 311
 312 string CompilerCPP::compile()
 313 {
 314         ir.fixup_reserved_names();
 315
 316         // Do not deal with ES-isms like precision, older extensions and such.
 317         options.es = false;
 318         options.version = 450;
 319         backend.float_literal_suffix = true;
 320         backend.double_literal_suffix = false;
 321         backend.long_long_literal_suffix = true;
 322         backend.uint32_t_literal_suffix = true;
 323         backend.basic_int_type = "int32_t";
 324         backend.basic_uint_type = "uint32_t";
 325         backend.swizzle_is_function = true;
 326         backend.shared_is_implied = true;
 327         backend.unsized_array_supported = false;
 328         backend.explicit_struct_type = true;
 329         backend.use_initializer_list = true;
 330
 331         fixup_type_alias();
 332         reorder_type_alias();
 333         build_function_control_flow_graphs_and_analyze();
 334         update_active_builtins();
 335
 336         uint32_t pass_count = 0;
 337         do
 338         {
 339                 resource_registrations.clear();
 340                 reset(pass_count);
 341
 342                 // Move constructor for this type is broken on GCC 4.9 ...
 343                 buffer.reset();
 344
 345                 emit_header();
 346                 emit_resources();
 347
 348                 emit_function(get<SPIRFunction>(ir.default_entry_point), Bitset());
 349
 350                 pass_count++;
 351         } while (is_forcing_recompilation());
 352
 353         // Match opening scope of emit_header().
 354         end_scope_decl();
 355         // namespace
 356         end_scope();
 357
 358         // Emit C entry points
 359         emit_c_linkage();
 360
 361         // Entry point in CPP is always main() for the time being.
 362         get_entry_point().name = "main";
 363
 364         return buffer.str();
 365 }
 366
 367 void CompilerCPP::emit_c_linkage()
 368 {
 369         statement("");
 370
 371         statement("spirv_cross_shader_t *spirv_cross_construct(void)");
 372         begin_scope();
 373         statement("return new ", impl_type, "();");
 374         end_scope();
 375
 376         statement("");
 377         statement("void spirv_cross_destruct(spirv_cross_shader_t *shader)");
 378         begin_scope();
 379         statement("delete static_cast<", impl_type, "*>(shader);");
 380         end_scope();
 381
 382         statement("");
 383         statement("void spirv_cross_invoke(spirv_cross_shader_t *shader)");
 384         begin_scope();
 385         statement("static_cast<", impl_type, "*>(shader)->invoke();");
 386         end_scope();
 387
 388         statement("");
 389         statement("static const struct spirv_cross_interface vtable =");
 390         begin_scope();
 391         statement("spirv_cross_construct,");
 392         statement("spirv_cross_destruct,");
 393         statement("spirv_cross_invoke,");
 394         end_scope_decl();
 395
 396         statement("");
 397         statement("const struct spirv_cross_interface *",
 398                   interface_name.empty() ? string("spirv_cross_get_interface") : interface_name, "(void)");
 399         begin_scope();
 400         statement("return &vtable;");
 401         end_scope();
 402 }
 403
 404 void CompilerCPP::emit_function_prototype(SPIRFunction &func, const Bitset &)
 405 {
 406         if (func.self != ir.default_entry_point)
 407                 add_function_overload(func);
 408
 409         local_variable_names = resource_names;
 410         string decl;
 411
 412         auto &type = get<SPIRType>(func.return_type);
 413         decl += "inline ";
 414         decl += type_to_glsl(type);
 415         decl += " ";
 416
 417         if (func.self == ir.default_entry_point)
 418         {
 419                 decl += "main";
 420                 processing_entry_point = true;
 421         }
 422         else
 423                 decl += to_name(func.self);
 424
 425         decl += "(";
 426         for (auto &arg : func.arguments)
 427         {
 428                 add_local_variable_name(arg.id);
 429
 430                 decl += argument_decl(arg);
 431                 if (&arg != &func.arguments.back())
 432                         decl += ", ";
 433
 434                 // Hold a pointer to the parameter so we can invalidate the readonly field if needed.
 435                 auto *var = maybe_get<SPIRVariable>(arg.id);
 436                 if (var)
 437                         var->parameter = &arg;
 438         }
 439
 440         decl += ")";
 441         statement(decl);
 442 }
 443
 444 string CompilerCPP::argument_decl(const SPIRFunction::Parameter &arg)
 445 {
 446         auto &type = expression_type(arg.id);
 447         bool constref = !type.pointer || arg.write_count == 0;
 448
 449         auto &var = get<SPIRVariable>(arg.id);
 450
 451         string base = type_to_glsl(type);
 452         string variable_name = to_name(var.self);
 453         remap_variable_type_name(type, variable_name, base);
 454
 455         for (uint32_t i = 0; i < type.array.size(); i++)
 456                 base = join("std::array<", base, ", ", to_array_size(type, i), ">");
 457
 458         return join(constref ? "const " : "", base, " &", variable_name);
 459 }
 460
 461 string CompilerCPP::variable_decl(const SPIRType &type, const string &name, uint32_t /* id */)
 462 {
 463         string base = type_to_glsl(type);
 464         remap_variable_type_name(type, name, base);
 465         bool runtime = false;
 466
 467         for (uint32_t i = 0; i < type.array.size(); i++)
 468         {
 469                 auto &array = type.array[i];
 470                 if (!array && type.array_size_literal[i])
 471                 {
 472                         // Avoid using runtime arrays with std::array since this is undefined.
 473                         // Runtime arrays cannot be passed around as values, so this is fine.
 474                         runtime = true;
 475                 }
 476                 else
 477                         base = join("std::array<", base, ", ", to_array_size(type, i), ">");
 478         }
 479         base += ' ';
 480         return base + name + (runtime ? "[1]" : "");
 481 }
 482
 483 void CompilerCPP::emit_header()
 484 {
 485         auto &execution = get_entry_point();
 486
 487         statement("// This C++ shader is autogenerated by spirv-cross.");
 488         statement("#include \"spirv_cross/internal_interface.hpp\"");
 489         statement("#include \"spirv_cross/external_interface.h\"");
 490         // Needed to properly implement GLSL-style arrays.
 491         statement("#include <array>");
 492         statement("#include <stdint.h>");
 493         statement("");
 494         statement("using namespace spirv_cross;");
 495         statement("using namespace glm;");
 496         statement("");
 497
 498         statement("namespace Impl");
 499         begin_scope();
 500
 501         switch (execution.model)
 502         {
 503         case ExecutionModelGeometry:
 504         case ExecutionModelTessellationControl:
 505         case ExecutionModelTessellationEvaluation:
 506         case ExecutionModelGLCompute:
 507         case ExecutionModelFragment:
 508         case ExecutionModelVertex:
 509                 statement("struct Shader");
 510                 begin_scope();
 511                 break;
 512
 513         default:
 514                 SPIRV_CROSS_THROW("Unsupported execution model.");
 515         }
 516
 517         switch (execution.model)
 518         {
 519         case ExecutionModelGeometry:
 520                 impl_type = "GeometryShader<Impl::Shader, Impl::Shader::Resources>";
 521                 resource_type = "GeometryResources";
 522                 break;
 523
 524         case ExecutionModelVertex:
 525                 impl_type = "VertexShader<Impl::Shader, Impl::Shader::Resources>";
 526                 resource_type = "VertexResources";
 527                 break;
 528
 529         case ExecutionModelFragment:
 530                 impl_type = "FragmentShader<Impl::Shader, Impl::Shader::Resources>";
 531                 resource_type = "FragmentResources";
 532                 break;
 533
 534         case ExecutionModelGLCompute:
 535                 impl_type = join("ComputeShader<Impl::Shader, Impl::Shader::Resources, ", execution.workgroup_size.x, ", ",
 536                                  execution.workgroup_size.y, ", ", execution.workgroup_size.z, ">");
 537                 resource_type = "ComputeResources";
 538                 break;
 539
 540         case ExecutionModelTessellationControl:
 541                 impl_type = "TessControlShader<Impl::Shader, Impl::Shader::Resources>";
 542                 resource_type = "TessControlResources";
 543                 break;
 544
 545         case ExecutionModelTessellationEvaluation:
 546                 impl_type = "TessEvaluationShader<Impl::Shader, Impl::Shader::Resources>";
 547                 resource_type = "TessEvaluationResources";
 548                 break;
 549
 550         default:
 551                 SPIRV_CROSS_THROW("Unsupported execution model.");
 552         }
 553 }