`tf32` matrix multiplication is much less accurate than `fp16` matrix multiplication (in relative error)

[lengstrom]

`tf32` matrix multiplication is much less accurate than `fp16` matrix multiplication (in relative error)

[alexsamardzic]

I think it should be $0, $1 instead of $0, $0.

[ThomasRaoux]

Thank you for looking into this!

@ThomasRaoux , what instruction are you using for ASM? I tried cvt.rna.tf32.f32 $0, $0; on my A100 and it failed (produced all 0s).

Wrt truncation: do you know why tf32 with truncation would be less accurate than fp16? They should ideally be using basically the same number of input bits of precision, so I wouldn't expect an almost order-of-magnitude accuracy loss.

Sorry I missed to copy it... Here is what I ran that matched exactly torch:

@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, N: tl.constexpr, allow_tf32: tl.constexpr):
    ASM: tl.constexpr = "cvt.rna.tf32.f32 $0, $1;"
    offs = tl.arange(0, N)[:, None] * N + tl.arange(0, N)
    a = tl.load(a_ptr + offs)
    b = tl.load(b_ptr + offs)
    if allow_tf32:
        a = tl.inline_asm_elementwise(ASM, "=r, r", [a], dtype=tl.float32, is_pure=True, pack=1)
        b = tl.inline_asm_elementwise(ASM, "=r, r", [b], dtype=tl.float32, is_pure=True, pack=1)
    c = tl.dot(a, b, allow_tf32=allow_tf32)
    tl.store(c_ptr + offs, c.to(a.dtype))

[ThomasRaoux]

This has to do with the rounding applied to go from f32 to tf32. By default the tensorcore will truncate the f32. If you apply rounding before by doing something like that then you get the same precision as torch:

    if allow_tf32:
        a = tl.inline_asm_elementwise(ASM, "=r, r", [a], dtype=tl.float32, is_pure=True, pack=1)
        b = tl.inline_asm_elementwise(ASM, "=r, r", [b], dtype=tl.float32, is_pure=True, pack=1)

It's not something we want to do implicitly do in triton as it can affect performance significantly and allowing user to control this sounds better.
On Hopper applying such rounding would break the pipeline as tensorcore inputs come from shared memory.

[alexsamardzic]

This has to do with the rounding applied to go from f32 to tf32.

I'm looking into Triton TF32 precision issues from a related context, and just "discovered" the same: my experimenting involves CUTLASS too, that gives the same error as cuBLAS when tensor cores (i.e. TF32) used, so I was able to compare PTX generated from Triton and from CUTLASS, and it could be seen in the PTX that CUTLASS is applying rounding, while Triton is not.

[lengstrom]

Thank you for looking into this!

@ThomasRaoux , what instruction are you using for ASM? I tried cvt.rna.tf32.f32 $0, $0; on my A100 and it failed (produced all 0s).

Wrt truncation: do you know why tf32 with truncation would be less accurate than fp16? They should ideally be using basically the same number of input bits of precision, so I wouldn't expect an almost order-of-magnitude accuracy loss.