浮点数的金属 SIMD 最小值和最大值操作失败

问：

问题简述

为什么我在 Metal 2.1 中使用浮点数从 Metal 2.1 中获得未定义的行为和函数？simd_minsimd_max

更新：似乎这只发生在 Radeon Pro 560X GPU 上，而不是在 Intel UHD Graphics 630 上。

背景

根据《金属着色语言指南》第 5.14 节，常见的标量或向量、整数或浮点类型都支持函数。simd_minsimd_max

对于 ，规范规定：simd_max

T simd_max(T data)

返回所有数据中的最大值 SIMD 组中的活动线程和 将结果广播到 SIMD组。

测试用例

为了测试这一点，我正在执行以下测试内核，其输入缓冲区为 0..100 范围内的 128 个随机浮点数：

kernel void simdMaxDebugKernel(
                          const device float *buffer [[ buffer(0) ]],
                          device float *output [[ buffer(1) ]],
                          uint id [[ thread_position_in_grid ]])
{
    output[id] = simd_max(buffer[id]);
}

通过检查，128 值的缓冲区分为两个 64 值的 SIMD 组。因此，我希望输出中的前 64 个值和后 64 个值分别设置为第一个和最后一个 SIMD 组的最大值。

测试结果

我得到了一些意想不到的结果：

inputs  [simd_float1]   128 values  
[0] Float   94.3006362
[1] Float   98.1107177
[2] Float   85.3725891
[3] Float   45.1457863
...
[63] Float  36.5486336
[64] Float  56.5494308
[65] Float  45.6249847
[66] Float  34.8077431

actual  [simd_float1]   128 values  
[0] Float   94.3006362
[1] Float   NaN
[2] Float   -3.80461845E+20
[3] Float   0.0000000000000000000000000000000000000212763294
...
[63] Float  0
[64] Float  56.5494308
[65] Float  -2467.3457
[66] Float  0.0000000000010178117
...

expectedMax simd_float1 99.4676971

在我看来，每个 SIMD 组的第一个 SIMD 通道的值只是被复制，其余的值是未定义的。

相比之下，如果按如下方式使用 to，则内核的行为与预期相同：uint

output[id] = (float)simd_max((uint)buffer[id]);

→

actual  [simd_float1]   128 values  
[0] Float   99
[1] Float   99
[2] Float   99
...
[63] Float  99
[64] Float  96
[65] Float  96
...

测试配置

• Mac OS 10.14.2 （18C54）
• MacBook Pro（15 英寸，2018 年机型）
• Radeon Pro 560X 4096 MB
• XCode 版本 10.1 （10B61）