【翻译】Metal 着色语言规范--数据类型

这份文档介绍了 Metal 统一图形和计算语言。Metal 是一种基于 C++ 的编程语言,开发者可以用它来编写在 GPU 上执行的图形和通用数据并行计算代码。这是第二章:数据类型。

August 24, 2017 -
ios metal 翻译

这份文档翻译自 Metal 的官方文档:Metal 着色语言规范(Metal Shading Language Specification)

Metal 着色语言规范

版本 2.0

2 数据类型

这一章是 Metal 数据类型的详细信息,包括表示矢量和矩阵的类型。还讨论了原子数据类型、缓冲区、纹理、取样器、数组和用户自定义结构体。也描述了类型对齐和类型转换。

2.1 标量数据类型

Metal 支持的标量类型列在了表 1。Metal 不支持 double、long、unsigned long、long long、unsigned long long 和 long double 数据类型。

表 1 Metal 标量数据类型

类型 描述
bool 一个有条件的数据类型,其值只能是 true 或 false。值 true 扩展为整型常量 1,值 false 扩展为整型常量 0。
char
int8_t
一个有符号的二进制补码 8 位整数
unsigned char
uchar
uint8_t
一个无符号 8 位整数
short
int16_t
一个有符号的二进制补码 16 位整数
unsigned short
ushort
unit16_t
一个无符号 16 位整数
int
int32_t
一个有符号二进制补码 32 位整数
unsigned int
unit
uint32_t
一个无符号 32 位整数
half 一个 16 位浮点数。half 数据类型必须符合 IEEE 754 binary16 存储格式
float 一个 32 位浮点数。float 数据类型必须符号 IEEE 754 单精度存储格式
size_t 一个无符号整数类型,是 sizeof 计算出的结果。这是一个 64 位无符号整数。
ptrdiff_t 一个有符号整数类型,是两个指针相减的结果。这是一个 64 位有符号整数。
void void 类型包含一组空值;这是一个还未完成的不完整的类型。

注意:Metal 支持标准的以 f 或 F 结尾指定单精度浮点文字值(如:0.5f 或 0.5F)。此外,Metal 还支持以 h 或 H 结尾指定半精度浮点文字值(如:0.5f 或 0.5F)。Metal 也支持以 u 或 U 结尾的无符号整数文字。

2.2 矢量和矩阵数据类型

Metal 支持由系统矢量数学库实现的矢量和矩阵数据类型的子集。

矢量类型名字支持:

booln,
charn, shortn, intn, ucharn, ushortn, uintn,
halfn  floatn

n 是 2、3 或 4 表示 2-、3- 或 4- 分量矢量类型。

矩阵类型名字支持:

halfnxm  floatnxm

n 和 m 是行和列的数字。n 和 m 可以是 2、3 或 4。矩阵类型 floatnxm 由 n 个 floatm 矢量组成。同样的,矩阵类型 halfnxm 由 n 个 halfm 矢量组成。

2.2.1 访问矢量分量

可以用数组索引访问矢量分量。数组索引 0 指向矢量的第一个分量,索引 1 指向第二个分量,依此类推。下面的例子展示了访问数组分量的各个方法:

pos = float4(1.0f, 2.0f, 3.0f, 4.0f);

float x = pos[0]; // x = 1.0
float z = pos[2]; // z = 3.0

float4 vA = float4(1.0f, 2.0f, 3.0f, 4.0f);
float4 vB;

for (int i=0; i<4; i++)
      vB[i] = vA[i] * 2.0f // vB = (2.0, 4.0, 6.0, 8.0);

Metal 支持使用点(.)作为选择符号来访问矢量分量,使用可能指示坐标或颜色的字母:

<vector_data_type>.xyzw 
<vector_data_type>.rgba

在接下来的代码里,矢量 test 被初始化了,接着使用 .xyzw 或 .rgba 选择语法来访问分量:

int4 test = int4(0, 1, 2, 3);
int a = test.x;  // a = 0
int b = test.y;  // b = 1
int c = test.z;  // c = 2
int d = test.w;  // d = 3
int e = test.r;  // e = 0
int f = test.g;  // f = 1
int g = test.b;  // g = 2
int h = test.a;  // h = 3

分量选择语法允许选择多个分量。

float4 c;
c.xyzw = float4(1.0f, 2.0f, 3.0f, 4.0f);
c.z = 1.0f;
c.xy = float2(3.0f, 4.0f);
c.xyz = float3(3.0f, 4.0f, 5.0f);

分量选择语法也允许替换或复制分量。

float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
float4 swiz = pos.wzyx; // swiz = (4.0f, 3.0f, 2.0f, 1.0f)
float4 dup = pos.xxyy;  // dup = (1.0f, 1.0f, 2.0f, 2.0f)

分量组符号可以出现在表达式的左侧。为了形成左值,可以应用交叉混合。作为结果的左值应该是标量或矢量类型,依赖于分量指定的数字。每个分量都必须支持标量或矢量类型。作为结果的左值绝对不能包含重复的分量。

float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
// pos = (5.0, 2.0, 3.0, 6.0)
pos.xw = float2(5.0f, 6.0f);

// pos = (8.0, 2.0, 3.0, 7.0)
pos.wx = float2(7.0f, 8.0f);

// pos = (3.0, 5.0, 9.0, 7.0)
pos.xyz = float3(3.0f, 5.0f, 9.0f);

以下访问矢量分量的方法不合法,并会导致编译期错误的结果:

  • 访问超出了为矢量类型声明的分量是一个错误。2 个分量矢量数据类型只能访问 .xy 或 .rg 元素。3 个分量矢量数据类型只能访问 .xyz 或 .rgb 元素。例如:
float2 pos;
pos.x = 1.0f; // 合法; y 也一样
pos.z = 1.0f; // 不合法;w 也一样

float3 pos;
pos.z = 1.0f; // 合法
pos.w = 1.0f; // 不合法
  • 在左侧两次访问同样分量是含混的;例如,
// 不合法 - 'x' 使用了两次
pos.xx = float2(3.0f, 4.0f);

// 不合法 - 没有匹配 float2 和 float4 之间
pos.xy = float4(1.0f, 2.0f, 3.0f, 4.0f);
  • .rgba 和 .xyzw 限定符不能在单次访问里混合使用;例如,
float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
pos.x = 1.0f;
pos.g = 2.0f;
pos.xg = float2(3.0f, 4.0f); // 不合法 - 混合使用限定符
float3 coord = pos.ryz;  // 不合法 - 混合使用限定符
  • 对矢量交叉使用指针和引用;例如:
float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
my_func(&pos.xy); // 不合法

矢量类型上的 sizeof 运算符返回矢量的大小,就是给定的分量数量 * 每个分量的大小。例如,sizeof(float4) 返回 16 和 sizeof(half4) 返回 8。

2.2.2 访问矩阵分量

floatnxm 和 halfnxm 矩阵可以被当作 n 个 floatm 数组或 n 个 halfm 数组来访问。

矩阵的分量可以用数组下标语法来访问。应用单个下标到矩阵,是将矩阵视作列矢量的数组。最上面一列是 0 列。第二个下标将对所得到的矢量进行操作,如前面定义的矢量。于是,两个下标选择一列和一行。

float4x4 m;

// sets the 2nd column to all 2.0
m[1] = float4(2.0f);

// sets the 1st element of the 1st column to 1.0
m[0][0] = 1.0f;

// sets the 4th elements of the 3rd column to 3.0
m[2][3] = 3.0f;

使用非常量表达式访问矩阵范围之外的分量会导致未定义的行为。使用常量表达式访问矩阵范围之外的分量会产生编译期错误。

2.2.3 矢量构造函数

构造函数可以用来从一组标量或矢量里创造矢量。当一个矢量被初始化时,它的参数签名确定了它的构造方式。例如,如果矢量只用单个标量参数来初始化,构造矢量的全部分量都会被设置成标量值。

如果矢量的构造来自于多个标量,一个或多个矢量,或是这些的合成,那么矢量的分量会按参数的组成顺序来构造。参数从左到右被消耗。在下一个参数的任何分量被消耗之前,按顺序,每个参数都消耗其全部分量。

这是一个可用于 float4 的构造函数的完整列表:

float4(float x);
float4(float x, float y, float z, float w);
float4(float2 a, float2 b);
float4(float2 a, float b, float c);
float4(float a, float b, float c);
float4(float a, float2 b, float c);
float4(float3 a, float b);
float4(float a, float3 b);
float4(float4 x);

这是一个可用于 float3 的构造函数的完整列表:

float3(float x);
float3(float x, float y, float z);
float3(float a, float2 b);
float3(float2 a, float b);
float3(float3 x);

这是一个可用于 float2 的构造函数的完整列表:

float2(float x);
float2(floatx, float y);
float2(float2 x);

下面的例子演示了如何使用构造函数:

float x = 1.0f, y = 2.0f, z = 3.0f, w = 4.0f;
float4 a = float4(0.0f);
float4 b = float4(x, y, z, w);
float2 c = float2(5.0f, 6.0f);

float2 a = float2(x, y);
float2 b = float2(z, w);
float4 x = float4(a.xy, b.xy);

在矢量构造函数里使用初始化是一个编译期错误。

2.2.4 矩阵构造函数

构造函数可用来从一组标量、矢量或矩阵里创建矩阵。当一个矩阵被初始化之后,它的参数签名确定了它是如何被构造的。例如,如果仅使用单个标量参数来初始化矩阵,结果就是矩阵对角线上的全部分量都是这个标量,其余的分量都被初始化为 0.0。例如,一次调用:

float4x4(fval);

其中 fval 是一个标量浮点值,使用以下初始内容来构造一个矩阵:

fval  0.0   0.0   0.0
0.0   fval  0.0   0.0
0.0   0.0   fval  0.0
0.0   0.0   0.0   fval

也可以使用一个矩阵来构造另一个矩阵,它们具有同样的大小,即有同样数量的行和列。例如,

float3x4(float3x4);
float3x4(half3x4);

矩阵的分量按列的主要顺序被构造和消耗。在构造的矩阵对象里,矩阵的构造函数必须有刚好足够的值来指定它的参数去初始化每个分量。提供多余所需的参数是一个错误。在矩阵构造函数里使用初始化也会造成编译期错误。

一个有着 n 列和 m 行的矩阵类型 T,也可以使用有 n 个矢量和 m 个分量的类型 T 来构造。下面的例子是合法的构造函数:

float2x2(float2, float2);
float3x3(float3, float3, float3);
float3x2(float2, float2, float2);

一个有着 n 列和 m 行的矩阵类型 T,也可以使用有 n * m 个标量的类型 T 来构造。下面的例子是合法的构造函数:

float2x2(float, float, float, float);
float3x2(float, float, float, float, float, float);

以下矩阵构造函数的例子是不被支持的。不能混合使用矢量和标量来构造矩阵。

// 不支持
// float2x3(float2 a, float b, float2 c, float d);

2.3 原子数据类型

Metal 原子数据类型仅限于在 Metal 编程语言在执行原子函数时使用,在 5.12 节会解释。这些原子函数是 C++14 原子和同步函数的子集。Metal 原子函数必须在 Metal 原子数据上运算。

Metal 原子类型的定义如下:

atomic<T>T 必须是 intuint  bool
atomic_intatomic_uint  atomic_bool

2.4 缓冲区

Metal 将缓冲区实现为一个指向内部的指针,或在设备、线程组或常量地址空间里描述地用户自定义数据类型。(参考 4.2 节对这些地址限定符的完整说明。)。这些缓冲区可以在程序范围内声明,也可以作为参数传递给函数。

例子:

device float4 *device_buffer;

struct my_user_data {
    float4 a;
    float b;
    int2 c;
};

一般的 Metal 缓冲区可能包含:

  • 基本类型如 float 和 int2
  • 矢量和矩阵类型
  • 缓冲区类型的数组
  • 缓冲区类型的结构体
  • 缓冲区类型的统一值 对于参数缓冲区,见 2.8 节。

2.5 纹理

纹理数据类型是处理对应于全部或部分纹理上的单个纹理映射层的一维、二维或三维纹理数据。下面的模版定义了具体的纹理数据类型:

enum class access { sample, read, write, read_write };

texture1d<T, access a = access::sample>
texture1d_array<T, access a = access::sample>
texture2d<T, access a = access::sample>
texture2d_array<T, access a = access::sample>
texture3d<T, access a = access::sample>
texturecube<T, access a = access::sample>
texturecube_array<T, access a = acess::sample>
texture2d_ms<T, access a = access::read>

深度格式(depth format)的纹理必须声明为以下纹理数据类型之一:

depth2d<T, access a = access::sample>
depth2d_array<T, access a = access::sample>
depthcube<T, access a = access::sample>
depthcube_array<T, access a = access::sample>
depth2d_ms<T, access a = access::read>

T 指定从纹理读取时返回的颜色类型,或者写入到纹理时指定的颜色类型。对于纹理类型(除了深度纹理类型),T 可以是 half、float、short、ushort、int 或 uint。对于深度纹理类型,T 必须是 float。

注意:如果 T 是 uint 或 ushort,纹理相关联的数据必须使用无符号整数格式。如果 T 是 half,纹理相关联的数据必须是普通的(有符号或无符号整数)或者半精度格式。如果 T 是 float,纹理相关联的数据必须是普通的(有符号或无符号整数),或者半精度或单进度格式。

访问限定符解释了纹理可以被访问的方式。支持的访问限定符是:

  • sample - 纹理对象可以被取样。sample 意味着使用或不使用取样器时从一个纹理里读取的能力。
  • read - 没有取样器时,图形或内核函数只能读取纹理对象。
  • write - 图形或内核函数可以写入纹理对象。
  • read_write - 图形或内核对象可以读取和写入纹理对象。 注意:
  • 对于多重取样纹理,只支持 read 限定符。
  • 对于深度纹理,只支持 sample 和 read 限定符。

下面的例子是在纹理对象的参数上使用这些访问限定符。

void foo (texture2d<float> imgA [[ texture(0) ]],
    texture2d<float, access::read> imgB [[ texture(1) ]],
    texture2d<float, access::write> imgC [[ texture(2) ]])
{
    ...
}

(纹理属性限定符的详情请见 4.3.1 节。)

纹理类型也可以用作在函数里声明的任何变量的变量类型。函数里声明的作为变量的纹理类型,使用的访问限定符,只能是 access::read 或 access::sample。在函数里声明作为变量的纹理类型,没有使用 access::read 或 access::sample 限定符,会导致编译错误。 例如:

void foo (texture2d<float> imgA [[ texture(0) ]],
    texture2d<float, access::read> imgB [[ texture(1) ]],
    texture2d<float, access::write> imgC [[ texture(2) ]])
{
    texture2d<float> x = imgA; <- 合法
    texture2d<float, access::read> y = imgB; <- 合法
    texture2d<float, access::write> z; < - 不合法
    ...
}

2.6 取样器

取样器类型识别如何取样一个纹理。Metal API 允许你创建一个取样器对象,并作为图形或核心函数的参数传递它。作为在 API 里描述的替代,一个取样器对象也可以在程序源码里进行描述。为了这些情况,我们只允许指定取样器状态的子集:寻址模式,滤波模式,归一化坐标和比较功能。

表 2 列出了支持的取样器状态枚举和它们的相关值(及默认值)。当一个取样器在 Metal 程序源码里被初始化时,可以指定这些状态。

表 2 取样器状态枚举值

枚举名称 有效值 描述
coord normalized (default)
pixel
指定从纹理采样时的纹理坐标是归一化的还是非标准的值。
address repeat
mirrored_repeat
clamp_to_edge (default)
clamp_to_zero
clamp_to_border
设置全部纹理坐标的寻址模式
s_address repeat
mirrored_repeat
clamp_to_edge (default)
clamp_to_zero
clamp_to_border
设置个别纹理坐标的寻址模式
filter nearest (default)
linear
设置纹理取样器的放大和缩小滤波的模式
mag_filter nearest (default)
linear
设置纹理取样器的放大滤波的模式
min_filter nearest (default)
linear
设置纹理取样器的缩小滤波的模式
mip_filter none (default)
nearest
linear
设置纹理取样器的纹理映射滤波模式。如果是 none,则只有一个细节层处于活动状态。
compare_func never (default)
less
less_equal
greater
greater_equal
equal
not_equal
always
设置由 sample_compare 和 gather_compare 纹理函数使用的比较测试

clamp_to_border 只支持 macOS。使用 clamp_to_border,在纹理之外进行取样时将只使用纹理坐标的边框颜色(并且不在纹理边缘使用任何颜色)。

clamp_to_zero 相当于 alpha 组件值来自于纹理,边框颜色为 transparent_black(0.0, 0.0, 0.0) 的 clamp_to_border。如果一个或多个纹理坐标的地址模式被指定为 clamp_to_zero,当且仅当边框颜色为 transparent_black 时,其它纹理坐标可以使用 clamp_to_border 地址模式。否则其行为是未定义的。

表 2 里说明的取样器数据类型使用的枚举类型的定义如下:

enum class coord<sup>1</sup>  { normalized, pixel };

enum class filter             { nearest, linear };
enum class min_filter         { nearest, linear };
enum class mag_filter         { nearest, linear };
enum class mip_filter         { none, nearest, linear };

enum class s_address          { clamp_to_zero, clamp_to_edge, clamp_to_border
                                repeat, mirrored_repeat };
enum class t_address          { clamp_to_zero, clamp_to_edge, clamp_to_border
                                repeat, mirrored_repeat };
enum class r_address          { clamp_to_zero, clamp_to_edge, clamp_to_border
                                repeat, mirrored_repeat };
enum class address            { clamp_to_zero, clamp_to_edge, clamp_to_border
                                repeat, mirrored_repeat };

enum class compare_func       { none, less, less_equal, greater, greater_equal, equal, not_equal };

enum class border_color<sup>2</sup> { transparent_black, opaque_black, opaque_white };

1 如果 coord 设置为 pixel,那么 min_filter 和 mag_filter 值必须相同,mip_filter 值必须是 none,并且 address 模式只能是 clamp_to_zero、clamp_to_border 或 clamp_to_edge。 2 当 address 模式是 clamp_to_border 时,border_color 是唯一有效的。

除了枚举类型,取样器也可以指定为以下类型:

max_anisotropy(int value)
lod_clamp(float min, float max)

Metal 执行的取样器对象如下:

struct sampler {
    public:
        // full version of sampler constructor
        template<typename... Ts>
        constexpr sampler(Ts... sampler_params){};
    private:
};

Ts 必须是上面列出的取样器数据类型可以使用的枚举类型。当在一个给定的取样器构造函数里多次声明同样的枚举类型时,最后列出的值生效。

以下 Metal 程序源码演示了几种声明取样器的方式。(出现在代码中的属性限定符(sampler(n),buffer(n) 和 texture(n))会在 4.3.1 节解释)。注意,程序源码里的取样器或常量缓冲区的声明不需要这些属性限定符。

constexpr sampler s(coord::pixel, address::clamp_to_zero, filter::linear);

constexpr sampler a(coord::normalized);

constexpr sampler b(address::repeat);

constexpr sampler s(address::clamp_to_zero, filter::linear, compare_func::less);

constexpr sampler s(address::clamp_to_zero, filter::linear, compare_func::less, max_anisotropy(10), lod_clamp(0.0f, MAXFLOAT));

注意:在 Metal 着色语言源码里初始化取样器时,必须使用 constexpr 来声明。

2.7 数组和结构体

数组和结构体只受以下有限的支持:

  • 图形和内核函数的参数不能被声明为 size_t、ptrdiff_t 或包含声明为这些内置标量类型之一成员的结构体和/或联合体。
  • 结构体的成员必须属于同样的地址空间。

在介绍 Metal 着色语言 2.0 版的参数缓冲区(见 2.8 节)之前,以下限制依然适用:

  • 如果纹理和取样器类型作为值传递给图形和内核函数,那么纹理和取样器类型只能在一个结构体里声明。(参数缓冲区支持引用纹理和取样器类型。)
  • 用来声明纹理的数组(见 2.7.1 节)的 array<T, N> 类型,不能在结构体中声明。(参数缓冲区支持声明纹理数组。)

#### 2.7.1 纹理数组和取样器数组

一个纹理数组的声明方式如:

    array<typename T, size_t N> 
    const array<typename T, size_t N>

T 必须是 2.5 节里描述的,使用 access::read 或 access::sample 限定符声明的纹理类型。

一个取样器数组的声明方式如:

    array<sampler, size_t N> 
    const array<sampler, size_t N>

纹理数组或取样器数组可以作为参数传递给函数(图形、内核或用户自定义函数),或被声明为函数内部的本地变量。取样器数组也可以在程序范围里声明。

Metal 着色语言还添加了对 array_ref 的支持。array_ref 表示一个 size() 个元素类型为 T 的不可变数组。T 必须是取样器类型或受支持的纹理类型。数组的存储 不是由 array_ref 对象所有。隐式转换操作是由具有连续迭代器(如 metal::array)的类型提供。array_ref 的常见用法是将纹理数组作为参数传递给需要接受各种数组类型的函数。array_ref 类型不能作为参数传递给图形和内核函数。不论怎样,array_ref 类型都可以作为参数传递给用户自定义函数。array_ref 类型不可以在函数内部声明为本地变量。

2.7.1.1 到 2.7.1.3 节列出的成员函数可用于纹理数组、取样器数组和 array_ref 类型:

2.7.1.1 元素访问

纹理数组或取样器数组的元素可以使用 [] 操作符来访问。以下 [] 操作符的变种都可使用:

  reference operator[] (size_t pos) const;
  constexpr const_reference operator[] (size_t pos) const;

模版类型 array_ref 的元素也可以使用以下 [] 操作符的变种来访问:

    constexpr const_reference operator[] (size_t pos) const;
2.7.1.2 容量
    constexpr size_t size();
    constexpr size_t size() const;

返回数组或 array_ref 里元素的数量。 例如:

    # include <metal_stdlib>
    using namespace metal;
    kernel void
    my_kernel(
      const array<texture2d<float>, 10> src [[ texture(0) ]],
      texture2d<float, access::write> dst [[ texture(10) ]],
      ... )
      {
          for (int i=0; i<src.size(); i++)
          {
            if (is_null_texture(src[i]))
                break;
            process_image(src[i], dst);
          }
      }
2.7.1.3 模版数组的构造函数
    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<size_t N>
    constexpr array_ref(const T(&a)[N]);

    template<typename T>
    constexpr array_ref<T> make_array_ref(const T * array, size_t length)

    template<typename T, size_t N>
    constexpr array_ref<T> make_array_ref(const T(&a)[N])

例如:

    #include <metal_stdlib>
    using namespace metal;

    float4
    foo(array_ref<texture2d<float>> src)
    {
        float4 clr(0.0f);
        for (int i=0; i<src.size; i++)
        {
            clr += process_texture(src[i]);
        }
        return clr;
    }

    kernel void
    my_kernel_A(
      const array<texture2d<float>, 10> src [[ texture(0) ]],
      texture2d<float, access::write> dst [[ texture(10) ]],
      ... )
      {
          float4 clr = foo(src);
          ...
      }

      kernel void
      my_kernel_B(
        const array<texture2d<float>, 20> src [[ texture(0) ]],
        texture2d<float, access::write dst [[ texture(10) ]],
        ... )
        {
          float4 clr = foo(src);
          ...
        }

下面是一个在程序范围里声明一个纹理数组的例子:

    #include <metal_stdlib>
    using namespace metal;
    constexpr array<sampler, 2> = { sampler(address::clamp_to_zero), sampler(coord::pixel) };

2.8 参数缓冲区

参数缓冲区扩冲了基本缓冲区,以包含指针(缓冲区)、纹理和取样器。但是,参数缓冲区不含有联合体。下面的例子演示了在着色器中指定一个叫做 Foo 的参数缓冲区结构体:

    struct Foo {
        texture2d<float, access::write> a; depth2d<float> b;
        sampler sam c;
        texture2d<float> d;
        device float4* e;
        texture2d<float> f;
        int g;
    };
    kernel void my_kernel(constant Foo & f [[buffer(0)]])
    {
        ...
    }

可用现存的 array<T, N> 模版类型来声明纹理数组和取样器数组。所有其它合法的缓冲区类型的数组应该也可以用 C 风格的数组语法来声明。

参数缓冲区的成员可以分配一个通用 [[id(n)]] 属性,n 是 32 位无符号整数,可用于识别 Metal API 中的缓冲区元素。参数缓冲区可以区别于常规缓冲区,如果它们含有缓冲区、纹理、取样器或任何 [[id]] 属性的元素。

在参数缓冲区里同样的索引不能被分配给多个成员。手动分配的索引不需要是连续的,但它们必须单调递增。在下面的例子里,索引 0 自动分配给 foo1。[[id(n)]] 属性限定符指定了结构体成员 t1 和 t2 的索引偏移量。由于没有为 foo2 指定索引,它自动分配了下一个索引,4,取决于前一个结构体成员使用的最大 ID 加 1。

    struct Foo {
        texture2d<float> t1 [[id(1)]];
        texture2d<float> 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
    };

如果省略了 [[id]] 属性限定符,ID 的自动分配依据以下规则:

  1. ID 按顺序分配给结构体成员,前一个结构体成员使用的最大 ID 加 1。在下面的例子里,没有提供索引,所以自动分配索引 0 和 1。
      struct MaterialTexture {
     texture2d<float> tex; // 分配索引 0
     float4 uvScaleOffset; // 分配索引 1
      };
    
  2. ID 按顺序分配给数组的元素,前一个数组元素使用的最大 ID 加 1。在下面的例子里,索引 1-3 是自动分配给了 texs1 的三个数组元素。索引 4-5 自动分配给 materials[0] 里的成员,索引 6-7 给 materials[1],索引 8-9 给 materials[2]。[[id(20)]] 属性限定符首先将索引 20 分配给常量。
      struct Material {
     float4 diffuse;                     // 分配索引 0
     array<texture2d<float>, 3> texs1;   // 分配索引 1-3
     MaterialTexture materials[3];       // 分配索引 4-9
     int constants [[id(20)]] [4];       // 分配索引 20-23
      };
    
  3. 如果一个结构体成员或数组元素 E 自身是结构体或数组,它的结构体成员或数组被分配的索引依据规则 1 和 2 的递归,从分配给 E 的 ID 开始。在下面的例子里,索引 4 被明确的提供给称为 normal 的嵌套结构体,所以它的元素(前面定义的 tex 和 uvScaleOffset)分配的 ID 是 4 和 5,个别的。称为 specular 的嵌套结构体的元素分配的 ID 是 6 和 7,由前一个成员使用的最大 ID(5)加 1。
      struct Material {
     MaterialTexture diffuse;            // 分配的索引 0,1
     MaterialTexture normal [[id(4)]];   // 分配的索引 4,5
     MaterialTexture specular;           // 分配的索引 6,7
      }
    
  4. 分配顶级参数缓冲区的参数的索引从 0 开始,满足规则 1-3。

在 macOS 上的二级硬件上,可以通过指针索引访问参数缓冲区。如下显示的这段语法,指的是连续的、独立编码参数缓冲区的数组:

    kernel void kern(constant Resources *resArray [[buffer(0)]])
    {
        constant Resources & resources = resArray[3];
    }
    kernel void kern(constant texture2d<float> *textures [[buffer(0)]]);

在支持的硬件上,为了支持 GPU 驱动的管线和间接绘制调用和调度,可以在着色器中的结构体和数组之间拷贝资源,如下显示:

    kernel void copy(constant Foo & src [[buffer(0)]], device Foo & dst [[buffer(1)]])
    {
        dst.a = src.d;
        ...
    }

取样器可能不会被从线程地址空间拷贝到设备地址空间。结果是,取样器可能只是直接从一个参数缓冲区拷贝到另一个参数缓冲区。下面显示了合法和不合法的拷贝例子:

    struct Resources {
        sampler sam;
    };

    kernel void copy(device Resources *src, device Resources *dst, sampler sam1)
    {
        constexpr sampler sam2;
        dst->sam = src->sam;    // 合法:设备 -> 设备
        dst->sam = sam1;        // 不合法:线程 -> 设备
        dst->sam = sam2;        // 不合法:线程 -> 设备
    }

在 macOS 上的二级硬件上,参数缓冲区可能含有指向其它参数缓冲区的指针:

    struct Textures {
        texture2d<float> diffuse;
        texture2d<float> specular;
    };
    struct Material {
        device Textures *textures;
    };
    fragment float4 fragFunc(device Material & material);

2.9 类型的大小和对齐

表 3 列出了标量和矢量数据类型的大小和对齐

表 3 标量和矢量数据类型的大小和对齐

类型 大小(byte) 对齐(byte)
bool 1 1
char,
uchar
1 1
short,
ushort
2 2
int,
uint
4 4
half 2 2
float 4 4
bool2 2 2
bool3 4 4
bool4 4 4
char2,
uchar2
2 2
char3,
uchar3
4 4
char4,
uchar4
4 4
short2,
ushort2
4 4
short3,
ushort3
8 8
short4,
ushort4
8 8
int2,
uint2
8 8
int3,
uint3
16 16
int4,
uint4
16 16
half2 4 4
half3 8 8
half4 8 8
float2 8 8
float3 16 16
float4 16 16

表 4 列出了矩阵数据类型的大小和对齐

表 4 矩阵数据类型的大小和对齐

类型 大小(byte) 对齐(byte)
half2x2 8 4
half2x3 16 8
half2x4 16 8
half3x2 12 4
half3x3 24 8
half3x4 24 8
half4x2 16 4
half4x3 32 8
half4x4 32 8
float2x2 16 8
float2x3 32 16
float2x4 32 16
float3x2 24 8
float3x3 48 16
float3x4 48 16
float4x2 32 8
float4x3 64 16
float4x4 64 16

可以使用对齐限定符 alignas 来指定类型或对象的对齐要求。alignas 限定符可以应用在结构体或类的变量或数据成员的声明上。它也可以应用在声明结构体、类或枚举类型。

Metal 编译器负责根据数据类型的要求,将数据项以合适的方式对齐。对于参数被声明为指针的图形或内核函数数据类型,Metal 编译器可以假定指针对象总是按照数据类型的要求进行适当对齐。

2.10 压缩矢量数据类型

在 2.2 节里详细解释了矢量数据类型是按矢量的大小来对齐的。这里有几个使用案例,开发者需要他们的矢量数据紧紧压缩。例如 - 一个顶点结构体可能包含位置、法线、切线向量和纹理坐标,紧紧的压缩并作为一个缓冲区传递给顶点函数。

支持的压缩矢量类型名字是:

    packed_charnpacked_shortnpacked_intn
    packed_ucharnpacked_ushortnpacked_uintn
    packed_halfn  packed_floatn

n 是 2、3 或 4 表示一个 2-、3- 或 4- 分量矢量类型。

表 5 列出了压缩矢量数据类型的大小和对齐方式。

表 5 压缩矢量数据类型的大小和对齐方式

类型 大小(byte) 对齐方式(byte)
packed_char2,
packed_uchar2
2 1
packed_char3,
packed_uchar3
3 1
packed_char4,
packed_uchar4
4 1
packed_short2,
packed_ushort2
4 2
packed_short3,
packed_ushort3
6 2
packed_short4,
packed_ushort4
8 2
packed_int2,
packed_uint2
8 4
packed_int3,
packed_uint3
12 4
packed_int4,
packed_uint4
16 4
packed_half2 4 2
packed_half3 6 2
packed_half4 8 2
packed_float2 8 4
packed_float3 12 4
packed_float4 16 4

3 保留 packed_booln 矢量类型名称

压缩矢量数据类型通常用作数据存储格式。双向支持在压缩矢量数据类型和对齐矢量数据类型之间加载和存储、复制结构函数和赋值运算。压缩矢量数据类型也支持算术、逻辑和关系运算符。

例如:

    device float4 *buffer;
    device packed_float4 *packed_buffer;

    int i;
    packed_float4 f ( buffer[i] );
    pack_buffer[i] = buffer[i];

    // 从 packed_float4 到 float4 的转换运算。
    buffer[i] = float4( packed_buffer[i] );

压缩矢量数据类型的分量可以使用数组索引来访问。但是,压缩矢量数据类型的分量不可以使用 .xyzw 或 .rgba 选择语法来访问。

例如:

    packed_float4 f;
    f[0] = 1.0f;    // OK
    f.x = 1.0f;     // 不合法 - 编译错误

2.11 隐式类型转换

支持在内置标量类型(除了 void)之间进行隐式转换。当隐式转换完成时,它不仅仅只是重新解释了表达式的值,而是将该值转换为新类型中的等效值。例如,整数值 5 转换为浮点值 5.0。

所有矢量类型被认为具有比标量类型更高的转换等级。不允许从一个矢量类型隐式转换成另一个矢量或标量类型,将导致编译错误。例如,下面尝试从 4 分量整数矢量转换到 4 分量浮点矢量,失败了。

    int4 i;
    float4 f = i;   // 编译错误

不支持从标量到矩阵类型和矢量到矩阵类型的隐式转换,将导致编译错误。不允许从一个矩阵类型到另一个矩阵、矢量或标量类型的隐式转换,将导致编译错误。

指针类型的隐式转换遵照 C++14 规范里描述的规则。

2.12 统一值类型

2.12.1 统一值类型的需要

在下面的着色器例子里,变量 i 是进入 texInput 给出的纹理数组的索引。变量 i 是非统一值;即,它可以在执行图形或内核函数的绘制或分发调用的线程上有不同的值,如下面的例子所示。因此,纹理取样硬件必须处理这种情况,一个示例请求可以引用用来执行图形或内核函数的绘制或分发调用的线程的不同纹理。

    kernel void
    foo(array<texture2d<float>, 10> texInput,
        array<texture2d<float>, 10> texOutput,
        sampler s,
        ...,
        uint2 gid [[ thread_position_in_grid ]])
    {
        int i = ...;
        float4 color = texList[i].sample(s, float2(coord.x, coord.y));
        ...;
        texOutput[i].write(color, coord);
    }

如果变量 i 在执行图形或内核函数的绘制或调用分发的全部线程上都有同样的值(即统一值),并且这个信息给传达给硬件,那么纹理取样硬件就可以应用适当优化。类似的参数可以用于纹理写入,其中在运行时计算的变量被用来作为纹理数组的索引,或一个或多个缓冲区的索引。

为了表明这个变量在执行图形或内核函数的绘制或分发调用的全部线程上都是统一值,Metal 着色语言加入了一个新的称为统一值(uniform)的模版类(在头 metal_uniform 里可用),可用于在图形或内核函数里声明变量。这个模版类只能使用算术类型(即 boolean、整型和浮点类型)和矢量类型来实例化。

下面的代码是之前例子的改进版,其中变量 i 被声明为统一值(uniform)类型:

    kernel void
    foo(array<texture2d<float>, 10> texInput,
        array<texture2d<float>, 10> texOutput,
        sampler s,
        ...,
        uint2 gid [[ thread_position_in_grid ]])
    {
        uniform<int> i = ...;
        float4 color = texList[i].sample(s, float2(coord.x, coord.y));
        ...;
        texOutput[i].write(color, coord);
    }

2.12.2 统一值类型的行为

如果一个变量是统一值类型,并且变量在执行内核或图形函数的全部线程上不是同样的值,那么其行为是未定义的。

统一值变量隐式类型转换到非统一值类型。将使用统一值变量的表达式计算的结果分配给统一值变量是合法的,但是分配非统一值变量给统一值变量会导致编译期错误。在下面的例子里,乘法运算符合法的将统一值变量 x 转换到了非统一值乘积 z。但是,分发非统一值变量 z 给统一值变量 b 导致编译期错误。

    uniform<int> x = ...;
    int y = ...;
    int z = x*y;         // x 被乘法转换成了非统一值
    uniform<int> b = z;  // 不合法;编译期错误

声明一个统一值元素的数组:

    uniform<float> bar[10]; // 数组 bar 里存储的元素是统一值

统一值类型可以合法的作为函数的参数和返回类型。例如:

    uniform<int> foo(...); // foo 返回一个统一值整数值
    int bar(uniform<int> a, ...);

将一个指针声明为统一值类型是合法的,但声明一个统一值指针是非法的。例如:

    device uniform<int> *ptr;   // 指向 ptr 的值是统一值
    uniform<device int *> ptr;  // 不合法;编译期错误

结合统一值与非统一值的表达式的值是非统一值。如果非统一值的结果被分配给一个统一值变量,如下面的例子,前端将产生一个编译期错误。如果编译期没有生成一个错误,后续的行为是未知。

    uniform<int> i = ...;
    int j = ...;
    if (i < j) {  // 表达式的结果是非统一值 (i < j)
        ...
        i++;      // 编译期错误,未定义的行为
    }

下面是类似的例子:

    bool p = ... // 非统一值条件判断。
    uniform<int> a = ..., b = ...;
    uniform<int> c = p ? a : b;   // 编译期错误,未定义的行为

2.12.4 统一值控制流

当控制流条件测试是基于统一值数量,全部程序实例在函数里的条件测试遵循同样的路径。基于统一值数量的控制流代码应该比基于非统一值数量的控制流代码更加高效。

2.13 类型转换和重新解析数据

static_cast4 运算符采用不饱和方式和默认舍入模式(即,当转换为浮点时,舍入到最近值;当转换到整数时,向零取整)将一个常量或矢量类型转换为其它常量或矢量类型。

Metal 加入了一个 as_type 运算符,允许任意常量或矢量数量类型重新解析为同样大小的其它常量或矢量数据类型。运算符里的位不修改为新类型,而是直接返回。不执行函数参数常用的类型提升。

例如,as_type(0x3f800000) 返回 1.0f,如果以 IEEE-754 单精度值来看,这个位模式 0x3f800000 的值。

使用 as_type 运算符将数据重新解析不同字节数的类型是一个错误。

例如:

    float f = 1.0f;
    // 合法。含有:0x3f800000
    uint u = as_type<uint>(f);
    // 合法。含有:
    // (int4)(0x3f800000, 0x40000000, 0x40400000, 0x40800000)
    float4 f = float4(1.0f, 2.0f, 3.0f, 4.0f);
    int4 i = as_type<int4>(f);

    int i;
    // 合法。
    short2 j = as_type<short2>(i);

    half4 f;
    // 错误。结果和运算有不同的大小
    float4 g = as_type<float4>(f);

    float4 f;
    // 合法。g.xyz 将与 f.xyz 有同样的值。
    // g.w 是未定义的
    float3 g = as_type<float3>(f);