what is "set.eq.s32.b32" assembly?

Question

I am trying to read a piece of CUDA code from this github repo and came across this:

__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;
}

I am not sure what the assembly instructions are, but form the context of the entire file, the function seems to do more than merely checking if the float is negative zero.

I would really appreciate a high-level explanation of the function.

Answer 1

I am not sure what the assembly instructions are...

The instructions come from here.

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

means "set the value of dest if valx is equal to valy". Here valx is the input value to the function and valy is 0x80000000, and dest is the return value of the function ret. The rest of the function is just standard gcc derived inline assembly syntax, also documented here.

You can check for yourself here that -0.0f is 10000000000000000000000000000000 or 0x80000000 in IEEE 754 binary32 representation.

Therefore the function returns true if a is negative 0, checking for bitwise equality.

... but form [sic] the context of the entire file..the function seems to do more than merely checking if the float is negative zero.

Hopefully it is now apparent that this is not the case.