Metal Shading Language Specification

NOTE:  To decide which address space (
device
 or 
constant
) a read-only buffer passed 
to a graphics or kernel function uses, look at how the buffer is accessed inside the 
graphics or kernel function. The 
constant
 address space is optimized for multiple 
instances executing a graphics or kernel function accessing the same location in the 
buffer. Some examples of this access pattern are accessing light or material properties 
for lighting / shading, matrix of a matrix array used for skinning, filter weight accessed 
from a filter weight array for convolution. If multiple executing instances of a graphics 
or kernel function are accessing the buffer using an index such as the vertex ID, 
fragment coordinate, or the thread position in grid, the buffer must be allocated in the 
device
 address space.  
4.2.5   
thread Address Space 
The 
thread
 address space refers to the per-thread memory address space. Variables allocated 
in this address space are not visible to other threads. Variables declared inside a graphics or 
kernel function are allocated in the 
thread
 address space.  
kernel void 
my_kernel(…) 
{ 
// A float allocated in the per-thread address space 
float x; 
// A pointer to variable x in per-thread address space 
thread float p = &x; 
… 
} 
4.3  Function Arguments and Variables 
Most inputs and outputs to a graphics and kernel functions are passed as arguments. 
(Exceptions are initialized variables in the constant address space and samplers declared in 
program scope.) Arguments to graphics (vertex and fragment) and kernel functions can be one 
of the following: 
• device buffer – a pointer or reference to any data type in the 
device
 address space (see 
section 2.6) 
• constant buffer – a pointer or reference to any data type in the 
constant
 address space 
(see section 2.6) 
•
texture
 object (see section 2.8) or an array of textures 
•
sampler
 object (see section 2.9) or an array of samplers 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
58
174

---

• a buffer shared between threads in a threadgroup – a pointer to a type in the 
threadgroup
 address space that can only be used as arguments for kernel functions. 
• imageblock (see section 2.10) 
• argument buffer (see section 2.12) 
• A struct with elements that are either buffers or textures.  
Buffers (device and constant) specified as argument values to a graphics or kernel function 
cannot alias; i.e., a buffer passed as an argument value cannot overlap another buffer passed to 
a separate argument of the same graphics or kernel function. 
Arguments to graphics and kernel functions cannot be declared to be of type 
size_t
, 
ptrdiff_t
, or a struct and/or union that contains members declared to be one of these built-in 
scalar types. 
The arguments to these functions are often specified with attributes to provide further guidance 
on their use. Attributes are used to specify: 
• the resource location for the argument (see section 4.3.1),  
• built-in variables that support communicating data between fixed-function and 
programmable pipeline stages (see section 4.3.3), 
• which data is sent down the pipeline from vertex function to fragment function (see 
section 4.3.5). 
4.3.1   
Attributes to Locate Buffers, Textures and Samplers  
For each argument, an attribute can be optionally specified to identify the location of a buffer, 
texture or sampler to use for this argument type. The Metal framework API uses this attribute to 
identify the location for these argument types.  
• device and constant buffers – 
[[buffer(index)]]
 
• texture and an array of textures – 
[[texture(index)]]
 
• sampler – 
[[sampler(index)]]
 
• threadgroup buffer – 
[[threadgroup(index)]]
 
The 
index
 value is an unsigned integer that identifies the location of a buffer, texture or sampler 
argument that is being assigned. The proper syntax is for the attribute to follow the argument/
variable name.  
The example below is a simple kernel function, 
add_vectors
, that adds an array of two buffers 
in the device address space, 
inA
 and 
inB
, and returns the result in the buffer out. The attributes 
(
buffer(index)
) specify the buffer locations for the function arguments.  
kernel void  
add_vectors(const device float4 *inA [[buffer(0)]], 
  const device float4 *inB [[buffer(1)]], 
  device float4 *out [[buffer(2)]], 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
59
174

---

  uint id [[thread_position_in_grid]]) 
{ 
out[id] = inA[id] + inB[id]; 
} 
The example below shows attributes used for function arguments of several different types (a 
buffer, a texture, and a sampler): 
kernel void 
my_kernel(device float4 *p [[buffer(0)]], 
texture2d img [[texture(0)]], 
sampler sam [[sampler(1)]]) 
{…} 
If the location indices are not specified the Metal compiler assigns them using the first available 
location index. In the following example, 
src
 is assigned texture index 0, 
dst
 texture index 1, 
s
 
sampler index 0 and 
u
 buffer index 0.: 
kernel void 
my_kernel(texture2d src, 
 
 
    texture2d dst, 
 
 
    sampler s, 
 
 
    device myUserInfo *u) 
{…} 
In the following example, some kernel arguments have explicitly assigned location indices and 
some do not. 
src
 is explicitly assigned texture index 0, and 
f
 is explicitly assigned buffer index 
10. The other arguments are assigned the first available location index: 
dst
 texture index 1, 
s
 
sampler index 0, and 
u
 buffer index 0.  
kernel void 
my_kernel(texture2d src [[texture(0)]], 
     texture2d dst, 
sampler s, 
device myUserInfo *u, 
device float *f [[buffer(10)]]) 
{…} 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
60
174

---