Metal Shading Language Specification

threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);  


// process pixels in imageblock - update the elements in slice 
process_pixels_in_imageblock(img_blk, gid, lid); 


// a barrier to make sure all threads have finished writing to the 
elements in the imageblock 
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock); 


// write a specific element in imageblock to output image 
// only 1 thread in the threadgroup performs the imageblock write 
if (lid.x == 0 && lid.y == 0)  
dst.write(img_blk.slice(f->a), gid); 
}


 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page  
 of  
157
174

---

6  Compiler Options 
The Metal compiler can be used online (i.e. using the appropriate APIs to compile Metal 
sources) or offline. Metal sources compiled offline can be loaded as binaries, using the 
appropriate Metal APIs.  
This chapter explains the compiler options supported by the Metal compiler, which are 
categorized as pre-processor options, options for math intrinsics, options that control 
optimization and miscellaneous options. The online and offline Metal compiler support these 
options. 
6.1  Pre-Processor Compiler Options 
These options control the Metal preprocessor that is run on each program source before actual 
compilation.  
 
-D name 
Predefine name as a macro, with definition 1. 
 
-D name=definition 
The contents of definition are tokenized and processed as if they appeared in a #define 
directive. This option may receive multiple options, which are processed in the order in which 
they appear. This option allows developers to compile Metal code to change which features are 
enabled or disabled. 
 
-I dir 
Add the directory dir to the list of directories to be searched for header files. This option is only 
available for the offline compiler. 
6.2  Math Intrinsics Compiler Options 
These options control compiler behavior regarding floating-point arithmetic. These options 
trade off between speed and correctness. 
-ffast-math (default)


-fno-fast-math 
These options enable (default) or disable the optimizations for floating-point arithmetic that 
may violate the IEEE 754 standard. They also enable or disable the high precision variant of 
math functions for single precision floating-point scalar and vector types. 
The optimizations for floating-point arithmetic include: 
• No NaNs – Allow optimizations to assume the arguments and result are not NaN. 
• No Infs – Allow optimizations to assume the arguments and result are not positive or 
negative infinity. 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page  
 of  
158
174

---

• No Signed Zeroes – Allow optimizations to treat the sign of zero argument or result as 
insignificant. 
• Allow Reciprocal – Allow optimizations to use the reciprocal of an argument rather than 
perform division 
• Fast – Allow algebraically equivalent transformations i.e. re-associate floating-point 
operations that may dramatically change results in floating-point.   
6.3  Compiler Options Controlling the Language Version 
The following option controls the version of unified graphics/compute language that the 
compiler accepts. 
 
-std= 
Determine the language revision to use. A value for this option must be provided. This can only 
be: 
•
ios-metal1.0
 – support the unified graphics / compute language revision 1.0 programs 
for iOS 8.0. 
•
ios-metal1.1
 – support the unified graphics / compute language revision 1.1 programs 
for iOS 9.0.  
•
ios-metal1.2
 – support the unified graphics / compute language revision 1.2 programs 
for iOS 10.0.  
•
ios-metal2.0
 – support the unified graphics / compute language revision 2.0 programs 
for iOS 11.0.  
•
osx-metal1.1
 – support the unified graphics / compute language revision 1.1 programs 
for macOS 10.11.  
•
osx-metal1.2
 – support the unified graphics / compute language revision 1.2 programs 
for macOS 10.12.  
•
osx-metal2.0
 – support the unified graphics / compute language revision 2.0 programs 
for macOS 10.13.  
6.4  Compiler Options to Request or Suppress Warnings 
The following options are available. 
 
-Werror 
Make all warnings into errors. 
 
-W 
Inhibit all warning messages. 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page  
 of  
159
174

---

7  Numerical Compliance 
This chapter covers how Metal represents floating-point numbers with regard to accuracy in 
mathematical operations. Metal is compliant to a subset of the IEEE 754 standard.  
7.1  INF, NaN, and Denormalized Numbers 
INF must be supported for single-precision and half-precision floating-point numbers. 
NaNs must be supported for single-precision and half-precision floating-point numbers (with 
fast math disabled). If fast math is enabled the behavior of handling NaN or INF (as inputs or 
outputs) is undefined. Signaling NaNs are not supported.  
Denormalized single-precision or half-precision floating-point numbers passed as input to or 
produced as the output of single-precision or half-precision floating-point arithmetic operations 
may be flushed to zero. 
7.2  Rounding Mode 
Either round to nearest even or round to zero rounding mode may be supported for single 
precision and half precision floating-point operations. 
7.3  Floating-Point Exceptions 
Floating-point exceptions are disabled in Metal. 
7.4  Relative Error as ULPs 
Table 36 describes the minimum accuracy of single-precision floating-point basic arithmetic 
operations and math functions given as ULP values. The reference value used to compute the 
ULP value of an arithmetic operation is the infinitely precise result. 
Table 36 Minimum Accuracy of Single-Precision Floating-Point Operations and 
Functions 
Math Function
Min Accuracy - ULP values 
x + y
Correctly rounded
x - y
Correctly rounded
x * y
Correctly rounded
1.0 / x
Correctly rounded
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page  
 of  
160
174

---