Floating-point `%` loses the sign of an exact zero remainder

[merlinsun]

Describe the bug

For floating-point operands, Triton's % lowers through the floating remainder operation. On the tested build, the exact-zero result loses the sign of the dividend: -0.0 % 3.0 is stored as +0.0 for both f32 and f64. Other finite remainder values in the same probe matched the reference.

Minimal f32 reproducer:

import torch
import triton
import triton.language as tl

@triton.jit
def kernel(x_ptr, y_ptr, out_ptr):
    i = tl.arange(0, 2)
    x = tl.load(x_ptr + i)
    y = tl.load(y_ptr + i)
    tl.store(out_ptr + i, x % y)

x = torch.tensor([-0.0, 0.0], device="cuda", dtype=torch.float32)
y = torch.tensor([3.0, 3.0], device="cuda", dtype=torch.float32)
out = torch.empty_like(x)
kernel[(1,)](x, y, out)
torch.cuda.synchronize()
print(out, out.view(torch.int32))

Expected first result: -0.0 (f32 bits 0x80000000). Observed: +0.0 (bits 0x00000000). The same mismatch was observed for f64. TRITON_INTERPRET=1 passes this isolated % test.

Environment details

Triton 4fd7cc5 on A800

[he-yufeng]

Confirmed the mechanism at the PTX level using Triton's own pinned LLVM build, no GPU needed. % on floats lowers arith.remf -> llvm.frem, and NVPTX has no frem instruction, so LLVM expands it inline:

div.rn.f32      // a / b
cvt.rzi.f32.f32 // trunc(...)
neg.f32
fma.rn.f32      // a - trunc(a/b) * b
testp.infinite.f32 + selp // inf divisor passes a through

For a = -0.0, b = 3.0: div gives -0.0, trunc keeps -0.0, and the fused multiply-add computes -0.0 + 0.0 = +0.0 per IEEE — the sign of the exact-zero remainder is lost right there. TRITON_INTERPRET passes because the interpreter uses Python semantics.

The fix I'm building: follow the llvm.frem with llvm.copysign(rem, dividend) in the elementwise conversion. fmod semantics already put the dividend's sign on every nonzero result, so the copysign only moves the zero case, and on backends with a real frem it is a no-op. Lit coverage for the conversion pattern included.