Metal SIMD的最小和最大操作无法进行浮点运算

Question

简短的问题

为什么我从带有浮动的Metal 2.1中得到未定义的行为simd_min并simd_max在其中起作用？

更新：似乎仅在Radeon Pro 560X GPU上发生，而在Intel UHD Graphics 630上则没有。

背景

根据金属着色语言指南第5.14节，simd_min和simd_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个值将分别设置为第一个和最后一个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通道上的值仅被复制，其余的未定义。

相比之下，如果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）