[hlsl-out] Incorrect behaviour with firstTrailingBit

[hasali19]

This program incorrectly outputs -1 with DirectX. I think this is because of the use of firstbitlow, since it works correctly with tint which uses its own polyfill for firstbitlow - perhaps naga needs to do the same?

WGSL:

@group(0)
@binding(1)
var<storage, read_write> k: i32;

@compute
@workgroup_size(1)
fn main() {
    k = max(~0, firstTrailingBit(1));
}

HLSL:

RWByteAddressBuffer k : register(u1);

[numthreads(1, 1, 1)]
void main()
{
    k.Store(0, asuint(max(~0, firstbitlow(1))));
    return;
}

[evahop]

I currently don't have a method for testing hlsl, but this is perplexing. At first glance this would imply firstbitlow returns -1 for a bit at the 0 index. That seems very counterintuitive.

[teoxoy]

Looking at the generated DXBC via https://shader-playground.timjones.io/ it seems that the umax instruction is generated instead of imax. DXC seems to also be consistent here; generating DXIL with an evaluated -1.

It seems to me that the return type of firstbitlow is always an uint and doesn't depend on the argument contrary to what the HLSL docs suggest.

Wrapping the call with an int cast (i.e. int(firstbitlow(x))) gets FXC and DXC to output the expected instructions.

[teoxoy]

Looking at the src of DXC, firstbithigh might be in the same boat but I'm not sure what the meaning of the syntax in that file is.

[evahop]

I’m not sure of the syntax either but firstbithigh does appear to produce the same umax in the OP example. I will update this where needed.