Изменение одной части кода влияет на другую часть при компиляции с флагом use_fast_math
У меня есть следующее ядро:
__global__ void kernel()
{
float loc{269.0f};
float s{(356.0f - loc) / 13.05f};
float a{pow(1.0f - 0.15f * s, 1.0f)};
float b{pow(1.0f - 0.15f * (356.0f - loc) / 13.05f, 1.0f)};
printf("%f\n", a);
printf("%f\n", b);
}
Он печатает
0.000000
0.000000
Но если я изменю способ расчета b:
__global__ void kernel()
{
float loc{269.0f};
float s{(356.0f - loc) / 13.05f};
float a{pow(1.0f - 0.15f * s, 1.0f)};
float b{pow(1.0f - 0.15f * ((356.0f - loc) / 13.05f), 1.0f)}; // notice the added braces
printf("%f\n", a);
printf("%f\n", b);
}
он печатает:
nan
nan
Почему добавление фигурных скобок для b также меняет a? Почему брекеты вообще дают эффект? Глядя на код на www.godbolt.org я вижу, что сгенерированная сборка отличается, но мне не хватает знаний, чтобы понять, что именно вызывает такое поведение.
Это моя конфигурация проекта:
set_target_properties(test PROPERTIES
CXX_STANDARD 14
CXX_STANDARD_REQUIRED YES
CXX_EXTENSIONS NO
CUDA_STANDARD 14
CUDA_STANDARD_REQUIRED YES
CUDA_EXTENSIONS NO
CUDA_SEPARABLE_COMPILATION ON
CUDA_ARCHITECTURES "61"
)
set(CUDA_FLAGS --use_fast_math)
set(CXX_FLAGS -O0)
Обратите внимание на флаг --use_fast_math - без него все работает нормально.
Мой графический процессор Quadro P1000.
Инструменты компиляции Cuda, выпуск 11.2, V11.2.152.
Ответы 1
Ваш код не численно стабильный: он содержит катастрофическая отмена. Действительно, 1.0f - 0.15f * s и 1.0f - 0.15f * ((356.0f - loc) / 13.05f) почти равны 0. Оно может быть равно нескольким единицам на последнем месте (ULP). Что касается округления, значение может быть положительным, отрицательным или равным 0. Когда базовое значение отрицательное, результат pow не определен.. Дело в том, что округление зависит от компилятора эвристический, поскольку вы явно включили быструю математику, которая отключает некоторые правила IEEE-754, такие как забота об ассоциативности с плавающей запятой и возможном наличии значений NaN. С быстрой математикой компилятор может использовать приближения как обратное (менее точное) вместо основного деления. На самом деле это действительно так в вашем случае. Следовательно, результаты не определены, и компилятор может установить вывод на 0 или NaN в зависимости от своей эвристики. Я думаю, что установление степени pow на 1.0f (что хорошо представлено) не избавит вас от такого поведения с pow (обратите внимание, что вы должны использовать powf для float значений, а не pow для double значений.
Правильным способом решения этой проблемы является использовать устойчивый численный расчет, который не содержит каких-либо катастрофических отмен, как в вашем коде. Фактическое решение очень зависит от общего вычислительного метода. Иногда нужно использовать совершенно другой алгоритм (например, QR-факторизация вместо LU-факторизации, двухэтапное вычисление дисперсии вместо одноэтапного).
Компиляторы используют оптимизацию под названием устранение общего подвыражения, чтобы оптимизировать математическое выражение, подобное вашему. Поскольку используется --use_fast_math, скобки не имеют значения (если они математически эквивалентны, но не численно). В результате изменение вычисления общего подвыражения в a может повлиять на b и наоборот. Здесь я ожидаю, что компилятор один раз вычислит все выражение.
Обратите внимание, что эвристика компилятора может вести себя по-разному в отношении несвязанных изменений. Это (к сожалению) довольно часто встречается в GCC и ICC. Это может вызвать неясное поведение, такое как добавление printf (в условии, которое никогда не истинно), что неожиданно ускоряет программу. Эта эвристика особенно сильно влияет на встраивание функций, выравнивание и компоновку кода.
Спасибо за подробное объяснение, оно объясняет, почему в результате есть
nanи как наличие фигурных скобок влияет на вычислениеb. Но я спросил еще об одном: если вы внимательно посмотрите на вычислениеa, вы увидите, что оно одинаково для обоих случаев, хотя его результат зависит от того, как было вычисленоb.