How SPIRV-Cross Translates SPIR-V to Metal MSL
This article provides an overview of how SPIRV-Cross decompiles and transforms SPIR-V bytecode—originally compiled from GLSL—into valid Metal Shading Language (MSL) code. The translation process relies on parsing the binary instructions into an Abstract Syntax Tree (AST), resolving structural and semantic mismatches between Vulkan-style shaders and Apple's Metal graphics API, remapping resource bindings, and emitting idiomatic, compile-ready C++-based MSL source code.
1. Bytecode Parsing and AST Construction
The translation pipeline begins by ingesting the raw SPIR-V binary, which consists of a stream of 32-bit words containing operation codes (opcodes) and operands. SPIRV-Cross iterates through these opcodes sequentially to identify data types, constants, global variables, function definitions, and execution flow instructions.
Instead of translating opcode-by-opcode directly into source text,
SPIRV-Cross builds an internal high-level representation (AST) along
with a Control Flow Graph (CFG). This intermediate model reconstructs
control-flow structures such as if-else branches, loops,
and switch statements from low-level SPIR-V branch
instructions (OpBranch,
OpBranchConditional).
2. Type System and Semantic Mapping
Once the AST is formed, SPIRV-Cross maps SPIR-V types and operations to their MSL equivalents:
- Basic Types and Vectors: SPIR-V vector and scalar
definitions are converted into Metal standard types (e.g.,
OpTypeVectorwith 32-bit floats becomesmetal::float4,OpTypeMatrixbecomesmetal::float4x4). - Built-in Variables: GLSL built-ins captured in
SPIR-V are mapped to Metal attributes. For example,
gl_Positiontranslates to the[[position]]attribute,gl_VertexIndexbecomes[[vertex_id]], andgl_FragCoordbecomes a parameter marked with[[position]]. - Standard Intrinsics: GLSL extended instructions
(GLSL.std.450), such as mathematical functions (
mix,clamp,dot), are converted into corresponding functions inside themetal::namespace.
3. Reconciling Resource Binding Differences
Vulkan and Metal use fundamentally different resource-binding models.
GLSL-derived SPIR-V uses descriptor sets and bindings (e.g.,
layout(set = X, binding = Y)), whereas MSL addresses
resources via explicit hardware binding tables:
[[buffer(n)]], [[texture(n)]], and
[[sampler(n)]].
SPIRV-Cross resolves this discrepancy by maintaining a resource
remapping table. It analyzes the descriptor sets present in the bytecode
and assigns them to specific linear indices compatible with Metal.
SPIRV-Cross can either flatten descriptor sets into flat hardware table
slots or synthesize Metal Argument Buffers (structures marked with
[[buffer(n)]] containing embedded resources) to preserve
the original descriptor layout.
4. Splitting Combined Image Samplers
In GLSL, textures and samplers are frequently declared as combined
objects (such as sampler2D). In contrast, Metal strictly
decouples sampled textures (texture2d<float>) from
the sampling state (sampler).
When SPIR-V contains OpTypeSampledImage, SPIRV-Cross
decomposes the combined object into two distinct MSL arguments: a
texture2d and a sampler. At the call sites
where the texture is read, SPIRV-Cross transforms operations like
texture() into Metal member function calls on the texture
object (e.g., tex.sample(smp, coords)), passing both
resolved resources.
5. Function Signature and Memory Layout Generation
MSL shaders are functions with specific stage attributes
(vertex, fragment, kernel) rather
than a single void main() entry point that relies on global
state.
SPIRV-Cross aggregates all uniform buffers, storage buffers, inputs,
outputs, and push constants referenced by the SPIR-V entry point and
restructures them into explicit parameters for the generated MSL
entry-point function. Stage inputs are packed into an MSL struct tagged
with the [[stage_in]] attribute, while outputs are returned
as a structured value. Additionally, SPIR-V push constant blocks are
converted to standard buffers bound to a designated
[[buffer(n)]] slot.
6. MSL Code Emission
The final phase uses the CompilerMSL backend to traverse
the finalized AST and emit valid MSL text. This stage ensures that
language-specific constraints are respected:
- Inserts required headers, primarily
#include <metal_stdlib>andusing namespace metal;. - Resolves naming collisions and sanitizes identifiers that clash with reserved Metal keywords.
- Enforces proper memory address space qualifiers
(
device,constant,thread, orthreadgroup) on pointers and references.
The resulting output is a self-contained, human-readable MSL text
file capable of being compiled directly by Apple's Metal shader compiler
(metallib).