[BUG] cvt int8/uint8 -> bfloat16 emits a sm_90-only instruction on Ampere/Ada instead of the software fallback (arch gate too loose)
Nobody has claimed this yet.
- 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 onmain@e8ecfad75b44d1ad56264f5001d877e9e47fe080). AffectsInt8/Uint8→BFloat16on sm_80/86/87/89; Hopper/Blackwell unaffected.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- 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