Металлические операции SIMD Min и Max терпят неудачу для поплавков - PullRequest
Новая
       79

Металлические операции SIMD Min и Max терпят неудачу для поплавков

0 голосов
/ 11 декабря 2018

Краткий вопрос

Почему я получаю неопределенное поведение от функций simd_min и simd_max в Metal 2.1 с плавающей запятой?

Обновление: похоже, только этопроисходит в графическом процессоре Radeon Pro 560X, но не в Intel UHD Graphics 630.

Фон

Согласно Руководство по языку металлического шейдинга Раздел 5.14Функции simd_min и simd_max поддерживаются для общих скалярных или векторных, целочисленных или с плавающей запятой типов.

Для simd_max спецификация гласит:

T simd_max(T data)

Возвращает максимальное значение в данных по всем активным потокам в SIMD-группе и передает результат всем активным потокам в SIMD-группе.

Test Case

Чтобы проверить это, я выполняю следующее тестовое ядро ​​с входным буфером из 128 случайных чисел с плавающей точкой в ​​диапазоне 0..100: 

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 МБ
  • Версия XCode 10.1 (10B61)
...