Что такое сборка "set.eq.s32.b32"?

Question

Я пытаюсь прочитать фрагмент кода CUDA из этого репозитория github и наткнулся на это:

__device__ __forceinline__ bool isNegativeZero(float a) {
    int ret;
    asm volatile("{  set.eq.s32.b32 %0, %1, %2;}\n" : "=r"(ret) : "f"(a), "r"(0x80000000));
    return ret;
}

Я не уверен, что такое инструкции по сборке, но из контекста всего файла, функция, кажется, делает больше, чем просто проверяет, является ли число с плавающей запятой отрицательным нулем.

Я был бы очень признателен за объяснение функции на высоком уровне.

Answer 1

Я не уверен, что это за инструкция по сборке ...

Инструкции взяты из здесь .

set.eq.s32.b32 dest, valx, valy

означает «установить значение dest, если valx равно valy».Здесь valx - это входное значение для функции, valy - это 0x80000000, а dest - это возвращаемое значение функции ret.Остальная часть функции - это просто стандартный встроенный синтаксис, полученный из gcc, также задокументированный здесь .

Вы можете проверить здесь , что -0.0f равно 10000000000000000000000000000000или 0x80000000 в двоичном представлении IEEE 754.

Поэтому функция возвращает true, если a имеет отрицательное значение 0, проверяя на битовое равенство.

... но форму [sic] контекст всего файла ... кажется, что функция делает больше, чем просто проверяет, является ли число с плавающей точкой отрицательным нулем.

Надеюсь, теперь очевидно, что это не так.