Metal Shading Language Specification

supported texture type. The storage for the array is not owned by the 
array_ref
 object. 
Implicit conversions are provided from types with contiguous iterators like 
metal::array
. A 
common use for 
array_ref
 is to pass an array of textures as an argument to functions so 
they can accept a variety of array types.  
The 
array_ref
 type cannot be passed as an argument to graphics and kernel functions. 
However, the 
array_ref
 type can be passed as an argument to user functions. The 
array_ref
 type cannot be declared as local variables inside functions. 
The member functions listed in sections 2.11.1.1 to 2.11.1.3 are available for the array of textures, 
array of samplers, and the 
array_ref
 types: 
2.11.1.1  
Element Access 
Elements of an array of textures or an array of samplers can be accessed using the 
[]
 operator. 
The following variants of the 
[]
 operator are available: 
reference operator[] (size_t pos) const; 
constexpr const_reference  operator[] (size_t pos) const; 
Elements of a templated type 
array_ref
 can also be accessed using the following variant 
of the 
[]
 operator: 
constexpr const_reference  operator[] (size_t pos) const; 
2.11.1.2  
Capacity 
constexpr size_t size(); 
constexpr size_t size() const;
 
Returns the number of elements in the array or the 
array_ref
. 
Examples: 
kernel void 
my_kernel(const array, 10> src [[texture(0)]], 
     texture2d dst [[texture(10)]], 
…) 
{ 
    for (int i=0; i    { 
        if (is_null_texture(src[i])) 
            break; 
        process_image(src[i], dst); 
    } 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
36
174

---

} 
2.11.1.3  
Constructors for Templated Arrays 
constexpr array_ref(); 
constexpr array_ref(const array_ref &); 
array_ref & operator=(const array_ref &); 
constexpr array_ref(const T * array, size_t length); 
template  
constexpr array_ref(const T(&a)[N]); 
template  
constexpr array_ref make_array_ref(const T * array, size_t length) 
template  
constexpr array_ref make_array_ref(const T(&a)[N]) 
Examples: 
float4 foo(array_ref> src) 
{ 
    float4 clr(0.0f); 
    for (int i=0; i    { 
        clr += process_texture(src[i]); 
    } 
    return clr;


} 
kernel void 
my_kernel_A(const array, 10> src [[texture(0)]], 
 texture2d dst [[texture(10)]], 
 …) 
{ 
    float4 clr = foo(src); 
    … 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
37
174

---

} 
kernel void 
my_kernel_B(const array, 20> src [[texture(0)]], 
 texture2d dst [[texture(10)]], 
  …) 
{ 
    float4 clr = foo(src); 
    … 
} 
Below is an example of an array of samplers declared in program scope:  
constexpr array =  { sampler(address::clamp_to_zero), 
 sampler(coord::pixel) };  
2.12 Argument Buffers 
Argument buffers extend the basic buffer types to include pointers (buffers), textures, and 
samplers. However, argument buffers cannot contain unions. The following example 
demonstrates an argument buffer structure called 
Foo
 that is specified in a function:  
struct Foo { 
texture2d a; depth2d b; 
sampler sam c; 
texture2d d; 
device float4* e; 
texture2d f; 
int g;  
};  
kernel void 
my_kernel(constant Foo & f [[buffer(0)]])  
{…}  
Arrays of textures and samplers can be declared using the existing 
array
 templated 
type. Arrays of all other legal buffer types can also be declared using C-style array syntax.  
Members of argument buffers can be assigned a generic 
[[id(n)]]
 attribute, where 
n
 is a 32-
bit unsigned integer that can be used to identify the buffer element from the Metal API. 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
38
174

---

Argument buffers can be distinguished from regular buffers if they contain buffers, textures, 
samplers, or any element with the 
[[id]]
 attribute.  
The same index may not be assigned to more than one member of an argument buffer. Manually 
assigned indices do not need to be contiguous, but they must be monotonically increasing. In 
the example below, index 0 is automatically assigned to 
foo1
. The 
[[id(n)]]
 attribute 
specifies the index offsets for the 
t1
 and 
t2
 struct members. Since no index is specified for 
foo2
, it is automatically assigned the next index, 4, which is determined by adding 1 to the 
maximum ID used by the previous struct member.  
struct Foo { 
texture2d t1 [[id(1)]]; 
texture2d t2 [[id(3)]];  
};  
struct Bar { 
Foo foo1; // foo1 assigned idx 0, t1 and t2 assigned idx 1 and 3 
Foo foo2; // foo2 assigned idx 4, t1 and t2 assigned idx 5 and 7  
}; 
If the 
[[id]]
 attribute is omitted, an ID is automatically assigned according to the following 
rules: 
1. IDs are assigned to struct members in order, by adding 1 to the maximum ID used by the 
previous struct member. In the example below, the indices are not provided, so indices 0 
and 1 are automatically assigned.  
struct MaterialTexture { 
texture2d tex; // Assigned index 0  
float4 uvScaleOffset; // Assigned index 1  
}; 
2. IDs are assigned to array elements in order, by adding 1 to the maximum ID used by the 
previous array element. In the example below, indices 1-3 are automatically assigned to 
the three array elements of 
texs1
. Indices 4-5 are automatically assigned to the fields in 
materials[0]
, indices 6-7 to 
materials[1]
, and indices 8-9 to 
materials[2]
. The 
[[id(20)]]
 attribute starts by assigning index 20 to constants.  
struct Material { 
float4 diffuse; 
 
 
 
// Assigned index 0 
array, 3> texs1; 
// Assigned indices 1-3  
MaterialTexture materials[3];   
// Assigned indices 4-9  
int constants [[id(20)]] [4];   
// Assigned indices 20-23 
};  
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page    of  
39
174

---