IR builder does not support FP8 dtypes

[danielvegamyhre]

Describe the bug

This was discovered while debugging a PyTorch bug report: pytorch/pytorch#141396 (comment)

The TL;DR is tl.full does not currently support FP8 dtypes, due to the IR builder lacking necessary methods for fp8 that the implementation of tl.full would depend on. Is this a known gap? If so, is there a timeline for supporting fp8 dtypes in the IR builder that can be shared?

The original error was from a call to tl.full in a triton kernel generated by torch inductor:

def triton_poi_fused_ones_1(out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 262144
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x0 = xindex
    tmp0 = tl.full([1], 1.0, tl.float8e4nv)
           ^

The error:

Traceback (most recent call last):
  File "/home/danvm/.conda/envs/pytorch/lib/python3.9/site-packages/triton/language/core.py", line 35, in wrapper
    return fn(*args, **kwargs)
  File "/home/danvm/.conda/envs/pytorch/lib/python3.9/site-packages/triton/language/core.py", line 1229, in full
    return semantic.full(shape, value, dtype, _builder)
  File "/home/danvm/.conda/envs/pytorch/lib/python3.9/site-packages/triton/language/semantic.py", line 630, in full
    get_value_fn = getattr(builder, f"get_{dtype.name}")
AttributeError: 'triton._C.libtriton.ir.builder' object has no attribute 'get_fp8e4nv'

This error is being thrown from Triton on this line.

It looks like here is where the triton._C.libtriton.ir.builder is created. I can see that get_fp8e4nv is indeed not defined on this IR builder.

There are fp8 data type getters (i.e., the methods which have a _ty suffix and return the data type itself, such as these), but no plain get_{fp8_dtype} methods for creating a value of fp8 dtype (such as these).

Environment details

Triton: pytorch-triton 3.2.0+git0d4682f0