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