[QST] Support for GEMM with arch 120

[jungi-lee]

What is your question?
Hello!

It seems that CUTLASS 3.8 only supports up to sm_100a. I recently got an RTX 5090 and would like to try the GEMM kernel using the new MMA instructions introduced in sm_120.

Do you have any rough estimates on when the next update will be released that includes support for GEMM with these new MMA instructions?

Thanks in advance!

[thakkarV]

see #2185

however, note that this is only for sm120a . sm120 is already supported by existing sm80, sm86, and sm89 instructions as they are all portable

[jungi-lee]

@thakkarV Hi, thanks for the great update!

I tried some of the recently updated Gemm kernel templates and noticed that FP6 performance is significantly slower than FP8 (around x1/2). Do you have any insights on why?

Unfortunately, I'm running in a docker instance and don't have access to hardware performance counters (Nsight Compute).

• Environment: RTX 5080 + CUDA 12.8

• Instantiated kernel: Block scaled FP6 x FP6 w/ FP32 accum

• Measured performance:
  • The FP6 kernel achieves ~230TFlops, while blockscaled FP8 x FP8 w/ FP32 accumulation achieves ~410 TFlops.
  • I expected FP6 and FP8 to have similar throughput based on this whitepaper, and also check that both kernels are based on the same PTX instruction mma.sync.aligned.m16n8k32.row.col.kind::mxf8f6f4.block_scale.scale_vec::1X.f32.{e2m3 or e4m3}.{e2m3 or e4m3}.f32.ue8m0.

I am assuming the need for padded layout in shared memory and registers for sub-byte FP6 might be causing the slowdown, but I'm not sure. Do you have any thoughts?

[thakkarV]

Please do not instantiate kernels manually, and instead use the profiler to find the best kernel. You are most likely not changing the tile sizes etc when changing data types.

The performance should be pretty good already. If a autotune over the profiler kernels using the new autotuning feature https://github.com/NVIDIA/cutlass/blob/main/media/docs/profiler.md#exhaustive-search-mode-and-top-k-output-ranking-according-to-performance-in-gflopss

shows kernels are still slower than fp8, then we can investigate.

[jungi-lee]

Thanks for pointing out the cutlass profiler.

I swept the sm120 block-scaled Tensor Core Gemm kernels, but FP6 performance still show ~1/2x performance of FP8.

Under the same environment, I built the cutlass profiler:

cmake .. \
  -DCUTLASS_NVCC_ARCHS="120a" \
  -DCUTLASS_LIBRARY_KERNELS="cutlass3x_sm120_bstensorop_*" \
  -DCUTLASS_LIBRARY_INSTANTIATION_LEVEL="max" \
  -DCUTLASS_UNITY_BUILD_ENABLED=ON

make cutlass_profiler -j 

Then I sweep the Gemm kernels as below:

./cutlass_profiler --operation=block_scaled_gemm --enable-best-kernel-for-fixed-shape --m=2048 --n=5120 --k=5120

The results shows that the best FP6 kernel achieves ~226 TFlops, while FP8 kernel reaches ~417 TFlops.
I attached the kernel sweep result below. Could you take a look?

m2048n5120k5120.xlsx

[thakkarV]

thanks a lot! we will take a look

[hwu36]

Do you have the kernel name of the best fp8 and fp6?

[jungi-lee]

Hello @hwu36, I highlighted them in the above comment excel file.

• FP8: cutlass3x_sm120_bstensorop_s16x8x32gemm_ue8m0xe4m3_ue8m0xe4m3_f32_void_e5m2_128x128x128_1x1x1_0_tnt_align128_pingpong_q

• FP6: cutlass3x_sm120_bstensorop_s16x8x32gemm_ue8m0xe2m3_ue8m0xe2m3_f32_void_e5m2_128x128x128_1x1x1_0_tnt_align128_pingpong_q

[edisonchan]

I had tried the instructions above with a GeForce RTX 50 card, I found the GPU usage is always 0%, but there are still very high TFLOPS results output.

[jungi-lee]

Hi @edisonchan, have you measured the throughput of the FP6 × FP6 GEMM kernel (--A=e2m3:row, --B=e2m3:col)? FP8 × FP6 and FP8 x FP8 kernels also achieves near-peak throughput in my benchmarks.
---edited---
I missed the above FP6 kernel results. It seems FP6 kernels achieve around 85% of FP8 throughput on the 5090, so the performance issue of FP6 kernels might be specific to the 5080 GPU.

[edisonchan]

Hi @edisonchan, have you measured the throughput of the FP6 × FP6 GEMM kernel (--A=e2m3:row, --B=e2m3:col)? FP8 × FP6 and FP8 x FP8 kernels also achieves near-peak throughput in my benchmarks. ---edited--- I missed the above FP6 kernel results. It seems FP6 kernels achieve around 85% of FP8 throughput on the 5090, so the performance issue of FP6 kernels might be specific to the 5080 GPU.

Based on the GPU utilization rates and power consumption readings, I suspect the test code isn't actually running on the GPU but is instead being executed through CPU emulation, unless the GPU has activated an alternative monitoring mechanism with inherent flaws during these operations.

[hwu36]

mxfp6 and mxfp8 share the same code path. mxfp6 sass code looks reasonable. One possibility of the perf diff is that mxfp6 hit the power limit. To verify it, you could either 1) use all 0 input to run the kernel. or 2) lock the clock at a very low freq.

The command to lock the freq is

sudo nvidia-smi -i 0 -pm 1
sudo nvidia-smi -lgc 1005 -i 0 #lock to 1005MHz

sudo nvidia-smi -rgc can unlock the lock

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.