Как предотвратить ЗСТ для одной строки в CUDA

Я работаю над кодом частиц, где обнуление широко используется для извлечения производительности. Однако есть одно выражение сравнения с плавающей запятой, которое я не хочу сбрасывать. Одним из решений является использование встроенного PTX, но оно вводит ненужные инструкции, поскольку в PTX нет логического типа, а только регистры предикатов:
Код C ++:

float a, b;
if ( a < b ) do_something;
// compiles into SASS:
//     FSETP.LT.FTZ.AND P0, PT, A, B, PT;
// @P0 DO_SOMETHING

PTX:

float a, b;
uint p;
asm("{.reg .pred p; setp.lt.f32 p, %1, %2; selp %0, 1, 0, p;}" : "=r"(p) : "f"(a), "f"(b) );
if (p) do_something;
// compiled into SASS:
//     FSETP.LT.AND P0, PT, A, B, PT;
//     SEL R2, RZ, 0x1, !P0;
//     ISETP.NE.AND P0, PT, R2, RZ, PT;
// @P0 DO_SOMETHING

Есть ли способ, которым я могу сделать сравнение без FTZ с одной инструкцией, не кодируя все это в PTX / SASS?

3

Решение

Задача ещё не решена.

Другие решения


По вопросам рекламы [email protected]