Is there exist difference between SM89(Like 4090) and SM90(Like H100/H800) device?

[MARD1NO]

I have a question is 4090 support FP8 Gemm, but it seems there is no specialization for SM89(Like GMMA), I only found SM90 specialization in cute. I wonder if sm89 use the same MMA instruction as SM90?

And I also want to ask whether CUDA MATH will support FP8 math instruction like add, multiply etc in future. I only found fp8 datatype convert function in CUDA Documentations.

[alecco]

sm_90/sm_90a has a lot more features than sm_89.

At the moment, CUTLASS only supports fp8 for sm_90/sm_90a due to nvcc lacking sm_89 fp8 support. See the previous discussion "it is coming".

[MARD1NO]

Thanks!

[hwu36]

And I also want to ask whether CUDA MATH will support FP8 math instruction like add, multiply etc in future

not on ada and hopper. you could convert them to fp16 first and use hadd2 and hmul2

[MARD1NO]

Thanks! I also want to know in ada and hopper, the minimum MMA operator is WGMMA? It works on the warp group level. Is there still exist WarpLevel mma operation Like ampere m16n8k16 mma in Ada/Hopper？

[hwu36]

ada is very similar to ampere, it does not use wgmma. all ampere kernels can still run on hopper/ada.

[alecco]

FWIW this seems to be supported in latest CUTLASS 3.5 (need CUDA 12.4+).