Description
Your current environment
PyTorch version: 2.8.0a0+gitb32b002
Is debug build: False
CUDA used to build PyTorch: 12.3
ROCM used to build PyTorch: N/A
OS: Ubuntu 22.04.3 LTS (aarch64)
GCC version: (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0
Clang version: 14.0.0-1ubuntu1.1
CMake version: version 3.27.9
Libc version: glibc-2.35
Python version: 3.11.7 (main, Apr 16 2025, 17:30:54) [GCC 11.4.0] (64-bit runtime)
Python platform: Linux-6.2.0-1008-nvidia-64k-aarch64-with-glibc2.35
Is CUDA available: True
CUDA runtime version: 12.3.107
CUDA_MODULE_LOADING set to: LAZY
GPU models and configuration: GPU 0: GH200 480GB
Nvidia driver version: 545.23.08
cuDNN version: Probably one of the following:
/usr/lib/aarch64-linux-gnu/libcudnn.so.9.2.0
/usr/lib/aarch64-linux-gnu/libcudnn_adv.so.9.2.0
/usr/lib/aarch64-linux-gnu/libcudnn_cnn.so.9.2.0
/usr/lib/aarch64-linux-gnu/libcudnn_engines_precompiled.so.9.2.0
/usr/lib/aarch64-linux-gnu/libcudnn_engines_runtime_compiled.so.9.2.0
/usr/lib/aarch64-linux-gnu/libcudnn_graph.so.9.2.0
/usr/lib/aarch64-linux-gnu/libcudnn_heuristic.so.9.2.0
/usr/lib/aarch64-linux-gnu/libcudnn_ops.so.9.2.0
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True
CPU:
Architecture: aarch64
CPU op-mode(s): 64-bit
Byte Order: Little Endian
CPU(s): 72
On-line CPU(s) list: 0-71
Vendor ID: ARM
Model: 0
Thread(s) per core: 1
Core(s) per socket: 72
Socket(s): 1
Stepping: r0p0
Frequency boost: disabled
CPU max MHz: 3447.0000
CPU min MHz: 81.0000
BogoMIPS: 2000.00
Flags: fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm ssbs sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti
L1d cache: 4.5 MiB (72 instances)
L1i cache: 4.5 MiB (72 instances)
L2 cache: 72 MiB (72 instances)
L3 cache: 114 MiB (1 instance)
NUMA node(s): 9
NUMA node0 CPU(s): 0-71
NUMA node1 CPU(s):
NUMA node2 CPU(s):
NUMA node3 CPU(s):
NUMA node4 CPU(s):
NUMA node5 CPU(s):
NUMA node6 CPU(s):
NUMA node7 CPU(s):
NUMA node8 CPU(s):
Vulnerability Itlb multihit: Not affected
Vulnerability L1tf: Not affected
Vulnerability Mds: Not affected
Vulnerability Meltdown: Not affected
Vulnerability Mmio stale data: Not affected
Vulnerability Retbleed: Not affected
Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1: Mitigation; __user pointer sanitization
Vulnerability Spectre v2: Not affected
Vulnerability Srbds: Not affected
Vulnerability Tsx async abort: Not affected
Versions of relevant libraries:
[pip3] numpy==1.26.4
[pip3] torch==2.8.0a0+gitb32b002
[pip3] torchvision==0.19.0a0+89d2b38
[pip3] triton==3.3.0+git3ebf1171
[conda] Could not collect
While attempting to install the latest version of vllm from source on a machine with Hopper architecture (SM90, CUDA 12.3, aarch64), I encountered compilation errors in the vllm-flash-attn module, specifically for all FlashAttention kernels using cutlass::float_e4m3_t.
The error appears to originate from a static_assert in cutlass/include/cute/atom/copy_traits.hpp:
static assertion failed with "Copy_Traits: dst failed to vectorize into registers. Layout is incompatible with this CopyOp."
The error is reproducible across all e4m3-based kernel instantiations under flash_fwd_hdim*e4m3*_sm90.cu
Compilation proceeds correctly for kernels using fp16 and bf16. The failure seems specific to e4m3 and may be related to CUTLASS assumptions not yet fully adapted for Hopper vectorization in this format.
Note: The exact same procedure using the v0.8.4 tag builds successfully, which strongly suggests a regression or change introduced in main (post-v0.8.4) that affects Hopper vectorization for e4m3.
Would appreciate guidance on:
-Whether this is a known issue for e4m3 on SM90
-If certain CUTLASS/CUTE definitions need to be modified for compatibility
-If e4m3 is currently officially supported for FlashAttention 3 kernels on Hopper
-how to solve the issue
How you are installing vllm
git clone https://github.com/vllm-project/vllm.git
cd vllm
git checkout main # or v0.8.5
python use_existing_torch.py
pip install . --no-build-isolation --verboseBefore submitting a new issue...
- Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.