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

Open
  • Is something wrong?
#3,712 1 comment 1 reaction 0 assignees View on GitHub

Maintainers usually reply within 4 days

@yunweili3 is already working on this.

Since Oct 5, 2026.

  • #3714 by @yunweili3 — open

Assessment

Difficulty
4/5
Estimated time
3-5 days
Newbie friendliness
25/100
Issue type
Bug
Clarity
Clearly specified
Activity status
Stale
Tech stack
cpp, python
Domain
backend

Research direction

The report includes a compile-only Python reproducer using cute.arch.store, TensorSSA, and fp8 types, and notes that Tensor.store() succeeds while direct fp8 vector stores fail during LLVM lowering. Start by running the reproducer with the stated CUTLASS DSL version and trace the store’s LLVM lowering against the working Uint8 bitcast path. Done means direct fp8 vector stores compile to packed-byte vector stores, or unsupported inputs fail clearly at trace time; linked PR #3714 is already open.

Written by the indexing model from the issue text.

Description

? - Needs Triage bug CuTe DSL
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.

Dominant language
C++
Stars
10.7k
Forks
2.2k
Avg merge
3d 20h
Merged PRs (30d)
4

Getting set up

This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

More from NVIDIA/cutlass

All issues in NVIDIA/cutlass

Similar issues

More C++ issues

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.