naga icon indicating copy to clipboard operation
naga copied to clipboard

[hlsl-out] Incorrect behaviour with firstTrailingBit

Open hasali19 opened this issue 2 years ago • 0 comments

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

hasali19 avatar May 23 '22 15:05 hasali19