Skip to content

[BUG] [BUG] cute.arch.store doesn't work with fp8 vector #3712

Open

Description

@gau-nernst

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug
A clear and concise description of what the bug is.

cute.arch.store documents val as "scalar Numeric or vector ir.Value". Storing a TensorSSA vector works for f32, bf16, f16, u8, i8 and i16, for example cute.arch.store(ptr, x.load().to(BFloat16)); each lowers to a single st.global.v8.b32. With an fp8 element type (Float8E4M3FN or Float8E5M2), the same call passes tracing, but backend compilation fails with an internal error:

error: NVVM backend compilation failed
error: libNVVM failed while compiling generated device IR.
       LLVM Translation failed for operation: builtin.unrealized_conversion_cast
       Failed creating the llvm::Module.

It is not about bit width: 8-bit integer vectors work, and so does the same fp8 vector after .bitcast(Uint8) (also a single st.global.v8.b32). It looks like a vector of fp8 has no type conversion in the LLVM lowering of the store. Tensor.store() of the same fp8 vector does work, but it emits extra byte shuffles (PRMT) compared with the bitcast path.

Steps/Code to reproduce bug
Follow this guide http://matthewrocklin.com/blog/work/2018/02/28/minimal-bug-reports to craft a minimal bug report. This helps us reproduce the issue you're having and resolve the issue more quickly.

Compile-only (no GPU run needed):

# CuTeDSL MRE: cute.arch.store of an fp8 vector fails in the NVVM backend.
#
# cute.arch.store documents `val` as "scalar Numeric or vector ir.Value". An
# f32 / bf16 / u8 / i8 TensorSSA stores fine (one st.global.v8.b32), but an
# fp8 one (x.to(Float8E4M3FN) or Float8E5M2) traces and then fails while
# lowering:
#   LLVM Translation failed for operation: builtin.unrealized_conversion_cast
# Not a bit-width issue: the same fp8 vector bitcast to Uint8 stores fine.
#
# Compile-only; prints one line per variant.
#   python arch_store_fp8_vector.py
import importlib.metadata

import cutlass
import cutlass.cute as cute
from cutlass import BFloat16, Float8E4M3FN, Float8E5M2, Float32, Int8, Uint8

N = 32  # elements per thread


def make(dtype, via_u8):
    @cute.kernel
    def kernel(out: cute.Tensor):
        tid, _, _ = cute.arch.thread_idx()
        x = cute.make_rmem_tensor(N, Float32)
        for i in cutlass.range_constexpr(N):
            x[i] = Float32(tid + i)
        v = x.load().to(dtype)
        dst = out[tid, None].iterator
        if cutlass.const_expr(via_u8):
            cute.arch.store(cute.recast_ptr(dst, dtype=Uint8), v.bitcast(Uint8))
        else:
            cute.arch.store(dst, v)

    @cute.jit
    def run(out: cute.Tensor):
        kernel(out).launch(grid=(1, 1, 1), block=(32, 1, 1))

    return run


print("nvidia-cutlass-dsl", importlib.metadata.version("nvidia-cutlass-dsl"))
for dtype, via_u8 in (
    (Float32, False),
    (BFloat16, False),
    (Uint8, False),
    (Int8, False),
    (Float8E4M3FN, False),
    (Float8E5M2, False),
    (Float8E4M3FN, True),
):
    t = cute.runtime.make_fake_tensor(dtype, (32, N), stride=(N, 1), assumed_align=32)
    name = f"{dtype.__name__:13s} {'bitcast->u8' if via_u8 else 'direct'}"
    try:
        cute.compile(make(dtype, via_u8), t)
        print(f"  {name}: ok")
    except Exception as e:  # noqa: BLE001
        msg = " ".join(str(e).split())
        key = "unrealized_conversion_cast"
        print(f"  {name}: FAILED ({key if key in msg else msg[:200]})")

Output on 4.8.0:

nvidia-cutlass-dsl 4.8.0
  Float32       direct: ok
  BFloat16      direct: ok
  Uint8         direct: ok
  Int8          direct: ok
  Float8E4M3FN  direct: FAILED (unrealized_conversion_cast)
  Float8E5M2    direct: FAILED (unrealized_conversion_cast)
  Float8E4M3FN  bitcast->u8: ok

Expected behavior
A clear and concise description of what you expected to happen.

cute.arch.store(ptr, fp8_vector) lowers to a vector store of the packed bytes, the same as the Uint8 bitcast path. If fp8 vectors are intentionally unsupported, the DSL should raise a clear error at trace time, not an ICE in libNVVM.

Environment details (please complete the following information):

  • Environment location: [Bare-metal, Docker, Cloud(specify cloud provider)]

  • Environment location: Bare-metal

  • GPU: NVIDIA GB300 (sm_103a), driver 610.43.02

  • OS: Ubuntu 24.04.4 LTS (aarch64), Python 3.12.13

  • nvidia-cutlass-dsl: 4.8.0

Additional context
Add any other context about the problem here.

Activity

  1. yunweili3 commented on Oct 5, 2026

    @yunweili3

    Reproduced on B200 (sm_100a) with 4.8.0, same output as above. The scope is a bit wider than the store: cute.arch.load of an fp8 vector (ir.VectorType.get([32], Float8E4M3FN.mlir_type)) crashes the compiler, and f16 / bf16 / fp8 scalars fail the verifier (Unsupported FP type for ExtStoreOp, operand #0 must be LLVM type with size, but got 'f8E4M3FN').

    Root cause: cute.arch.store / load hand the value type straight to nvvm.store.ext / nvvm.load.ext, and those ops only carry integer, f32 and f64 values, vectors of those, and f16/bf16 vectors. An fp8 vector passes the verifier but has no LLVM type, which is the unrealized_conversion_cast failure you saw. Your .bitcast(Uint8) workaround works precisely because it never shows a float8 type to the ext op.

    PR #3714 does that bitcast inside the wrappers: other float types are moved as a same-width integer (scalar → iN, vector → vector<bits/8 x i8>, which also covers f6/f4 vectors) and load bitcasts back to the requested type. The fp8 vector store then lowers to one st.global.v8.b32 with SASS identical to the Uint8 path (no extra PRMT); previously working types are unchanged.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions