Attempting to use SM89_16x8x32_F32E4M3E4M3F32_TN without CUTE_ARCH_MMA_F16_SM89_ENABLED"` failed

[kunibald413]

running a workflow after python setup.py install gives this error:

ComfyUI/custom_nodes/LTX-Video-Q8-Kernels/csrc/gemm/mma_sm89_fp16.hpp:80: static void cute::SM89_16x8x32_F16E4M3E4M3F16_TN::fma(unsigned int &, unsigned int &, const unsigned int &, const unsigned int &, const unsigned int &, const unsigned int &, const unsigned int &, const unsigned int &, const unsigned int &, const unsigned int &): block: [3,1,0], thread: [223,0,0] Assertion `0 && "Attempting to use SM89_16x8x32_F32E4M3E4M3F32_TN without CUTE_ARCH_MMA_F16_SM89_ENABLED"` failed.
terminate called after throwing an instance of 'c10::Error'

python 3.11.10

torch==2.6.0+cu124
torchaudio==2.6.0+cu124
torchsde==0.2.6
torchvision==0.21.0+cu124

4090

NVIDIA-SMI 565.57.01              Driver Version: 565.57.01      CUDA Version: 12.7

[frankchieng1982]

If your GPU is not SM89 (Hopper), you need to modify the code or build configuration to target your GPU’s architecture (e.g., SM80 for Ampere, SM90 for Ada Lovelace). first check if the mma_sm80.hpp or mma_sm90.hpp exist under third_party/cutlass/include/cute/arch directory,if so,the SM90::GMMA::MMA_64x8x32_F16E4M3E4M3_SS_TN uses float_e4m3_t. Confirm that fp8_gemm_cuda.cu passes compatible data (E4M3 for A/B, FP16 for C/D),the next step is modify the mma_sm89_fp16.hpp under csrc/gemm directory

struct SM89_16x8x32_F16E4M3E4M3F16_TN
{
  using DRegisters = uint32_t[2];
  using ARegisters = uint32_t[4];
  using BRegisters = uint32_t[2];
  using CRegisters = uint32_t[2];

  CUTE_HOST_DEVICE static void
  fma(uint32_t      & d0, uint32_t      & d1,
      uint32_t const& a0, uint32_t const& a1, uint32_t const& a2, uint32_t const& a3,
      uint32_t const& b0, uint32_t const& b1,
      uint32_t const& c0, uint32_t const& c1)
  {
#if defined(CUTE_ARCH_MMA_F16_SM89_ENABLED) && defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 890)
    asm(
      "mma.sync.aligned.m16n8k32.row.col.f16.e4m3.e4m3.f16 "
      "{%0,%1}, {%2,%3,%4,%5}, {%6,%7}, {%8,%9};\n"
      : "=r"(d0), "=r"(d1)
      :
        "r"(a0), "r"(a1), "r"(a2), "r"(a3),
        "r"(b0), "r"(b1),
        "r"(c0), "r"(c1)
  );
#elif defined(CUTE_ARCH_MMA_SM90_ENABLED) && defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
    // SM90 MMA for Ada Lovelace (RTX 4090)
    asm volatile(
      "mma.sync.aligned.m16n8k32.row.col.f16.e4m3.e4m3.f16 "
      "{%0,%1}, {%2,%3,%4,%5}, {%6,%7}, {%8,%9};\n"
      : "=r"(d0), "=r"(d1)
      : "r"(a0), "r"(a1), "r"(a2), "r"(a3),
        "r"(b0), "r"(b1),
        "r"(c0), "r"(c1)
    );
#else
    // CUTE_INVALID_CONTROL_PATH("Attempting to use SM89_16x8x32_F32E4M3E4M3F32_TN without CUTE_ARCH_MMA_F16_SM89_ENABLED");
    // Fallback for non-SM89 architectures (e.g., SM90, CUDA < 12.8)
    d0 = c0; // Preserve accumulator inputs
    d1 = c1;
#endif
  }
};

[Kozatiju]

Thanks FrankChieng (@frankchieng1982) ! I couldn't get fix the issue with the modifications you suggested. I'm currently using a 4080 so maybe it could be related ( it's also using a ADA Lovelace architecure ). I wonder if anyone with a 4090 or less has managed to use the FP8 version so far ?

[Opeth66]

make sure you have C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8 in you're Environment Variables --> System variables, before running ..\python_embeded\python.exe setup.py install

[KannManMachen]

Same Problem here:

.../LTX-Video-Q8-Kernels\csrc\gemm\mma_sm89_fp16.hpp:80: block: [12,5,0], thread: [255,0,0] Assertion 0 && "Attempting to use SM89_16x8x32_F32E4M3E4M3F32_TN without CUTE_ARCH_MMA_F16_SM89_ENABLED" failed.

Running a 4090. ChatGPT said:

„That error message means that the kernel code you're using tries to use a specific CUDA tensor core instruction (SM89 for FP8/FP16) — but the corresponding architecture macro (CUTE_ARCH_MMA_F16_SM89_ENABLED) is not defined, so it fails intentionally.

You're likely running on an Ada Lovelace GPU (e.g. RTX 40-series) with compute capability 8.9 (SM89), and the CUTLASS-based kernel you compiled expects this feature to be explicitly enabled.“

ChatGPT suggested adding code corrections inside of the setup.py (LTX-Video-Q8-Kernels folder):
"-DCUTE_ARCH_MMA_SM89_ENABLED",
"-DCUTE_ARCH_MMA_F16_SM89_ENABLED",
"-DCUTE_USE_DEVICE_ATOMICS=1",

... that didnt work neither so I got stuck with the following statement from Chatty:

Your issue is that:

- The GEMM-related CUDA code (in mma_sm89_fp16.hpp) uses inline PTX assembly or syntax not supported under Windows NVCC/MSVC.

- That means the Windows toolchain can’t compile this — even though the hardware (RTX 4090) fully supports it.

What You Can Do
Use a Linux or WSL2 Environment

Can anyone make sense of this statement?

[Kozatiju]

make sure you have C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8 in you're Environment Variables --> System variables, before running ..\python_embeded\python.exe setup.py install

Thanks ! Installing the proper CUDA clearly made the difference ! that was so obvious 😅

[KannManMachen]

make sure you have C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8 in you're Environment Variables --> System variables, before running ..\python_embeded\python.exe setup.py install

I did set CUDA Toolkit 12.9 ... but it still did not work for me.

[garrick99]

I've submitted a fix for this in PR #22.

Root cause: The CUTE_ARCH_MMA_F16_SM89_ENABLED macro requires CUDA >= 12.8 for the F16 accumulator variant. Users with CUDA 12.4-12.7 get the assertion failure because the macro isn't defined, but the code path is still reached at runtime.

The F32 accumulator variant (mma.sync...f32.e4m3.e4m3.f32) works with CUDA >= 12.4.

The fix (in PR #22):

• Keeps the CUDA 12.8 requirement for the F16 variant (confirmed: CUDA 12.4 ptxas rejects the F16 FP8 MMA instruction)
• Adds __CUDA_ARCH__ < 1000 guard to prevent Blackwell (SM100/SM120) from hitting SM89-specific MMA atoms
• Fixes Blackwell GPU detection in setup.py so RTX 5090 users can install

Tested on both RTX 4090 (SM89, CUDA 12.4) and RTX 5090 (SM120, CUDA 13.0). RTX 4090 gets native FP8 GEMM with F32 accumulator; RTX 5090 gracefully falls back to torch._scaled_mm.