1 /* 2 * Copyright 2016-2021 The Brenwill Workshop Ltd. 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 #ifndef SPIRV_CROSS_MSL_HPP 25 #define SPIRV_CROSS_MSL_HPP 26 27 #include "spirv_glsl.hpp" 28 #include <map> 29 #include <set> 30 #include <stddef.h> 31 #include <unordered_map> 32 #include <unordered_set> 33 34 namespace SPIRV_CROSS_NAMESPACE 35 { 36 37 // Indicates the format of a shader input. Currently limited to specifying 38 // if the input is an 8-bit unsigned integer, 16-bit unsigned integer, or 39 // some other format. 40 enum MSLShaderInputFormat 41 { 42 MSL_SHADER_INPUT_FORMAT_OTHER = 0, 43 MSL_SHADER_INPUT_FORMAT_UINT8 = 1, 44 MSL_SHADER_INPUT_FORMAT_UINT16 = 2, 45 MSL_SHADER_INPUT_FORMAT_ANY16 = 3, 46 MSL_SHADER_INPUT_FORMAT_ANY32 = 4, 47 48 // Deprecated aliases. 49 MSL_VERTEX_FORMAT_OTHER = MSL_SHADER_INPUT_FORMAT_OTHER, 50 MSL_VERTEX_FORMAT_UINT8 = MSL_SHADER_INPUT_FORMAT_UINT8, 51 MSL_VERTEX_FORMAT_UINT16 = MSL_SHADER_INPUT_FORMAT_UINT16, 52 53 MSL_SHADER_INPUT_FORMAT_INT_MAX = 0x7fffffff 54 }; 55 56 // Defines MSL characteristics of an input variable at a particular location. 57 // After compilation, it is possible to query whether or not this location was used. 58 // If vecsize is nonzero, it must be greater than or equal to the vecsize declared in the shader, 59 // or behavior is undefined. 60 struct MSLShaderInput 61 { 62 uint32_t location = 0; 63 MSLShaderInputFormat format = MSL_SHADER_INPUT_FORMAT_OTHER; 64 spv::BuiltIn builtin = spv::BuiltInMax; 65 uint32_t vecsize = 0; 66 }; 67 68 // Matches the binding index of a MSL resource for a binding within a descriptor set. 69 // Taken together, the stage, desc_set and binding combine to form a reference to a resource 70 // descriptor used in a particular shading stage. The count field indicates the number of 71 // resources consumed by this binding, if the binding represents an array of resources. 72 // If the resource array is a run-time-sized array, which are legal in GLSL or SPIR-V, this value 73 // will be used to declare the array size in MSL, which does not support run-time-sized arrays. 74 // If pad_argument_buffer_resources is enabled, the base_type and count values are used to 75 // specify the base type and array size of the resource in the argument buffer, if that resource 76 // is not defined and used by the shader. With pad_argument_buffer_resources enabled, this 77 // information will be used to pad the argument buffer structure, in order to align that 78 // structure consistently for all uses, across all shaders, of the descriptor set represented 79 // by the arugment buffer. If pad_argument_buffer_resources is disabled, base_type does not 80 // need to be populated, and if the resource is also not a run-time sized array, the count 81 // field does not need to be populated. 82 // If using MSL 2.0 argument buffers, the descriptor set is not marked as a discrete descriptor set, 83 // and (for iOS only) the resource is not a storage image (sampled != 2), the binding reference we 84 // remap to will become an [[id(N)]] attribute within the "descriptor set" argument buffer structure. 85 // For resources which are bound in the "classic" MSL 1.0 way or discrete descriptors, the remap will 86 // become a [[buffer(N)]], [[texture(N)]] or [[sampler(N)]] depending on the resource types used. 87 struct MSLResourceBinding 88 { 89 spv::ExecutionModel stage = spv::ExecutionModelMax; 90 SPIRType::BaseType basetype = SPIRType::Unknown; 91 uint32_t desc_set = 0; 92 uint32_t binding = 0; 93 uint32_t count = 0; 94 uint32_t msl_buffer = 0; 95 uint32_t msl_texture = 0; 96 uint32_t msl_sampler = 0; 97 }; 98 99 enum MSLSamplerCoord 100 { 101 MSL_SAMPLER_COORD_NORMALIZED = 0, 102 MSL_SAMPLER_COORD_PIXEL = 1, 103 MSL_SAMPLER_INT_MAX = 0x7fffffff 104 }; 105 106 enum MSLSamplerFilter 107 { 108 MSL_SAMPLER_FILTER_NEAREST = 0, 109 MSL_SAMPLER_FILTER_LINEAR = 1, 110 MSL_SAMPLER_FILTER_INT_MAX = 0x7fffffff 111 }; 112 113 enum MSLSamplerMipFilter 114 { 115 MSL_SAMPLER_MIP_FILTER_NONE = 0, 116 MSL_SAMPLER_MIP_FILTER_NEAREST = 1, 117 MSL_SAMPLER_MIP_FILTER_LINEAR = 2, 118 MSL_SAMPLER_MIP_FILTER_INT_MAX = 0x7fffffff 119 }; 120 121 enum MSLSamplerAddress 122 { 123 MSL_SAMPLER_ADDRESS_CLAMP_TO_ZERO = 0, 124 MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE = 1, 125 MSL_SAMPLER_ADDRESS_CLAMP_TO_BORDER = 2, 126 MSL_SAMPLER_ADDRESS_REPEAT = 3, 127 MSL_SAMPLER_ADDRESS_MIRRORED_REPEAT = 4, 128 MSL_SAMPLER_ADDRESS_INT_MAX = 0x7fffffff 129 }; 130 131 enum MSLSamplerCompareFunc 132 { 133 MSL_SAMPLER_COMPARE_FUNC_NEVER = 0, 134 MSL_SAMPLER_COMPARE_FUNC_LESS = 1, 135 MSL_SAMPLER_COMPARE_FUNC_LESS_EQUAL = 2, 136 MSL_SAMPLER_COMPARE_FUNC_GREATER = 3, 137 MSL_SAMPLER_COMPARE_FUNC_GREATER_EQUAL = 4, 138 MSL_SAMPLER_COMPARE_FUNC_EQUAL = 5, 139 MSL_SAMPLER_COMPARE_FUNC_NOT_EQUAL = 6, 140 MSL_SAMPLER_COMPARE_FUNC_ALWAYS = 7, 141 MSL_SAMPLER_COMPARE_FUNC_INT_MAX = 0x7fffffff 142 }; 143 144 enum MSLSamplerBorderColor 145 { 146 MSL_SAMPLER_BORDER_COLOR_TRANSPARENT_BLACK = 0, 147 MSL_SAMPLER_BORDER_COLOR_OPAQUE_BLACK = 1, 148 MSL_SAMPLER_BORDER_COLOR_OPAQUE_WHITE = 2, 149 MSL_SAMPLER_BORDER_COLOR_INT_MAX = 0x7fffffff 150 }; 151 152 enum MSLFormatResolution 153 { 154 MSL_FORMAT_RESOLUTION_444 = 0, 155 MSL_FORMAT_RESOLUTION_422, 156 MSL_FORMAT_RESOLUTION_420, 157 MSL_FORMAT_RESOLUTION_INT_MAX = 0x7fffffff 158 }; 159 160 enum MSLChromaLocation 161 { 162 MSL_CHROMA_LOCATION_COSITED_EVEN = 0, 163 MSL_CHROMA_LOCATION_MIDPOINT, 164 MSL_CHROMA_LOCATION_INT_MAX = 0x7fffffff 165 }; 166 167 enum MSLComponentSwizzle 168 { 169 MSL_COMPONENT_SWIZZLE_IDENTITY = 0, 170 MSL_COMPONENT_SWIZZLE_ZERO, 171 MSL_COMPONENT_SWIZZLE_ONE, 172 MSL_COMPONENT_SWIZZLE_R, 173 MSL_COMPONENT_SWIZZLE_G, 174 MSL_COMPONENT_SWIZZLE_B, 175 MSL_COMPONENT_SWIZZLE_A, 176 MSL_COMPONENT_SWIZZLE_INT_MAX = 0x7fffffff 177 }; 178 179 enum MSLSamplerYCbCrModelConversion 180 { 181 MSL_SAMPLER_YCBCR_MODEL_CONVERSION_RGB_IDENTITY = 0, 182 MSL_SAMPLER_YCBCR_MODEL_CONVERSION_YCBCR_IDENTITY, 183 MSL_SAMPLER_YCBCR_MODEL_CONVERSION_YCBCR_BT_709, 184 MSL_SAMPLER_YCBCR_MODEL_CONVERSION_YCBCR_BT_601, 185 MSL_SAMPLER_YCBCR_MODEL_CONVERSION_YCBCR_BT_2020, 186 MSL_SAMPLER_YCBCR_MODEL_CONVERSION_INT_MAX = 0x7fffffff 187 }; 188 189 enum MSLSamplerYCbCrRange 190 { 191 MSL_SAMPLER_YCBCR_RANGE_ITU_FULL = 0, 192 MSL_SAMPLER_YCBCR_RANGE_ITU_NARROW, 193 MSL_SAMPLER_YCBCR_RANGE_INT_MAX = 0x7fffffff 194 }; 195 196 struct MSLConstexprSampler 197 { 198 MSLSamplerCoord coord = MSL_SAMPLER_COORD_NORMALIZED; 199 MSLSamplerFilter min_filter = MSL_SAMPLER_FILTER_NEAREST; 200 MSLSamplerFilter mag_filter = MSL_SAMPLER_FILTER_NEAREST; 201 MSLSamplerMipFilter mip_filter = MSL_SAMPLER_MIP_FILTER_NONE; 202 MSLSamplerAddress s_address = MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE; 203 MSLSamplerAddress t_address = MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE; 204 MSLSamplerAddress r_address = MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE; 205 MSLSamplerCompareFunc compare_func = MSL_SAMPLER_COMPARE_FUNC_NEVER; 206 MSLSamplerBorderColor border_color = MSL_SAMPLER_BORDER_COLOR_TRANSPARENT_BLACK; 207 float lod_clamp_min = 0.0f; 208 float lod_clamp_max = 1000.0f; 209 int max_anisotropy = 1; 210 211 // Sampler Y'CbCr conversion parameters 212 uint32_t planes = 0; 213 MSLFormatResolution resolution = MSL_FORMAT_RESOLUTION_444; 214 MSLSamplerFilter chroma_filter = MSL_SAMPLER_FILTER_NEAREST; 215 MSLChromaLocation x_chroma_offset = MSL_CHROMA_LOCATION_COSITED_EVEN; 216 MSLChromaLocation y_chroma_offset = MSL_CHROMA_LOCATION_COSITED_EVEN; 217 MSLComponentSwizzle swizzle[4]; // IDENTITY, IDENTITY, IDENTITY, IDENTITY 218 MSLSamplerYCbCrModelConversion ycbcr_model = MSL_SAMPLER_YCBCR_MODEL_CONVERSION_RGB_IDENTITY; 219 MSLSamplerYCbCrRange ycbcr_range = MSL_SAMPLER_YCBCR_RANGE_ITU_FULL; 220 uint32_t bpc = 8; 221 222 bool compare_enable = false; 223 bool lod_clamp_enable = false; 224 bool anisotropy_enable = false; 225 bool ycbcr_conversion_enable = false; 226 MSLConstexprSamplerSPIRV_CROSS_NAMESPACE::MSLConstexprSampler227 MSLConstexprSampler() 228 { 229 for (uint32_t i = 0; i < 4; i++) 230 swizzle[i] = MSL_COMPONENT_SWIZZLE_IDENTITY; 231 } swizzle_is_identitySPIRV_CROSS_NAMESPACE::MSLConstexprSampler232 bool swizzle_is_identity() const 233 { 234 return (swizzle[0] == MSL_COMPONENT_SWIZZLE_IDENTITY && swizzle[1] == MSL_COMPONENT_SWIZZLE_IDENTITY && 235 swizzle[2] == MSL_COMPONENT_SWIZZLE_IDENTITY && swizzle[3] == MSL_COMPONENT_SWIZZLE_IDENTITY); 236 } swizzle_has_one_or_zeroSPIRV_CROSS_NAMESPACE::MSLConstexprSampler237 bool swizzle_has_one_or_zero() const 238 { 239 return (swizzle[0] == MSL_COMPONENT_SWIZZLE_ZERO || swizzle[0] == MSL_COMPONENT_SWIZZLE_ONE || 240 swizzle[1] == MSL_COMPONENT_SWIZZLE_ZERO || swizzle[1] == MSL_COMPONENT_SWIZZLE_ONE || 241 swizzle[2] == MSL_COMPONENT_SWIZZLE_ZERO || swizzle[2] == MSL_COMPONENT_SWIZZLE_ONE || 242 swizzle[3] == MSL_COMPONENT_SWIZZLE_ZERO || swizzle[3] == MSL_COMPONENT_SWIZZLE_ONE); 243 } 244 }; 245 246 // Special constant used in a MSLResourceBinding desc_set 247 // element to indicate the bindings for the push constants. 248 // Kinda deprecated. Just use ResourceBindingPushConstant{DescriptorSet,Binding} directly. 249 static const uint32_t kPushConstDescSet = ResourceBindingPushConstantDescriptorSet; 250 251 // Special constant used in a MSLResourceBinding binding 252 // element to indicate the bindings for the push constants. 253 // Kinda deprecated. Just use ResourceBindingPushConstant{DescriptorSet,Binding} directly. 254 static const uint32_t kPushConstBinding = ResourceBindingPushConstantBinding; 255 256 // Special constant used in a MSLResourceBinding binding 257 // element to indicate the buffer binding for swizzle buffers. 258 static const uint32_t kSwizzleBufferBinding = ~(1u); 259 260 // Special constant used in a MSLResourceBinding binding 261 // element to indicate the buffer binding for buffer size buffers to support OpArrayLength. 262 static const uint32_t kBufferSizeBufferBinding = ~(2u); 263 264 // Special constant used in a MSLResourceBinding binding 265 // element to indicate the buffer binding used for the argument buffer itself. 266 // This buffer binding should be kept as small as possible as all automatic bindings for buffers 267 // will start at max(kArgumentBufferBinding) + 1. 268 static const uint32_t kArgumentBufferBinding = ~(3u); 269 270 static const uint32_t kMaxArgumentBuffers = 8; 271 272 // The arbitrary maximum for the nesting of array of array copies. 273 static const uint32_t kArrayCopyMultidimMax = 6; 274 275 // Decompiles SPIR-V to Metal Shading Language 276 class CompilerMSL : public CompilerGLSL 277 { 278 public: 279 // Options for compiling to Metal Shading Language 280 struct Options 281 { 282 typedef enum 283 { 284 iOS = 0, 285 macOS = 1 286 } Platform; 287 288 Platform platform = macOS; 289 uint32_t msl_version = make_msl_version(1, 2); 290 uint32_t texel_buffer_texture_width = 4096; // Width of 2D Metal textures used as 1D texel buffers 291 uint32_t r32ui_linear_texture_alignment = 4; 292 uint32_t r32ui_alignment_constant_id = 65535; 293 uint32_t swizzle_buffer_index = 30; 294 uint32_t indirect_params_buffer_index = 29; 295 uint32_t shader_output_buffer_index = 28; 296 uint32_t shader_patch_output_buffer_index = 27; 297 uint32_t shader_tess_factor_buffer_index = 26; 298 uint32_t buffer_size_buffer_index = 25; 299 uint32_t view_mask_buffer_index = 24; 300 uint32_t dynamic_offsets_buffer_index = 23; 301 uint32_t shader_input_buffer_index = 22; 302 uint32_t shader_index_buffer_index = 21; 303 uint32_t shader_input_wg_index = 0; 304 uint32_t device_index = 0; 305 uint32_t enable_frag_output_mask = 0xffffffff; 306 // Metal doesn't allow setting a fixed sample mask directly in the pipeline. 307 // We can evade this restriction by ANDing the internal sample_mask output 308 // of the shader with the additional fixed sample mask. 309 uint32_t additional_fixed_sample_mask = 0xffffffff; 310 bool enable_point_size_builtin = true; 311 bool enable_frag_depth_builtin = true; 312 bool enable_frag_stencil_ref_builtin = true; 313 bool disable_rasterization = false; 314 bool capture_output_to_buffer = false; 315 bool swizzle_texture_samples = false; 316 bool tess_domain_origin_lower_left = false; 317 bool multiview = false; 318 bool multiview_layered_rendering = true; 319 bool view_index_from_device_index = false; 320 bool dispatch_base = false; 321 bool texture_1D_as_2D = false; 322 323 // Enable use of MSL 2.0 indirect argument buffers. 324 // MSL 2.0 must also be enabled. 325 bool argument_buffers = false; 326 327 // Ensures vertex and instance indices start at zero. This reflects the behavior of HLSL with SV_VertexID and SV_InstanceID. 328 bool enable_base_index_zero = false; 329 330 // Fragment output in MSL must have at least as many components as the render pass. 331 // Add support to explicit pad out components. 332 bool pad_fragment_output_components = false; 333 334 // Specifies whether the iOS target version supports the [[base_vertex]] and [[base_instance]] attributes. 335 bool ios_support_base_vertex_instance = false; 336 337 // Use Metal's native frame-buffer fetch API for subpass inputs. 338 bool use_framebuffer_fetch_subpasses = false; 339 340 // Enables use of "fma" intrinsic for invariant float math 341 bool invariant_float_math = false; 342 343 // Emulate texturecube_array with texture2d_array for iOS where this type is not available 344 bool emulate_cube_array = false; 345 346 // Allow user to enable decoration binding 347 bool enable_decoration_binding = false; 348 349 // Requires MSL 2.1, use the native support for texel buffers. 350 bool texture_buffer_native = false; 351 352 // Forces all resources which are part of an argument buffer to be considered active. 353 // This ensures ABI compatibility between shaders where some resources might be unused, 354 // and would otherwise declare a different IAB. 355 bool force_active_argument_buffer_resources = false; 356 357 // Aligns each resource in an argument buffer to its assigned index value, id(N), 358 // by adding synthetic padding members in the argument buffer struct for any resources 359 // in the argument buffer that are not defined and used by the shader. This allows 360 // the shader to index into the correct argument in a descriptor set argument buffer 361 // that is shared across shaders, where not all resources in the argument buffer are 362 // defined in each shader. For this to work, an MSLResourceBinding must be provided for 363 // all descriptors in any descriptor set held in an argument buffer in the shader, and 364 // that MSLResourceBinding must have the basetype and count members populated correctly. 365 // The implementation here assumes any inline blocks in the argument buffer is provided 366 // in a Metal buffer, and doesn't take into consideration inline blocks that are 367 // optionally embedded directly into the argument buffer via add_inline_uniform_block(). 368 bool pad_argument_buffer_resources = false; 369 370 // Forces the use of plain arrays, which works around certain driver bugs on certain versions 371 // of Intel Macbooks. See https://github.com/KhronosGroup/SPIRV-Cross/issues/1210. 372 // May reduce performance in scenarios where arrays are copied around as value-types. 373 bool force_native_arrays = false; 374 375 // If a shader writes clip distance, also emit user varyings which 376 // can be read in subsequent stages. 377 bool enable_clip_distance_user_varying = true; 378 379 // In a tessellation control shader, assume that more than one patch can be processed in a 380 // single workgroup. This requires changes to the way the InvocationId and PrimitiveId 381 // builtins are processed, but should result in more efficient usage of the GPU. 382 bool multi_patch_workgroup = false; 383 384 // If set, a vertex shader will be compiled as part of a tessellation pipeline. 385 // It will be translated as a compute kernel, so it can use the global invocation ID 386 // to index the output buffer. 387 bool vertex_for_tessellation = false; 388 389 // Assume that SubpassData images have multiple layers. Layered input attachments 390 // are addressed relative to the Layer output from the vertex pipeline. This option 391 // has no effect with multiview, since all input attachments are assumed to be layered 392 // and will be addressed using the current ViewIndex. 393 bool arrayed_subpass_input = false; 394 395 // Whether to use SIMD-group or quadgroup functions to implement group nnon-uniform 396 // operations. Some GPUs on iOS do not support the SIMD-group functions, only the 397 // quadgroup functions. 398 bool ios_use_simdgroup_functions = false; 399 400 // If set, the subgroup size will be assumed to be one, and subgroup-related 401 // builtins and operations will be emitted accordingly. This mode is intended to 402 // be used by MoltenVK on hardware/software configurations which do not provide 403 // sufficient support for subgroups. 404 bool emulate_subgroups = false; 405 406 // If nonzero, a fixed subgroup size to assume. Metal, similarly to VK_EXT_subgroup_size_control, 407 // allows the SIMD-group size (aka thread execution width) to vary depending on 408 // register usage and requirements. In certain circumstances--for example, a pipeline 409 // in MoltenVK without VK_PIPELINE_SHADER_STAGE_CREATE_ALLOW_VARYING_SUBGROUP_SIZE_BIT_EXT-- 410 // this is undesirable. This fixes the value of the SubgroupSize builtin, instead of 411 // mapping it to the Metal builtin [[thread_execution_width]]. If the thread 412 // execution width is reduced, the extra invocations will appear to be inactive. 413 // If zero, the SubgroupSize will be allowed to vary, and the builtin will be mapped 414 // to the Metal [[thread_execution_width]] builtin. 415 uint32_t fixed_subgroup_size = 0; 416 417 enum class IndexType 418 { 419 None = 0, 420 UInt16 = 1, 421 UInt32 = 2 422 }; 423 424 // The type of index in the index buffer, if present. For a compute shader, Metal 425 // requires specifying the indexing at pipeline creation, rather than at draw time 426 // as with graphics pipelines. This means we must create three different pipelines, 427 // for no indexing, 16-bit indices, and 32-bit indices. Each requires different 428 // handling for the gl_VertexIndex builtin. We may as well, then, create three 429 // different shaders for these three scenarios. 430 IndexType vertex_index_type = IndexType::None; 431 432 // If set, a dummy [[sample_id]] input is added to a fragment shader if none is present. 433 // This will force the shader to run at sample rate, assuming Metal does not optimize 434 // the extra threads away. 435 bool force_sample_rate_shading = false; 436 is_iosSPIRV_CROSS_NAMESPACE::CompilerMSL::Options437 bool is_ios() const 438 { 439 return platform == iOS; 440 } 441 is_macosSPIRV_CROSS_NAMESPACE::CompilerMSL::Options442 bool is_macos() const 443 { 444 return platform == macOS; 445 } 446 set_msl_versionSPIRV_CROSS_NAMESPACE::CompilerMSL::Options447 void set_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0) 448 { 449 msl_version = make_msl_version(major, minor, patch); 450 } 451 supports_msl_versionSPIRV_CROSS_NAMESPACE::CompilerMSL::Options452 bool supports_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0) const 453 { 454 return msl_version >= make_msl_version(major, minor, patch); 455 } 456 make_msl_versionSPIRV_CROSS_NAMESPACE::CompilerMSL::Options457 static uint32_t make_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0) 458 { 459 return (major * 10000) + (minor * 100) + patch; 460 } 461 }; 462 get_msl_options() const463 const Options &get_msl_options() const 464 { 465 return msl_options; 466 } 467 set_msl_options(const Options & opts)468 void set_msl_options(const Options &opts) 469 { 470 msl_options = opts; 471 } 472 473 // Provide feedback to calling API to allow runtime to disable pipeline 474 // rasterization if vertex shader requires rasterization to be disabled. get_is_rasterization_disabled() const475 bool get_is_rasterization_disabled() const 476 { 477 return is_rasterization_disabled && (get_entry_point().model == spv::ExecutionModelVertex || 478 get_entry_point().model == spv::ExecutionModelTessellationControl || 479 get_entry_point().model == spv::ExecutionModelTessellationEvaluation); 480 } 481 482 // Provide feedback to calling API to allow it to pass an auxiliary 483 // swizzle buffer if the shader needs it. needs_swizzle_buffer() const484 bool needs_swizzle_buffer() const 485 { 486 return used_swizzle_buffer; 487 } 488 489 // Provide feedback to calling API to allow it to pass a buffer 490 // containing STORAGE_BUFFER buffer sizes to support OpArrayLength. needs_buffer_size_buffer() const491 bool needs_buffer_size_buffer() const 492 { 493 return !buffers_requiring_array_length.empty(); 494 } 495 496 // Provide feedback to calling API to allow it to pass a buffer 497 // containing the view mask for the current multiview subpass. needs_view_mask_buffer() const498 bool needs_view_mask_buffer() const 499 { 500 return msl_options.multiview && !msl_options.view_index_from_device_index; 501 } 502 503 // Provide feedback to calling API to allow it to pass a buffer 504 // containing the dispatch base workgroup ID. needs_dispatch_base_buffer() const505 bool needs_dispatch_base_buffer() const 506 { 507 return msl_options.dispatch_base && !msl_options.supports_msl_version(1, 2); 508 } 509 510 // Provide feedback to calling API to allow it to pass an output 511 // buffer if the shader needs it. needs_output_buffer() const512 bool needs_output_buffer() const 513 { 514 return capture_output_to_buffer && stage_out_var_id != ID(0); 515 } 516 517 // Provide feedback to calling API to allow it to pass a patch output 518 // buffer if the shader needs it. needs_patch_output_buffer() const519 bool needs_patch_output_buffer() const 520 { 521 return capture_output_to_buffer && patch_stage_out_var_id != ID(0); 522 } 523 524 // Provide feedback to calling API to allow it to pass an input threadgroup 525 // buffer if the shader needs it. needs_input_threadgroup_mem() const526 bool needs_input_threadgroup_mem() const 527 { 528 return capture_output_to_buffer && stage_in_var_id != ID(0); 529 } 530 531 explicit CompilerMSL(std::vector<uint32_t> spirv); 532 CompilerMSL(const uint32_t *ir, size_t word_count); 533 explicit CompilerMSL(const ParsedIR &ir); 534 explicit CompilerMSL(ParsedIR &&ir); 535 536 // input is a shader input description used to fix up shader input variables. 537 // If shader inputs are provided, is_msl_shader_input_used() will return true after 538 // calling ::compile() if the location was used by the MSL code. 539 void add_msl_shader_input(const MSLShaderInput &input); 540 541 // resource is a resource binding to indicate the MSL buffer, 542 // texture or sampler index to use for a particular SPIR-V description set 543 // and binding. If resource bindings are provided, 544 // is_msl_resource_binding_used() will return true after calling ::compile() if 545 // the set/binding combination was used by the MSL code. 546 void add_msl_resource_binding(const MSLResourceBinding &resource); 547 548 // desc_set and binding are the SPIR-V descriptor set and binding of a buffer resource 549 // in this shader. index is the index within the dynamic offset buffer to use. This 550 // function marks that resource as using a dynamic offset (VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC 551 // or VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC). This function only has any effect if argument buffers 552 // are enabled. If so, the buffer will have its address adjusted at the beginning of the shader with 553 // an offset taken from the dynamic offset buffer. 554 void add_dynamic_buffer(uint32_t desc_set, uint32_t binding, uint32_t index); 555 556 // desc_set and binding are the SPIR-V descriptor set and binding of a buffer resource 557 // in this shader. This function marks that resource as an inline uniform block 558 // (VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT). This function only has any effect if argument buffers 559 // are enabled. If so, the buffer block will be directly embedded into the argument 560 // buffer, instead of being referenced indirectly via pointer. 561 void add_inline_uniform_block(uint32_t desc_set, uint32_t binding); 562 563 // When using MSL argument buffers, we can force "classic" MSL 1.0 binding schemes for certain descriptor sets. 564 // This corresponds to VK_KHR_push_descriptor in Vulkan. 565 void add_discrete_descriptor_set(uint32_t desc_set); 566 567 // If an argument buffer is large enough, it may need to be in the device storage space rather than 568 // constant. Opt-in to this behavior here on a per set basis. 569 void set_argument_buffer_device_address_space(uint32_t desc_set, bool device_storage); 570 571 // Query after compilation is done. This allows you to check if an input location was used by the shader. 572 bool is_msl_shader_input_used(uint32_t location); 573 574 // If not using add_msl_shader_input, it's possible 575 // that certain builtin attributes need to be automatically assigned locations. 576 // This is typical for tessellation builtin inputs such as tess levels, gl_Position, etc. 577 // This returns k_unknown_location if the location was explicitly assigned with 578 // add_msl_shader_input or the builtin is not used, otherwise returns N in [[attribute(N)]]. 579 uint32_t get_automatic_builtin_input_location(spv::BuiltIn builtin) const; 580 581 // NOTE: Only resources which are remapped using add_msl_resource_binding will be reported here. 582 // Constexpr samplers are always assumed to be emitted. 583 // No specific MSLResourceBinding remapping is required for constexpr samplers as long as they are remapped 584 // by remap_constexpr_sampler(_by_binding). 585 bool is_msl_resource_binding_used(spv::ExecutionModel model, uint32_t set, uint32_t binding) const; 586 587 // This must only be called after a successful call to CompilerMSL::compile(). 588 // For a variable resource ID obtained through reflection API, report the automatically assigned resource index. 589 // If the descriptor set was part of an argument buffer, report the [[id(N)]], 590 // or [[buffer/texture/sampler]] binding for other resources. 591 // If the resource was a combined image sampler, report the image binding here, 592 // use the _secondary version of this call to query the sampler half of the resource. 593 // If no binding exists, uint32_t(-1) is returned. 594 uint32_t get_automatic_msl_resource_binding(uint32_t id) const; 595 596 // Same as get_automatic_msl_resource_binding, but should only be used for combined image samplers, in which case the 597 // sampler's binding is returned instead. For any other resource type, -1 is returned. 598 // Secondary bindings are also used for the auxillary image atomic buffer. 599 uint32_t get_automatic_msl_resource_binding_secondary(uint32_t id) const; 600 601 // Same as get_automatic_msl_resource_binding, but should only be used for combined image samplers for multiplanar images, 602 // in which case the second plane's binding is returned instead. For any other resource type, -1 is returned. 603 uint32_t get_automatic_msl_resource_binding_tertiary(uint32_t id) const; 604 605 // Same as get_automatic_msl_resource_binding, but should only be used for combined image samplers for triplanar images, 606 // in which case the third plane's binding is returned instead. For any other resource type, -1 is returned. 607 uint32_t get_automatic_msl_resource_binding_quaternary(uint32_t id) const; 608 609 // Compiles the SPIR-V code into Metal Shading Language. 610 std::string compile() override; 611 612 // Remap a sampler with ID to a constexpr sampler. 613 // Older iOS targets must use constexpr samplers in certain cases (PCF), 614 // so a static sampler must be used. 615 // The sampler will not consume a binding, but be declared in the entry point as a constexpr sampler. 616 // This can be used on both combined image/samplers (sampler2D) or standalone samplers. 617 // The remapped sampler must not be an array of samplers. 618 // Prefer remap_constexpr_sampler_by_binding unless you're also doing reflection anyways. 619 void remap_constexpr_sampler(VariableID id, const MSLConstexprSampler &sampler); 620 621 // Same as remap_constexpr_sampler, except you provide set/binding, rather than variable ID. 622 // Remaps based on ID take priority over set/binding remaps. 623 void remap_constexpr_sampler_by_binding(uint32_t desc_set, uint32_t binding, const MSLConstexprSampler &sampler); 624 625 // If using CompilerMSL::Options::pad_fragment_output_components, override the number of components we expect 626 // to use for a particular location. The default is 4 if number of components is not overridden. 627 void set_fragment_output_components(uint32_t location, uint32_t components); 628 629 void set_combined_sampler_suffix(const char *suffix); 630 const char *get_combined_sampler_suffix() const; 631 632 protected: 633 // An enum of SPIR-V functions that are implemented in additional 634 // source code that is added to the shader if necessary. 635 enum SPVFuncImpl 636 { 637 SPVFuncImplNone, 638 SPVFuncImplMod, 639 SPVFuncImplRadians, 640 SPVFuncImplDegrees, 641 SPVFuncImplFindILsb, 642 SPVFuncImplFindSMsb, 643 SPVFuncImplFindUMsb, 644 SPVFuncImplSSign, 645 SPVFuncImplArrayCopyMultidimBase, 646 // Unfortunately, we cannot use recursive templates in the MSL compiler properly, 647 // so stamp out variants up to some arbitrary maximum. 648 SPVFuncImplArrayCopy = SPVFuncImplArrayCopyMultidimBase + 1, 649 SPVFuncImplArrayOfArrayCopy2Dim = SPVFuncImplArrayCopyMultidimBase + 2, 650 SPVFuncImplArrayOfArrayCopy3Dim = SPVFuncImplArrayCopyMultidimBase + 3, 651 SPVFuncImplArrayOfArrayCopy4Dim = SPVFuncImplArrayCopyMultidimBase + 4, 652 SPVFuncImplArrayOfArrayCopy5Dim = SPVFuncImplArrayCopyMultidimBase + 5, 653 SPVFuncImplArrayOfArrayCopy6Dim = SPVFuncImplArrayCopyMultidimBase + 6, 654 SPVFuncImplTexelBufferCoords, 655 SPVFuncImplImage2DAtomicCoords, // Emulate texture2D atomic operations 656 SPVFuncImplFMul, 657 SPVFuncImplFAdd, 658 SPVFuncImplFSub, 659 SPVFuncImplCubemapTo2DArrayFace, 660 SPVFuncImplUnsafeArray, // Allow Metal to use the array<T> template to make arrays a value type 661 SPVFuncImplInverse4x4, 662 SPVFuncImplInverse3x3, 663 SPVFuncImplInverse2x2, 664 // It is very important that this come before *Swizzle and ChromaReconstruct*, to ensure it's 665 // emitted before them. 666 SPVFuncImplForwardArgs, 667 // Likewise, this must come before *Swizzle. 668 SPVFuncImplGetSwizzle, 669 SPVFuncImplTextureSwizzle, 670 SPVFuncImplGatherSwizzle, 671 SPVFuncImplGatherCompareSwizzle, 672 SPVFuncImplSubgroupBroadcast, 673 SPVFuncImplSubgroupBroadcastFirst, 674 SPVFuncImplSubgroupBallot, 675 SPVFuncImplSubgroupBallotBitExtract, 676 SPVFuncImplSubgroupBallotFindLSB, 677 SPVFuncImplSubgroupBallotFindMSB, 678 SPVFuncImplSubgroupBallotBitCount, 679 SPVFuncImplSubgroupAllEqual, 680 SPVFuncImplSubgroupShuffle, 681 SPVFuncImplSubgroupShuffleXor, 682 SPVFuncImplSubgroupShuffleUp, 683 SPVFuncImplSubgroupShuffleDown, 684 SPVFuncImplQuadBroadcast, 685 SPVFuncImplQuadSwap, 686 SPVFuncImplReflectScalar, 687 SPVFuncImplRefractScalar, 688 SPVFuncImplFaceForwardScalar, 689 SPVFuncImplChromaReconstructNearest2Plane, 690 SPVFuncImplChromaReconstructNearest3Plane, 691 SPVFuncImplChromaReconstructLinear422CositedEven2Plane, 692 SPVFuncImplChromaReconstructLinear422CositedEven3Plane, 693 SPVFuncImplChromaReconstructLinear422Midpoint2Plane, 694 SPVFuncImplChromaReconstructLinear422Midpoint3Plane, 695 SPVFuncImplChromaReconstructLinear420XCositedEvenYCositedEven2Plane, 696 SPVFuncImplChromaReconstructLinear420XCositedEvenYCositedEven3Plane, 697 SPVFuncImplChromaReconstructLinear420XMidpointYCositedEven2Plane, 698 SPVFuncImplChromaReconstructLinear420XMidpointYCositedEven3Plane, 699 SPVFuncImplChromaReconstructLinear420XCositedEvenYMidpoint2Plane, 700 SPVFuncImplChromaReconstructLinear420XCositedEvenYMidpoint3Plane, 701 SPVFuncImplChromaReconstructLinear420XMidpointYMidpoint2Plane, 702 SPVFuncImplChromaReconstructLinear420XMidpointYMidpoint3Plane, 703 SPVFuncImplExpandITUFullRange, 704 SPVFuncImplExpandITUNarrowRange, 705 SPVFuncImplConvertYCbCrBT709, 706 SPVFuncImplConvertYCbCrBT601, 707 SPVFuncImplConvertYCbCrBT2020, 708 SPVFuncImplDynamicImageSampler, 709 }; 710 711 // If the underlying resource has been used for comparison then duplicate loads of that resource must be too 712 // Use Metal's native frame-buffer fetch API for subpass inputs. 713 void emit_texture_op(const Instruction &i, bool sparse) override; 714 void emit_binary_unord_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op); 715 void emit_instruction(const Instruction &instr) override; 716 void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args, 717 uint32_t count) override; 718 void emit_spv_amd_shader_trinary_minmax_op(uint32_t result_type, uint32_t result_id, uint32_t op, 719 const uint32_t *args, uint32_t count) override; 720 void emit_header() override; 721 void emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) override; 722 void emit_sampled_image_op(uint32_t result_type, uint32_t result_id, uint32_t image_id, uint32_t samp_id) override; 723 void emit_subgroup_op(const Instruction &i) override; 724 std::string to_texture_op(const Instruction &i, bool sparse, bool *forward, 725 SmallVector<uint32_t> &inherited_expressions) override; 726 void emit_fixup() override; 727 std::string to_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index, 728 const std::string &qualifier = ""); 729 void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index, 730 const std::string &qualifier = "", uint32_t base_offset = 0) override; 731 void emit_struct_padding_target(const SPIRType &type) override; 732 std::string type_to_glsl(const SPIRType &type, uint32_t id = 0) override; 733 void emit_block_hints(const SPIRBlock &block) override; 734 735 // Allow Metal to use the array<T> template to make arrays a value type 736 std::string type_to_array_glsl(const SPIRType &type) override; 737 738 // Threadgroup arrays can't have a wrapper type 739 std::string variable_decl(const SPIRVariable &variable) override; 740 741 bool variable_decl_is_remapped_storage(const SPIRVariable &variable, spv::StorageClass storage) const override; 742 743 // GCC workaround of lambdas calling protected functions (for older GCC versions) 744 std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0) override; 745 746 std::string image_type_glsl(const SPIRType &type, uint32_t id = 0) override; 747 std::string sampler_type(const SPIRType &type, uint32_t id); 748 std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage) override; 749 std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id) override; 750 std::string to_name(uint32_t id, bool allow_alias = true) const override; 751 std::string to_function_name(const TextureFunctionNameArguments &args) override; 752 std::string to_function_args(const TextureFunctionArguments &args, bool *p_forward) override; 753 std::string to_initializer_expression(const SPIRVariable &var) override; 754 std::string to_zero_initialized_expression(uint32_t type_id) override; 755 756 std::string unpack_expression_type(std::string expr_str, const SPIRType &type, uint32_t physical_type_id, 757 bool is_packed, bool row_major) override; 758 759 // Returns true for BuiltInSampleMask because gl_SampleMask[] is an array in SPIR-V, but [[sample_mask]] is a scalar in Metal. 760 bool builtin_translates_to_nonarray(spv::BuiltIn builtin) const override; 761 762 std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override; 763 bool emit_complex_bitcast(uint32_t result_id, uint32_t id, uint32_t op0) override; 764 bool skip_argument(uint32_t id) const override; 765 std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) override; 766 std::string to_qualifiers_glsl(uint32_t id) override; 767 void replace_illegal_names() override; 768 void declare_undefined_values() override; 769 void declare_constant_arrays(); 770 771 void replace_illegal_entry_point_names(); 772 void sync_entry_point_aliases_and_names(); 773 774 static const std::unordered_set<std::string> &get_reserved_keyword_set(); 775 static const std::unordered_set<std::string> &get_illegal_func_names(); 776 777 // Constant arrays of non-primitive types (i.e. matrices) won't link properly into Metal libraries 778 void declare_complex_constant_arrays(); 779 780 bool is_patch_block(const SPIRType &type); 781 bool is_non_native_row_major_matrix(uint32_t id) override; 782 bool member_is_non_native_row_major_matrix(const SPIRType &type, uint32_t index) override; 783 std::string convert_row_major_matrix(std::string exp_str, const SPIRType &exp_type, uint32_t physical_type_id, 784 bool is_packed) override; 785 786 void preprocess_op_codes(); 787 void localize_global_variables(); 788 void extract_global_variables_from_functions(); 789 void mark_packable_structs(); 790 void mark_as_packable(SPIRType &type); 791 792 std::unordered_map<uint32_t, std::set<uint32_t>> function_global_vars; 793 void extract_global_variables_from_function(uint32_t func_id, std::set<uint32_t> &added_arg_ids, 794 std::unordered_set<uint32_t> &global_var_ids, 795 std::unordered_set<uint32_t> &processed_func_ids); 796 uint32_t add_interface_block(spv::StorageClass storage, bool patch = false); 797 uint32_t add_interface_block_pointer(uint32_t ib_var_id, spv::StorageClass storage); 798 799 struct InterfaceBlockMeta 800 { 801 struct LocationMeta 802 { 803 uint32_t base_type_id = 0; 804 uint32_t num_components = 0; 805 bool flat = false; 806 bool noperspective = false; 807 bool centroid = false; 808 bool sample = false; 809 }; 810 std::unordered_map<uint32_t, LocationMeta> location_meta; 811 bool strip_array = false; 812 bool allow_local_declaration = false; 813 }; 814 815 std::string to_tesc_invocation_id(); 816 void emit_local_masked_variable(const SPIRVariable &masked_var, bool strip_array); 817 void add_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, SPIRType &ib_type, 818 SPIRVariable &var, InterfaceBlockMeta &meta); 819 void add_composite_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, 820 SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta); 821 void add_plain_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, 822 SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta); 823 bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, 824 SPIRVariable &var, const SPIRType &type, 825 InterfaceBlockMeta &meta); 826 void add_plain_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, 827 SPIRType &ib_type, SPIRVariable &var, uint32_t index, 828 InterfaceBlockMeta &meta); 829 void add_composite_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, 830 SPIRType &ib_type, SPIRVariable &var, uint32_t index, 831 InterfaceBlockMeta &meta); 832 void add_tess_level_input_to_interface_block(const std::string &ib_var_ref, SPIRType &ib_type, SPIRVariable &var); 833 834 void fix_up_interface_member_indices(spv::StorageClass storage, uint32_t ib_type_id); 835 836 void mark_location_as_used_by_shader(uint32_t location, const SPIRType &type, 837 spv::StorageClass storage, bool fallback = false); 838 uint32_t ensure_correct_builtin_type(uint32_t type_id, spv::BuiltIn builtin); 839 uint32_t ensure_correct_input_type(uint32_t type_id, uint32_t location, 840 uint32_t num_components, bool strip_array); 841 842 void emit_custom_templates(); 843 void emit_custom_functions(); 844 void emit_resources(); 845 void emit_specialization_constants_and_structs(); 846 void emit_interface_block(uint32_t ib_var_id); 847 bool maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs); 848 uint32_t get_resource_array_size(uint32_t id) const; 849 850 void fix_up_shader_inputs_outputs(); 851 852 std::string func_type_decl(SPIRType &type); 853 std::string entry_point_args_classic(bool append_comma); 854 std::string entry_point_args_argument_buffer(bool append_comma); 855 std::string entry_point_arg_stage_in(); 856 void entry_point_args_builtin(std::string &args); 857 void entry_point_args_discrete_descriptors(std::string &args); 858 std::string to_qualified_member_name(const SPIRType &type, uint32_t index); 859 std::string ensure_valid_name(std::string name, std::string pfx); 860 std::string to_sampler_expression(uint32_t id); 861 std::string to_swizzle_expression(uint32_t id); 862 std::string to_buffer_size_expression(uint32_t id); 863 bool is_sample_rate() const; 864 bool is_direct_input_builtin(spv::BuiltIn builtin); 865 std::string builtin_qualifier(spv::BuiltIn builtin); 866 std::string builtin_type_decl(spv::BuiltIn builtin, uint32_t id = 0); 867 std::string built_in_func_arg(spv::BuiltIn builtin, bool prefix_comma); 868 std::string member_attribute_qualifier(const SPIRType &type, uint32_t index); 869 std::string argument_decl(const SPIRFunction::Parameter &arg); 870 std::string round_fp_tex_coords(std::string tex_coords, bool coord_is_fp); 871 uint32_t get_metal_resource_index(SPIRVariable &var, SPIRType::BaseType basetype, uint32_t plane = 0); 872 uint32_t get_member_location(uint32_t type_id, uint32_t index, uint32_t *comp = nullptr) const; 873 uint32_t get_or_allocate_builtin_input_member_location(spv::BuiltIn builtin, 874 uint32_t type_id, uint32_t index, uint32_t *comp = nullptr); 875 876 uint32_t get_physical_tess_level_array_size(spv::BuiltIn builtin) const; 877 878 // MSL packing rules. These compute the effective packing rules as observed by the MSL compiler in the MSL output. 879 // These values can change depending on various extended decorations which control packing rules. 880 // We need to make these rules match up with SPIR-V declared rules. 881 uint32_t get_declared_type_size_msl(const SPIRType &type, bool packed, bool row_major) const; 882 uint32_t get_declared_type_array_stride_msl(const SPIRType &type, bool packed, bool row_major) const; 883 uint32_t get_declared_type_matrix_stride_msl(const SPIRType &type, bool packed, bool row_major) const; 884 uint32_t get_declared_type_alignment_msl(const SPIRType &type, bool packed, bool row_major) const; 885 886 uint32_t get_declared_struct_member_size_msl(const SPIRType &struct_type, uint32_t index) const; 887 uint32_t get_declared_struct_member_array_stride_msl(const SPIRType &struct_type, uint32_t index) const; 888 uint32_t get_declared_struct_member_matrix_stride_msl(const SPIRType &struct_type, uint32_t index) const; 889 uint32_t get_declared_struct_member_alignment_msl(const SPIRType &struct_type, uint32_t index) const; 890 891 uint32_t get_declared_input_size_msl(const SPIRType &struct_type, uint32_t index) const; 892 uint32_t get_declared_input_array_stride_msl(const SPIRType &struct_type, uint32_t index) const; 893 uint32_t get_declared_input_matrix_stride_msl(const SPIRType &struct_type, uint32_t index) const; 894 uint32_t get_declared_input_alignment_msl(const SPIRType &struct_type, uint32_t index) const; 895 896 const SPIRType &get_physical_member_type(const SPIRType &struct_type, uint32_t index) const; 897 SPIRType get_presumed_input_type(const SPIRType &struct_type, uint32_t index) const; 898 899 uint32_t get_declared_struct_size_msl(const SPIRType &struct_type, bool ignore_alignment = false, 900 bool ignore_padding = false) const; 901 902 std::string to_component_argument(uint32_t id); 903 void align_struct(SPIRType &ib_type, std::unordered_set<uint32_t> &aligned_structs); 904 void mark_scalar_layout_structs(const SPIRType &ib_type); 905 void mark_struct_members_packed(const SPIRType &type); 906 void ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t index); 907 bool validate_member_packing_rules_msl(const SPIRType &type, uint32_t index) const; 908 std::string get_argument_address_space(const SPIRVariable &argument); 909 std::string get_type_address_space(const SPIRType &type, uint32_t id, bool argument = false); 910 const char *to_restrict(uint32_t id, bool space = true); 911 SPIRType &get_stage_in_struct_type(); 912 SPIRType &get_stage_out_struct_type(); 913 SPIRType &get_patch_stage_in_struct_type(); 914 SPIRType &get_patch_stage_out_struct_type(); 915 std::string get_tess_factor_struct_name(); 916 SPIRType &get_uint_type(); 917 uint32_t get_uint_type_id(); 918 void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, uint32_t mem_order_1, 919 uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0, 920 bool op1_is_pointer = false, bool op1_is_literal = false, uint32_t op2 = 0); 921 const char *get_memory_order(uint32_t spv_mem_sem); 922 void add_pragma_line(const std::string &line); 923 void add_typedef_line(const std::string &line); 924 void emit_barrier(uint32_t id_exe_scope, uint32_t id_mem_scope, uint32_t id_mem_sem); 925 void emit_array_copy(const std::string &lhs, uint32_t lhs_id, uint32_t rhs_id, 926 spv::StorageClass lhs_storage, spv::StorageClass rhs_storage) override; 927 void build_implicit_builtins(); 928 uint32_t build_constant_uint_array_pointer(); 929 void emit_entry_point_declarations() override; 930 uint32_t builtin_frag_coord_id = 0; 931 uint32_t builtin_sample_id_id = 0; 932 uint32_t builtin_sample_mask_id = 0; 933 uint32_t builtin_vertex_idx_id = 0; 934 uint32_t builtin_base_vertex_id = 0; 935 uint32_t builtin_instance_idx_id = 0; 936 uint32_t builtin_base_instance_id = 0; 937 uint32_t builtin_view_idx_id = 0; 938 uint32_t builtin_layer_id = 0; 939 uint32_t builtin_invocation_id_id = 0; 940 uint32_t builtin_primitive_id_id = 0; 941 uint32_t builtin_subgroup_invocation_id_id = 0; 942 uint32_t builtin_subgroup_size_id = 0; 943 uint32_t builtin_dispatch_base_id = 0; 944 uint32_t builtin_stage_input_size_id = 0; 945 uint32_t builtin_local_invocation_index_id = 0; 946 uint32_t builtin_workgroup_size_id = 0; 947 uint32_t swizzle_buffer_id = 0; 948 uint32_t buffer_size_buffer_id = 0; 949 uint32_t view_mask_buffer_id = 0; 950 uint32_t dynamic_offsets_buffer_id = 0; 951 uint32_t uint_type_id = 0; 952 uint32_t argument_buffer_padding_buffer_type_id = 0; 953 uint32_t argument_buffer_padding_image_type_id = 0; 954 uint32_t argument_buffer_padding_sampler_type_id = 0; 955 956 bool does_shader_write_sample_mask = false; 957 958 void cast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override; 959 void cast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) override; 960 void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) override; 961 962 void analyze_sampled_image_usage(); 963 964 bool access_chain_needs_stage_io_builtin_translation(uint32_t base) override; 965 void prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type, spv::StorageClass storage, 966 bool &is_packed) override; 967 void fix_up_interpolant_access_chain(const uint32_t *ops, uint32_t length); 968 bool emit_tessellation_access_chain(const uint32_t *ops, uint32_t length); 969 bool emit_tessellation_io_load(uint32_t result_type, uint32_t id, uint32_t ptr); 970 bool is_out_of_bounds_tessellation_level(uint32_t id_lhs); 971 972 void ensure_builtin(spv::StorageClass storage, spv::BuiltIn builtin); 973 974 void mark_implicit_builtin(spv::StorageClass storage, spv::BuiltIn builtin, uint32_t id); 975 976 std::string convert_to_f32(const std::string &expr, uint32_t components); 977 978 Options msl_options; 979 std::set<SPVFuncImpl> spv_function_implementations; 980 // Must be ordered to ensure declarations are in a specific order. 981 std::map<uint32_t, MSLShaderInput> inputs_by_location; 982 std::unordered_map<uint32_t, MSLShaderInput> inputs_by_builtin; 983 std::unordered_set<uint32_t> location_inputs_in_use; 984 std::unordered_set<uint32_t> location_inputs_in_use_fallback; 985 std::unordered_map<uint32_t, uint32_t> fragment_output_components; 986 std::unordered_map<uint32_t, uint32_t> builtin_to_automatic_input_location; 987 std::set<std::string> pragma_lines; 988 std::set<std::string> typedef_lines; 989 SmallVector<uint32_t> vars_needing_early_declaration; 990 991 std::unordered_map<StageSetBinding, std::pair<MSLResourceBinding, bool>, InternalHasher> resource_bindings; 992 std::unordered_map<StageSetBinding, uint32_t, InternalHasher> resource_arg_buff_idx_to_binding_number; 993 994 uint32_t next_metal_resource_index_buffer = 0; 995 uint32_t next_metal_resource_index_texture = 0; 996 uint32_t next_metal_resource_index_sampler = 0; 997 // Intentionally uninitialized, works around MSVC 2013 bug. 998 uint32_t next_metal_resource_ids[kMaxArgumentBuffers]; 999 1000 VariableID stage_in_var_id = 0; 1001 VariableID stage_out_var_id = 0; 1002 VariableID patch_stage_in_var_id = 0; 1003 VariableID patch_stage_out_var_id = 0; 1004 VariableID stage_in_ptr_var_id = 0; 1005 VariableID stage_out_ptr_var_id = 0; 1006 VariableID stage_out_masked_builtin_type_id = 0; 1007 1008 // Handle HLSL-style 0-based vertex/instance index. 1009 enum class TriState 1010 { 1011 Neutral, 1012 No, 1013 Yes 1014 }; 1015 TriState needs_base_vertex_arg = TriState::Neutral; 1016 TriState needs_base_instance_arg = TriState::Neutral; 1017 1018 bool has_sampled_images = false; 1019 bool builtin_declaration = false; // Handle HLSL-style 0-based vertex/instance index. 1020 1021 bool is_using_builtin_array = false; // Force the use of C style array declaration. 1022 bool using_builtin_array() const; 1023 1024 bool is_rasterization_disabled = false; 1025 bool capture_output_to_buffer = false; 1026 bool needs_swizzle_buffer_def = false; 1027 bool used_swizzle_buffer = false; 1028 bool added_builtin_tess_level = false; 1029 bool needs_subgroup_invocation_id = false; 1030 bool needs_subgroup_size = false; 1031 bool needs_sample_id = false; 1032 std::string qual_pos_var_name; 1033 std::string stage_in_var_name = "in"; 1034 std::string stage_out_var_name = "out"; 1035 std::string patch_stage_in_var_name = "patchIn"; 1036 std::string patch_stage_out_var_name = "patchOut"; 1037 std::string sampler_name_suffix = "Smplr"; 1038 std::string swizzle_name_suffix = "Swzl"; 1039 std::string buffer_size_name_suffix = "BufferSize"; 1040 std::string plane_name_suffix = "Plane"; 1041 std::string input_wg_var_name = "gl_in"; 1042 std::string input_buffer_var_name = "spvIn"; 1043 std::string output_buffer_var_name = "spvOut"; 1044 std::string patch_output_buffer_var_name = "spvPatchOut"; 1045 std::string tess_factor_buffer_var_name = "spvTessLevel"; 1046 std::string index_buffer_var_name = "spvIndices"; 1047 spv::Op previous_instruction_opcode = spv::OpNop; 1048 1049 // Must be ordered since declaration is in a specific order. 1050 std::map<uint32_t, MSLConstexprSampler> constexpr_samplers_by_id; 1051 std::unordered_map<SetBindingPair, MSLConstexprSampler, InternalHasher> constexpr_samplers_by_binding; 1052 const MSLConstexprSampler *find_constexpr_sampler(uint32_t id) const; 1053 1054 std::unordered_set<uint32_t> buffers_requiring_array_length; 1055 SmallVector<uint32_t> buffer_arrays; 1056 std::unordered_set<uint32_t> atomic_image_vars; // Emulate texture2D atomic operations 1057 std::unordered_set<uint32_t> pull_model_inputs; 1058 1059 // Must be ordered since array is in a specific order. 1060 std::map<SetBindingPair, std::pair<uint32_t, uint32_t>> buffers_requiring_dynamic_offset; 1061 1062 SmallVector<uint32_t> disabled_frag_outputs; 1063 1064 std::unordered_set<SetBindingPair, InternalHasher> inline_uniform_blocks; 1065 1066 uint32_t argument_buffer_ids[kMaxArgumentBuffers]; 1067 uint32_t argument_buffer_discrete_mask = 0; 1068 uint32_t argument_buffer_device_storage_mask = 0; 1069 1070 void analyze_argument_buffers(); 1071 bool descriptor_set_is_argument_buffer(uint32_t desc_set) const; 1072 MSLResourceBinding &get_argument_buffer_resource(uint32_t desc_set, uint32_t arg_idx); 1073 void add_argument_buffer_padding_buffer_type(SPIRType &struct_type, uint32_t &mbr_idx, uint32_t &arg_buff_index, MSLResourceBinding &rez_bind); 1074 void add_argument_buffer_padding_image_type(SPIRType &struct_type, uint32_t &mbr_idx, uint32_t &arg_buff_index, MSLResourceBinding &rez_bind); 1075 void add_argument_buffer_padding_sampler_type(SPIRType &struct_type, uint32_t &mbr_idx, uint32_t &arg_buff_index, MSLResourceBinding &rez_bind); 1076 void add_argument_buffer_padding_type(uint32_t mbr_type_id, SPIRType &struct_type, uint32_t &mbr_idx, uint32_t &arg_buff_index, uint32_t count); 1077 1078 uint32_t get_target_components_for_fragment_location(uint32_t location) const; 1079 uint32_t build_extended_vector_type(uint32_t type_id, uint32_t components, 1080 SPIRType::BaseType basetype = SPIRType::Unknown); 1081 uint32_t build_msl_interpolant_type(uint32_t type_id, bool is_noperspective); 1082 1083 bool suppress_missing_prototypes = false; 1084 1085 void add_spv_func_and_recompile(SPVFuncImpl spv_func); 1086 1087 void activate_argument_buffer_resources(); 1088 1089 bool type_is_msl_framebuffer_fetch(const SPIRType &type) const; 1090 bool type_is_pointer(const SPIRType &type) const; 1091 bool type_is_pointer_to_pointer(const SPIRType &type) const; 1092 bool is_supported_argument_buffer_type(const SPIRType &type) const; 1093 1094 bool variable_storage_requires_stage_io(spv::StorageClass storage) const; 1095 has_additional_fixed_sample_mask() const1096 bool has_additional_fixed_sample_mask() const { return msl_options.additional_fixed_sample_mask != 0xffffffff; } 1097 std::string additional_fixed_sample_mask_str() const; 1098 1099 // OpcodeHandler that handles several MSL preprocessing operations. 1100 struct OpCodePreprocessor : OpcodeHandler 1101 { OpCodePreprocessorSPIRV_CROSS_NAMESPACE::CompilerMSL::OpCodePreprocessor1102 OpCodePreprocessor(CompilerMSL &compiler_) 1103 : compiler(compiler_) 1104 { 1105 } 1106 1107 bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override; 1108 CompilerMSL::SPVFuncImpl get_spv_func_impl(spv::Op opcode, const uint32_t *args); 1109 void check_resource_write(uint32_t var_id); 1110 1111 CompilerMSL &compiler; 1112 std::unordered_map<uint32_t, uint32_t> result_types; 1113 std::unordered_map<uint32_t, uint32_t> image_pointers; // Emulate texture2D atomic operations 1114 bool suppress_missing_prototypes = false; 1115 bool uses_atomics = false; 1116 bool uses_resource_write = false; 1117 bool needs_subgroup_invocation_id = false; 1118 bool needs_subgroup_size = false; 1119 bool needs_sample_id = false; 1120 }; 1121 1122 // OpcodeHandler that scans for uses of sampled images 1123 struct SampledImageScanner : OpcodeHandler 1124 { SampledImageScannerSPIRV_CROSS_NAMESPACE::CompilerMSL::SampledImageScanner1125 SampledImageScanner(CompilerMSL &compiler_) 1126 : compiler(compiler_) 1127 { 1128 } 1129 1130 bool handle(spv::Op opcode, const uint32_t *args, uint32_t) override; 1131 1132 CompilerMSL &compiler; 1133 }; 1134 1135 // Sorts the members of a SPIRType and associated Meta info based on a settable sorting 1136 // aspect, which defines which aspect of the struct members will be used to sort them. 1137 // Regardless of the sorting aspect, built-in members always appear at the end of the struct. 1138 struct MemberSorter 1139 { 1140 enum SortAspect 1141 { 1142 LocationThenBuiltInType, 1143 Offset 1144 }; 1145 1146 void sort(); 1147 bool operator()(uint32_t mbr_idx1, uint32_t mbr_idx2); 1148 MemberSorter(SPIRType &t, Meta &m, SortAspect sa); 1149 1150 SPIRType &type; 1151 Meta &meta; 1152 SortAspect sort_aspect; 1153 }; 1154 }; 1155 } // namespace SPIRV_CROSS_NAMESPACE 1156 1157 #endif 1158