Metal Shading Language Specification

atomic_bool 
atomic
 
atomic
 represents templated types, where 
T
 can be 
int
, 
uint
, or 
bool
.  
The Metal atomic data type is restricted for use by Metal atomic functions, as described in 
section 5.12. These atomic functions are a subset of the C++14 atomic and synchronization 
functions. 
2.6  Pixel Data Types 
The Metal pixel data type is a templated type that describes the pixel format type and its 

corresponding ALU type. The ALU type represents the type returned by a load operation 

and the input type specified for a store operation. The pixel data types are generally 

available in all address spaces (for details on address spaces, see section 4.2).  
Table 6 lists supported pixel data types in the Metal shading language, as well as their size and 
alignment. 

Table 6 Metal Pixel Data Types 
Pixel Data Type
Supported values 
of T
Size (in bytes)
Alignment (in 
bytes)
r8unorm
half
 or 
float
1
1
r8snorm
half
 or 
float
1
1
r16unorm
float
2
2
r16snorm
float
2
2
rg8unorm
half2
 or 
float2
2
1
rg8snorm
half2
 or 
float2
2
1
rg16unorm
float2
4
2
rg16snorm
float2
4
2
rgba8unorm
half4
 or 
float4
4
1
srgba8unorm
half4
 or 
float4
4
1
rgba8snorm
half4
 or 
float4
4
1
rgba16unorm
float4
8
2
rgba16snorm
float4
8
2
rgb10a2
half4
 or 
float4
4
4
rg11b10f
half3
 or 
float3
4
4
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
23
174

---

Only assignments and equality/inequality comparisons between the pixel data types and their 
corresponding ALU types are allowed. 
Examples:  
kernel void 
my_kernel(device rgba8unorm *p [[buffer(0)]],  
     uint gid [[thread_position_in_grid]], …)  
{ 
rgba8unorm x = p[index]; half4 val = p[gid]; 
… 
p[gid] = val; 
p[index] = x;  
}  
struct Foo {  
rgba8unorm a;  
};  
kernel void 
my_kernel(device Foo *p [[buffer(0)]],  
     uint gid [[thread_position_in_grid]], …)  
{ 
half4 a = p[gid].a; 
… 
p[gid].a = a;  
}  
2.7  Buffers 
Metal implements buffers as a pointer to a built-in or user defined data type described in the 
device
, 
threadgroup
, or 
constant
 address space. (Refer to section 4.2 for a full description 
of these address attributes.). These buffers can be declared in program scope or passed as 
arguments to a function. 
rgb9e5
half3
 or 
float3
4
4
Pixel Data Type
Supported values 
of T
Size (in bytes)
Alignment (in 
bytes)
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
24
174

---

Examples: 
device float4 *device_buffer; 
struct my_user_data { 
float4 a; 
float b; 
int2 c; 
}; 
constant my_user_data *user_data; 
Ordinary Metal buffers may contain: 
• Basic types such as 
float
 and 
int
 
• Vector and matrix types 
• Arrays of buffer types 
• Structs of buffer types 
• Unions of buffer types 
For argument buffers, see section 2.12. 
2.8  Textures  
The texture data type is a handle to one-, two-, or three-dimensional texture data that 
corresponds to all or a portion of a single mipmap level of a texture. The following templates 
define specific texture data types:  
enum class access { sample, read, write, read_write }; 
texture1d 
texture1d_array 
texture2d 
texture2d_array 
texture3d 
texturecube 
texturecube_array 
texture2d_ms 
Textures with depth formats must be declared as one of the following texture data types: 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
25
174

---

depth2d 
depth2d_array 
depthcube 
depthcube_array 
depth2d_ms 
T
 specifies the color type returned when reading from a texture or the color type specified when 
writing to the texture. For texture types (except depth texture types), 
T
 can be 
half
, 
float
, 
short
, 
ushort
, 
int
, or 
uint
. For depth texture types, 
T
 must be 
float
.      
NOTE:  If 
T
 is 
int
 or 
short
, the data associated with the texture must use a signed 
integer format. If 
T
 is 
uint
 or 
ushort
, the data associated with the texture must use an 
unsigned integer format. If 
T
 is 
half
, the data associated with the texture must either be 
a normalized (signed or unsigned integer) or half precision format. If 
T
 is 
float
, the data 
associated with the texture must either be a normalized (signed or unsigned integer), 
half or single precision format.  
The 
access
 attribute describes how the texture can be accessed. The supported access 
attributes are: 
•
sample
 – The texture object can be sampled. 
sample
 implies the ability to read from a 
texture with and without a sampler.  
•
read
 – Without a sampler, a graphics or kernel function can only read the texture object.  
•
write
 – A graphics or kernel function can write to the texture object. 
•
read_write
 – A graphics or kernel function can read and write to the texture object. 
NOTE: 
• For multisampled textures, only the 
read
 attribute is supported. 
• For depth textures, only the 
sample
 and 
read
 attributes are supported. 
The following example uses these access attributes with texture object arguments.  
void foo (texture2d imgA [[texture(0)]],  
     texture2d imgB [[texture(1)]], 
texture2d imgC [[texture(2)]])  
{…} 
(See section 4.3.1 for a description of the 
texture
 attribute.) 
A texture type can also be used as the variable type for any variables declared inside a function. 
The 
access
 attribute for variables of texture type declared inside a function must be 
access::read
 or 
access:sample
. Declaring variables inside a function to be a texture type but 
which do not use 
access::read
 or 
access:sample
 attributes causes a compilation error. 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
26
174

---