Metal-Specific Functionalities

This chapter provides information for Metal-specific functionalities and behaviors in Slang.

Entry Point Parameter Handling

Slang performs several transformations on entry point parameters when targeting Metal:

  • Struct parameters are flattened to eliminate nested structures

  • Input parameters with varying inputs are packed into a single struct

  • System value semantics are translated to Metal attributes

  • Parameters without semantics are given automatic attribute indices

Limitation: isParameterLocationUsed for varying inputs

Because Metal entry point varying inputs are packed into a single [[stage_in]] struct parameter during legalization, IMetadata::isParameterLocationUsed cannot distinguish between used and unused individual varying inputs. The metadata is recorded at the struct level, so if any varying input is used, all varying input locations covered by the struct will be reported as used. This differs from SPIR-V, where each varying input becomes a separate global parameter that can be individually tracked.

This limitation does not affect non-varying resource types (e.g. descriptor table slots), which are tracked individually.

System-Value semantics

The system-value semantics are translated to the following Metal attributes:

SV semantic name

Metal attribute

SV_Position

[[position]]

SV_Coverage

[[sample_mask]]

SV_Depth

[[depth(any)]]

SV_DepthGreaterEqual

[[depth(greater)]]

SV_DepthLessEqual

[[depth(less)]]

SV_DispatchThreadID

[[thread_position_in_grid]]

SV_FragInvocationCount

(Not supported)

SV_FragSize

(Not supported)

SV_GroupID

[[threadgroup_position_in_grid]]

SV_GroupThreadID

[[thread_position_in_threadgroup]]

SV_GroupIndex

Calculated from SV_GroupThreadID and group extents

SV_InstanceID

[[instance_id]]

SV_IsFrontFace

[[front_facing]]

SV_PointSize

[[point_size]]

SV_PointCoord

[[point_coord]]

SV_PrimitiveID

[[primitive_id]]

SV_RenderTargetArrayIndex

[[render_target_array_index]]

SV_SampleIndex

[[sample_id]]

SV_Target<N>

[[color(N)]]

SV_VertexID

[[vertex_id]]

SV_ViewportArrayIndex

[[viewport_array_index]]

SV_StartVertexLocation

[[base_vertex]]

SV_StartInstanceLocation

[[base_instance]]

SV_VulkanInstanceID

[[instance_id]]

SV_VulkanSamplePosition

(Not supported)

SV_VulkanVertexID

[[vertex_id]]

Custom semantics are mapped to user attributes:

  • [[user(SEMANTIC_NAME)]] For non-system value semantics

  • [[user(SEMANTIC_NAME_INDEX)]] When semantic has an index

Interpolation Modifiers

Slang maps interpolation modifiers to Metal’s interpolation attributes:

Slang Interpolation

Metal Attribute

nointerpolation

[[flat]]

noperspective

[[center_no_perspective]]

linear

[[sample_no_perspective]]

sample

[[sample_perspective]]

centroid

[[center_perspective]]

Resource Types

Resource types are translated with appropriate Metal qualifiers:

Slang Type

Metal Translation

Texture2D

texture2d

RWTexture2D

texture2d

ByteAddressBuffer

uint32_t device*

StructuredBuffer<T>

device* T

ConstantBuffer<T>

constant* T

Slang Type

Metal Translation

Texture1D

texture1d

Texture1DArray

texture1d_array

RWTexture1D

texture1d

RWTexture1DArray

texture1d_array

Texture2D

texture2d

Texture2DArray

texture2d_array

RWTexture2D

texture2d

RWTexture2DArray

texture2d_array

Texture3D

texture3d

RWTexture3D

texture3d

TextureCube

texturecube

TextureCubeArray

texturecube_array

Buffer<T>

device* T

RWBuffer<T>

device* T

ByteAddressBuffer

device* uint32_t

RWByteAddressBuffer

device* uint32_t

StructuredBuffer<T>

device* T

RWStructuredBuffer<T>

device* T

AppendStructuredBuffer<T>

device* T

ConsumeStructuredBuffer<T>

device* T

ConstantBuffer<T>

constant* T

SamplerState

sampler

SamplerComparisonState

sampler

RaytracingAccelerationStructure

(Not supported)

RasterizerOrderedTexture2D

texture2d [[raster_order_group(0)]]

RasterizerOrderedBuffer<T>

device* T [[raster_order_group(0)]]

Raster-ordered access resources receive the [[raster_order_group(0)]] attribute, for example texture2d<float, access::read_write> tex [[raster_order_group(0)]].

Array Types

Array types in Metal are declared using the array template:

Slang Type

Metal Translation

ElementType[Size]

array<ElementType, Size>

Matrix Layout

Metal exclusively uses column-major matrix layout. Slang automatically handles the translation of matrix operations to maintain correct semantics:

  • Matrix multiplication is transformed to account for layout differences

  • Matrix types are declared as matrix<T, Columns, Rows>, for example float3x4 is represented as matrix<float, 3, 4>

Mesh Shader Support

Mesh shaders can be targeted using the following types and syntax, which is the same as for task/mesh shaders generally in Slang.

[outputtopology("triangle")]
[numthreads(12, 1, 1)]
void meshMain(
    in uint tig: SV_GroupIndex,
    in payload MeshPayload meshPayload,
    OutputVertices<Vertex, MAX_VERTS> verts,
    OutputIndices<uint3, MAX_PRIMS> triangles,
    OutputPrimitives<Primitive, MAX_PRIMS> primitives
    )

Header Inclusions and Namespace

When targeting Metal, Slang automatically includes the following headers, these are available to any intrinsic code.

#include <metal_stdlib>
#include <metal_math>
#include <metal_texture>
using namespace metal;

Parameter blocks and Argument Buffers

ParameterBlock values are translated into Argument Buffers potentially containing nested resources. For example, this Slang code…

struct MyParameters
{
    int x;
    int y;
    StructuredBuffer<float> buffer1;
    RWStructuredBuffer<uint3> buffer2;
}

ParameterBlock<MyParameters> gObj;

void main(){ ... gObj ... }

… results in this Metal output:

struct MyParameters
{
    int x;
    int y;
    float device* buffer1;
    uint3 device* buffer2;
};

[[kernel]] void main(MyParameters constant* gObj [[buffer(1)]])

Struct Parameter Flattening

When targeting Metal, top-level nested struct parameters are automatically flattened. For example:

struct NestedStruct
{
    float2 uv;
};
struct InputStruct
{
    float4 position;
    float3 normal;
    NestedStruct nested;
};

Will be flattened to:

struct InputStruct
{
    float4 position;
    float3 normal;
    float2 uv;
};

Return Value Handling

Non-struct return values from entry points are automatically wrapped in a struct with appropriate semantics. For example:

float4 main() : SV_Target
{
    return float4(1,2,3,4);
}

becomes:

struct FragmentOutput
{
    float4 value : SV_Target;
};
FragmentOutput main()
{
    return { float4(1,2,3,4) };
}

Value Type Conversion

Metal enforces strict type requirements for certain operations. Slang automatically performs the following conversions:

  • Vector size expansion (e.g., float2 to float4), for example when the user specified float2 but the semantic type in Metal is float4.

  • Image store value expansion to 4-components

For example:

RWTexture2D<float2> tex;
tex[coord] = float2(1,2);  // Automatically expanded to float4(1,2,0,0)

Conservative Rasterization

Since Metal doesn’t support conservative rasterization, SV_InnerCoverage is always false.

SubpassInput

Slang supports SubpassInput and SubpassInputMS on Metal, lowered to Metal’s framebuffer fetch. The [[vk::input_attachment_index(N)]] attribute maps to Metal’s [[color(N)]] fragment input, allowing a fragment shader to read the current value in the render target within the same render pass.

[[vk::input_attachment_index(0)]] SubpassInput<float4> color;
[[vk::input_attachment_index(1)]] SubpassInput<float4> prevBloom;

[shader("fragment")]
float4 main() : SV_Target
{
    return color.SubpassLoad() + prevBloom.SubpassLoad();
}

This generates a Metal fragment function with [[color(N)]] parameters. The SV_Target return value is wrapped in a Slang-generated output struct:

struct pixelOutput { float4 output [[color(0)]]; };

[[fragment]] pixelOutput main(float4 color [[color(0)]], float4 prevBloom [[color(1)]]);

Note that [[vk::input_attachment_index(N)]] always maps to [[color(N)]] regardless of what the attachment semantically holds — Metal framebuffer fetch reads from color attachments only, not from [[depth]]/[[stencil]] attachments.

Restrictions

  • Fragment stage only. SubpassInput can only be used in fragment entry points.

  • Direct entry-point parameters only. SubpassInput cannot be placed inside a ParameterBlock. It must be a global-scope declaration directly referenced by the entry point.

  • No cross-function references. SubpassInput must be referenced only from the fragment entry point or from helper functions that are inlined into it. Mark any helper that calls SubpassLoad with [ForceInline] to guarantee inlining; [noinline] helpers will fail with error E56108.

  • SubpassInputMS per-sample limitation. Metal does not support per-sample reads. When using SubpassInputMS, the sample index passed to SubpassLoad is ignored and the resolved value is returned. The compiler emits warning E56107 at each call site where a sample index is provided.

Address Space Assignment

Metal requires explicit address space qualifiers. Slang automatically assigns appropriate address spaces:

Variable Type

Metal Address Space

Local Variables

thread

Global Variables

device

Uniform Buffers

constant

RW/Structured Buffers

device

Group Shared

threadgroup

Parameter Blocks

constant

Explicit Parameter Binding

The HLSL :register() semantic is respected when emitting Metal code.

Since Metal does not differentiate between a constant buffer, a shader resource (read-only) buffer and an unordered access buffer, Slang will map register(tN), register(uN) and register(bN) to [[buffer(N)]] when such register semantic is declared on a buffer-typed parameter.

spaceN specifiers inside register semantics are ignored.

The [vk::location(N)] attributes on stage input/output parameters are respected.

Specialization Constants

Specialization constants declared with the [SpecializationConstant] or [vk::constant_id] attribute will be translated into a function_constant when generating Metal source. For example:

[vk::constant_id(7)]
const int a = 2;

Translates to:

constant int fc_a_0 [[function_constant(7)]];
constant int a_0 = is_function_constant_defined(fc_a_0) ? fc_a_0 : 2;