1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
|
---
layout: user-guide
permalink: /user-guide/metal-target-specific
---
# 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
## 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. The same as task/mesh shaders generally in Slang.
```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.
```cpp
#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...
```slang
struct MyParameters
{
int x;
int y;
StructuredBuffer<float> buffer1;
RWStructuredBuffer<uint3> buffer2;
}
ParameterBlock<MyParameters> gObj;
void main(){ ... gObj ... }
```
... results in this Metal output:
```cpp
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:
```slang
struct NestedStruct
{
float2 uv;
};
struct InputStruct
{
float4 position;
float3 normal;
NestedStruct nested;
};
```
Will be flattened to:
```cpp
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:
```slang
float4 main() : SV_Target
{
return float4(1,2,3,4);
}
```
becomes:
```c++
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:
```slang
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.
## 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:
```csharp
[vk::constant_id(7)]
const int a = 2;
```
Translates to:
```metal
constant int fc_a_0 [[function_constant(7)]];
constant int a_0 = is_function_constant_defined(fc_a_0) ? fc_a_0 : 2;
```
|
