NVIDIA / NVIDIA/cutlass

[BUG] cvt int8/uint8 -> bfloat16 emits a sm_90-only instruction on Ampere/Ada instead of the software fallback (arch gate too loose)

Open Beginner friendly
#3,354 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage bug CuTe DSL inactive-30d
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug

cvt_i8_bf16_intrinsic.supported_archs lists Ampere (*Arch.AmpereArchs()) and Ada (*Arch.AdaArchs()), but it lowers to cvt.rn.bf16.s8, which the PTX ISA requires sm_90 or higher. So on a pre-Hopper target, int8/uint8 → bf16 emits an instruction ptxas rejects, instead of taking the else: cutlass_arith.itofp software fallback the dispatch already has (the too-broad allow-list makes the arch in supported_archs check pass, so the fallback is never reached).

Steps/Code to reproduce bug

import cutlass
import cutlass.cute as cute

@cute.jit
def convert(src: cute.Tensor, dst: cute.Tensor):
    dst.store(src.load().to(cutlass.BFloat16))   # Int8 -> BFloat16

# compile for Ampere/Ada, e.g. CUTE_DSL_ARCH=sm_80

The lowered module contains cvt.rn.bf16.s8, which ptxas rejects on sm_80/86/87/89.

Expected behavior

Drop *Arch.AmpereArchs() / *Arch.AdaArchs() from cvt_i8_bf16_intrinsic.supported_archs so pre-sm_90 targets fall through to the existing itofp fallback — matching the sibling cvt_i4_bf16_intrinsic, which correctly restricts itself.

Environment details (please complete the following information):

  • nvidia-cutlass-dsl (also on main @ e8ecfad75b44d1ad56264f5001d877e9e47fe080). Affects Int8/Uint8BFloat16 on sm_80/86/87/89; Hopper/Blackwell unaffected.

Contributor guide

No contributing guide indexed for this repository

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.

Research direction

Start in python/CuTeDSL/cutlass/cute/arch/numeric_conversion.py at cvt_i8_bf16_intrinsic.supported_archs, and compare it with cvt_i4_bf16_intrinsic. Check the related lowering in arch/nvvm_wrappers.py and dispatch in cute/tensor.py, then reproduce for sm_80, sm_86, sm_87, or sm_89. Done means pre-sm_90 targets use the existing itofp fallback while Hopper and newer targets retain the intrinsic.

Written by the indexing model from the issue text.

Assessment

Tech stack
python
Domain
compilers
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Quiet
Clarity
Clearly specified
Newbie friendliness
78/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.