deepseek-ai/DeepGEMM
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepGEMM.git synced 2026-11-05 00:47:26 +00:00
Code Issues 162 Projects Packages Wiki Activity

[GH-ISSUE #160] v2 seems to be slower than v1 on hopper on some cases #92

New issue
Open
opened 2026-10-26 21:12:17 +00:00 by bsakaguchi · 18 comments
bsakaguchi commented 2026-10-26 21:12:17 +00:00
Owner
Copy link

Originally created by @fzyzcjy on GitHub (Aug 14, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/160

related: https://github.com/sgl-project/sglang/pull/9167

I briefly look at profile, and see some deepgemm kernels have slower speed. before digging deeper, I create an issue to report to know e.g. is it on the roadmap and is not yet optimized.

speed:

Image

baseline:

1755163919.4467359-TP-6.trace.json.gz

pr:

1755161669.4129217-TP-6.trace.json.gz

example screenshot (for a stable-shape gemm. the most time-consuming gemm is moe, but that changes a lot since expert is imbalanced, so I do not screenshot here)

baseline:

Image

pr:

Image
Originally created by @fzyzcjy on GitHub (Aug 14, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/160 related: https://github.com/sgl-project/sglang/pull/9167 I briefly look at profile, and see some deepgemm kernels have slower speed. before digging deeper, I create an issue to report to know e.g. is it on the roadmap and is not yet optimized. speed: <img width="589" height="241" alt="Image" src="https://github.com/user-attachments/assets/c9954731-909e-4035-92a2-164aa3c257cb" /> baseline: [1755163919.4467359-TP-6.trace.json.gz](https://github.com/user-attachments/files/21770886/1755163919.4467359-TP-6.trace.json.gz) pr: [1755161669.4129217-TP-6.trace.json.gz](https://github.com/user-attachments/files/21770892/1755161669.4129217-TP-6.trace.json.gz) example screenshot (for a stable-shape gemm. the most time-consuming gemm is moe, but that changes a lot since expert is imbalanced, so I do not screenshot here) baseline: <img width="1563" height="530" alt="Image" src="https://github.com/user-attachments/assets/d05eadab-1d6c-432d-81a1-a818d7f05f8d" /> pr: <img width="1597" height="525" alt="Image" src="https://github.com/user-attachments/assets/3c8e8912-0859-4a7c-bd2a-f4bc7332f557" />
bsakaguchi commented 2026-10-26 21:12:19 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 15, 2025):

Made a docker-based sample that I test on both H200 and H100 as follows. I think probably some of my env has something wrong, especially given the H100 v2 is so slow...

docker rm --force tom_cuda_20250815_v1
docker rm --force tom_cuda_20250815_v2

docker run -it --name tom_cuda_20250815_v1 --gpus all --privileged --cap-add SYS_PTRACE --cap-add SYS_ADMIN nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 /bin/bash
docker run -it --name tom_cuda_20250815_v2 --gpus all --privileged --cap-add SYS_PTRACE --cap-add SYS_ADMIN nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 /bin/bash

# ================ configure common (src: sgl Dockerfile) ================

apt update && apt install -y git vim

apt update && apt install wget -y && apt install software-properties-common -y \
 && add-apt-repository ppa:deadsnakes/ppa -y \
  && apt install python3.12-full python3.12-dev python3.10-venv -y \
 && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.10 1 \
 && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.12 2 \
 && update-alternatives --set python3 /usr/bin/python3.12 \
 && wget https://bootstrap.pypa.io/get-pip.py \
 && python3 get-pip.py

pip install torch

ln -s $(which python3) /usr/bin/python

pip install cuda-python

# ================ configure v1 ================

(cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && git checkout 391755ada0ffefa9a6a52b6f14dcaf22d1a463e0 && pip install .)

# ================ configure v2 ================

(cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && ./install.sh)

# ================ test v1 ================

(cd /tmp/DeepGEMM/tests && python3 test_core.py)

# ================ test v2 ================

(cd /tmp/DeepGEMM/tests && python3 test_fp8.py)

H200, v1

root@c58518f70edf:/# (cd /tmp/DeepGEMM/tests && CUDA_VISIBLE_DEVICES=7 python3 test_core.py)
NVRTC version: (12, 8)
/usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.)
  cpu = _conversion_method_template(device=torch.device("cpu"))
Library path:
 > ['/usr/local/lib/python3.12/dist-packages/deep_gemm']

Testing grouped contiguous GEMM:
 > Performance (num_groups=4, m=36224, n=4096, k=7168): 1564 us | throughput: 1360 TFLOPS,  431 GB/s
 > Performance (num_groups=4, m=30976, n=7168, k=2048):  707 us | throughput: 1286 TFLOPS,  801 GB/s
 > Performance (num_groups=8, m=32896, n=4096, k=7168): 1438 us | throughput: 1343 TFLOPS,  515 GB/s
 > Performance (num_groups=8, m=32128, n=7168, k=2048):  742 us | throughput: 1271 TFLOPS,  867 GB/s

H200, v2

root@f23fc66d7f58:/# (cd /tmp/DeepGEMM/tests && CUDA_VISIBLE_DEVICES=7 python3 test_fp8.py)
/usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.)
  cpu = _conversion_method_template(device=torch.device("cpu"))
Library path:
 > ['/usr/local/lib/python3.12/dist-packages/deep_gemm']

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1781 us | 1169 TFLOPS |  376 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT):  845 us | 1255 TFLOPS |  772 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1643 us | 1157 TFLOPS |  450 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  745 us | 1230 TFLOPS |  847 GB/s

H100, v1

root@89a7a98c3f30:/# (cd /tmp/DeepGEMM/tests && python3 test_core.py)
NVRTC version: (12, 8)
/usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.)
  cpu = _conversion_method_template(device=torch.device("cpu"))
Library path:
 > ['/usr/local/lib/python3.12/dist-packages/deep_gemm']

Testing grouped contiguous GEMM:
 > Performance (num_groups=4, m=36224, n=4096, k=7168): 1560 us | throughput: 1364 TFLOPS,  432 GB/s
 > Performance (num_groups=4, m=30976, n=7168, k=2048):  710 us | throughput: 1280 TFLOPS,  797 GB/s
 > Performance (num_groups=8, m=32896, n=4096, k=7168): 1438 us | throughput: 1343 TFLOPS,  515 GB/s
 > Performance (num_groups=8, m=32128, n=7168, k=2048):  747 us | throughput: 1263 TFLOPS,  862 GB/s

H100, v2

root@8dcc4d1d7e21:/# (cd /tmp/DeepGEMM/tests && python3 test_fp8.py)
/usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.)
  cpu = _conversion_method_template(device=torch.device("cpu"))
Library path:
 > ['/usr/local/lib/python3.12/dist-packages/deep_gemm']

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 2064 us | 1009 TFLOPS |  325 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 1103 us |  961 TFLOPS |  592 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1919 us |  991 TFLOPS |  385 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  898 us | 1021 TFLOPS |  702 GB/s
<!-- gh-comment-id:3191755674 --> @fzyzcjy commented on GitHub (Aug 15, 2025): Made a docker-based sample that I test on both H200 and H100 as follows. I think probably some of my env has something wrong, especially given the H100 v2 is so slow... ``` docker rm --force tom_cuda_20250815_v1 docker rm --force tom_cuda_20250815_v2 docker run -it --name tom_cuda_20250815_v1 --gpus all --privileged --cap-add SYS_PTRACE --cap-add SYS_ADMIN nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 /bin/bash docker run -it --name tom_cuda_20250815_v2 --gpus all --privileged --cap-add SYS_PTRACE --cap-add SYS_ADMIN nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 /bin/bash # ================ configure common (src: sgl Dockerfile) ================ apt update && apt install -y git vim apt update && apt install wget -y && apt install software-properties-common -y \ && add-apt-repository ppa:deadsnakes/ppa -y \ && apt install python3.12-full python3.12-dev python3.10-venv -y \ && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.10 1 \ && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.12 2 \ && update-alternatives --set python3 /usr/bin/python3.12 \ && wget https://bootstrap.pypa.io/get-pip.py \ && python3 get-pip.py pip install torch ln -s $(which python3) /usr/bin/python pip install cuda-python # ================ configure v1 ================ (cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && git checkout 391755ada0ffefa9a6a52b6f14dcaf22d1a463e0 && pip install .) # ================ configure v2 ================ (cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && ./install.sh) # ================ test v1 ================ (cd /tmp/DeepGEMM/tests && python3 test_core.py) # ================ test v2 ================ (cd /tmp/DeepGEMM/tests && python3 test_fp8.py) ``` H200, v1 ``` root@c58518f70edf:/# (cd /tmp/DeepGEMM/tests && CUDA_VISIBLE_DEVICES=7 python3 test_core.py) NVRTC version: (12, 8) /usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.) cpu = _conversion_method_template(device=torch.device("cpu")) Library path: > ['/usr/local/lib/python3.12/dist-packages/deep_gemm'] Testing grouped contiguous GEMM: > Performance (num_groups=4, m=36224, n=4096, k=7168): 1564 us | throughput: 1360 TFLOPS, 431 GB/s > Performance (num_groups=4, m=30976, n=7168, k=2048): 707 us | throughput: 1286 TFLOPS, 801 GB/s > Performance (num_groups=8, m=32896, n=4096, k=7168): 1438 us | throughput: 1343 TFLOPS, 515 GB/s > Performance (num_groups=8, m=32128, n=7168, k=2048): 742 us | throughput: 1271 TFLOPS, 867 GB/s ``` H200, v2 ``` root@f23fc66d7f58:/# (cd /tmp/DeepGEMM/tests && CUDA_VISIBLE_DEVICES=7 python3 test_fp8.py) /usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.) cpu = _conversion_method_template(device=torch.device("cpu")) Library path: > ['/usr/local/lib/python3.12/dist-packages/deep_gemm'] Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1781 us | 1169 TFLOPS | 376 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 845 us | 1255 TFLOPS | 772 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1643 us | 1157 TFLOPS | 450 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 745 us | 1230 TFLOPS | 847 GB/s ``` H100, v1 ``` root@89a7a98c3f30:/# (cd /tmp/DeepGEMM/tests && python3 test_core.py) NVRTC version: (12, 8) /usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.) cpu = _conversion_method_template(device=torch.device("cpu")) Library path: > ['/usr/local/lib/python3.12/dist-packages/deep_gemm'] Testing grouped contiguous GEMM: > Performance (num_groups=4, m=36224, n=4096, k=7168): 1560 us | throughput: 1364 TFLOPS, 432 GB/s > Performance (num_groups=4, m=30976, n=7168, k=2048): 710 us | throughput: 1280 TFLOPS, 797 GB/s > Performance (num_groups=8, m=32896, n=4096, k=7168): 1438 us | throughput: 1343 TFLOPS, 515 GB/s > Performance (num_groups=8, m=32128, n=7168, k=2048): 747 us | throughput: 1263 TFLOPS, 862 GB/s ``` H100, v2 ``` root@8dcc4d1d7e21:/# (cd /tmp/DeepGEMM/tests && python3 test_fp8.py) /usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.) cpu = _conversion_method_template(device=torch.device("cpu")) Library path: > ['/usr/local/lib/python3.12/dist-packages/deep_gemm'] Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 2064 us | 1009 TFLOPS | 325 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 1103 us | 961 TFLOPS | 592 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1919 us | 991 TFLOPS | 385 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 898 us | 1021 TFLOPS | 702 GB/s ```
bsakaguchi commented 2026-10-26 21:12:19 +00:00
Author
Owner
Copy link

@coconut902 commented on GitHub (Aug 18, 2025):

As NVCC 12.9 will automatically do the FFMA interleaving, all post optimizations will be no longer supported

This might be the reason. Try NVCC 12.9?

<!-- gh-comment-id:3196510455 --> @coconut902 commented on GitHub (Aug 18, 2025): > As NVCC 12.9 will automatically do the FFMA interleaving, all post optimizations will be no longer supported This might be the reason. Try NVCC 12.9?
bsakaguchi commented 2026-10-26 21:12:20 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 18, 2025):

The docker image is: nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 thus I think it seems to be 12.9.

The host driver is a bit old though, and I am wondering whether that may be the case or not... (do not have time to check it yet)

<!-- gh-comment-id:3196554932 --> @fzyzcjy commented on GitHub (Aug 18, 2025): The docker image is: `nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04` thus I think it seems to be 12.9. The host driver is a bit old though, and I am wondering whether that may be the case or not... (do not have time to check it yet)
bsakaguchi commented 2026-10-26 21:12:20 +00:00
Author
Owner
Copy link

@coconut902 commented on GitHub (Aug 19, 2025):

The docker image is: nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 thus I think it seems to be 12.9.

The host driver is a bit old though, and I am wondering whether that may be the case or not... (do not have time to check it yet)

But your v1 log shows you were using NVRTC version: (12, 8)...

<!-- gh-comment-id:3199405548 --> @coconut902 commented on GitHub (Aug 19, 2025): > The docker image is: `nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04` thus I think it seems to be 12.9. > > The host driver is a bit old though, and I am wondering whether that may be the case or not... (do not have time to check it yet) But your v1 log shows you were using `NVRTC version: (12, 8)`...
bsakaguchi commented 2026-10-26 21:12:20 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 19, 2025):

Yes I am also confused... nvcc is new though. indeed if nvcc is old, deepgemm will provide an warning message.

Image
<!-- gh-comment-id:3199444236 --> @fzyzcjy commented on GitHub (Aug 19, 2025): Yes I am also confused... nvcc is new though. indeed if nvcc is old, deepgemm will provide an warning message. <img width="548" height="353" alt="Image" src="https://github.com/user-attachments/assets/678773e9-1b1f-487d-b77b-c91b63282d36" />
bsakaguchi commented 2026-10-26 21:12:21 +00:00
Author
Owner
Copy link

@coconut902 commented on GitHub (Aug 19, 2025):

You can update your nvrtc like:

pip install nvidia-cuda-nvrtc-cu12==12.9.86
<!-- gh-comment-id:3199483536 --> @coconut902 commented on GitHub (Aug 19, 2025): You can update your nvrtc like: ``` pip install nvidia-cuda-nvrtc-cu12==12.9.86 ```
bsakaguchi commented 2026-10-26 21:12:21 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 19, 2025):

Good point. (but it seems v2 does not use nvrtc)

Image
<!-- gh-comment-id:3199549944 --> @fzyzcjy commented on GitHub (Aug 19, 2025): Good point. (but it seems v2 does not use nvrtc) <img width="1064" height="249" alt="Image" src="https://github.com/user-attachments/assets/d7569d96-b928-4619-a927-b50562ee8cb5" />
bsakaguchi commented 2026-10-26 21:12:22 +00:00
Author
Owner
Copy link

@coconut902 commented on GitHub (Aug 19, 2025):

Recent PRs have supported NVRTC. README is outdated.

<!-- gh-comment-id:3199572553 --> @coconut902 commented on GitHub (Aug 19, 2025): Recent PRs have supported NVRTC. README is outdated.
bsakaguchi commented 2026-10-26 21:12:22 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 19, 2025):

looks reasonable, I will upgrade and try again

<!-- gh-comment-id:3199582571 --> @fzyzcjy commented on GitHub (Aug 19, 2025): looks reasonable, I will upgrade and try again
bsakaguchi commented 2026-10-26 21:12:22 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 20, 2025):

@zhonghui-J

Commands

docker rm --force tom_cuda_20250815_v2
docker run -it --name tom_cuda_20250815_v2 --gpus all --privileged --cap-add SYS_PTRACE --cap-add SYS_ADMIN nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 /bin/bash

====================================

apt update && apt install -y git vim

apt update && apt install wget -y && apt install software-properties-common -y \
 && add-apt-repository ppa:deadsnakes/ppa -y \
  && apt install python3.12-full python3.12-dev python3.10-venv -y \
 && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.10 1 \
 && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.12 2 \
 && update-alternatives --set python3 /usr/bin/python3.12 \
 && wget https://bootstrap.pypa.io/get-pip.py \
 && python3 get-pip.py

pip install torch

ln -s $(which python3) /usr/bin/python

pip install cuda-python==12.9

(cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && ./install.sh)

# NOTE ADD
pip install nvidia-cuda-nvrtc-cu12==12.9.86

pip install nvidia-cuda-runtime-cu12==12.9.79

# remove mm test
vim /tmp/DeepGEMM/tests/test_fp8.py

(cd /tmp/DeepGEMM/tests && python3 test_fp8.py)

Results

cuda-bindings            12.9.2
cuda-pathfinder          1.1.0
cuda-python              12.9.0
nvidia-cuda-cupti-cu12   12.8.90
nvidia-cuda-nvrtc-cu12   12.9.86
nvidia-cuda-runtime-cu12 12.9.79

nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Apr__9_19:24:57_PDT_2025
Cuda compilation tools, release 12.9, V12.9.41
Build cuda_12.9.r12.9/compiler.35813241_0

root@9cf31f6d831c:/# (cd /tmp/DeepGEMM/tests && python3 test_fp8.py)
/usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.)
  cpu = _conversion_method_template(device=torch.device("cpu"))
Library path:
 > ['/usr/local/lib/python3.12/dist-packages/deep_gemm']

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 2064 us | 1009 TFLOPS |  325 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 1101 us |  963 TFLOPS |  593 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1917 us |  992 TFLOPS |  386 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  895 us | 1025 TFLOPS |  705 GB/s
<!-- gh-comment-id:3203354009 --> @fzyzcjy commented on GitHub (Aug 20, 2025): @zhonghui-J Commands <details> ``` docker rm --force tom_cuda_20250815_v2 docker run -it --name tom_cuda_20250815_v2 --gpus all --privileged --cap-add SYS_PTRACE --cap-add SYS_ADMIN nvidia/cuda:12.9.0-cudnn-devel-ubuntu22.04 /bin/bash ==================================== apt update && apt install -y git vim apt update && apt install wget -y && apt install software-properties-common -y \ && add-apt-repository ppa:deadsnakes/ppa -y \ && apt install python3.12-full python3.12-dev python3.10-venv -y \ && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.10 1 \ && update-alternatives --install /usr/bin/python3 python3 /usr/bin/python3.12 2 \ && update-alternatives --set python3 /usr/bin/python3.12 \ && wget https://bootstrap.pypa.io/get-pip.py \ && python3 get-pip.py pip install torch ln -s $(which python3) /usr/bin/python pip install cuda-python==12.9 (cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && ./install.sh) # NOTE ADD pip install nvidia-cuda-nvrtc-cu12==12.9.86 pip install nvidia-cuda-runtime-cu12==12.9.79 # remove mm test vim /tmp/DeepGEMM/tests/test_fp8.py (cd /tmp/DeepGEMM/tests && python3 test_fp8.py) ``` </details> Results ``` cuda-bindings 12.9.2 cuda-pathfinder 1.1.0 cuda-python 12.9.0 nvidia-cuda-cupti-cu12 12.8.90 nvidia-cuda-nvrtc-cu12 12.9.86 nvidia-cuda-runtime-cu12 12.9.79 nvcc --version nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2025 NVIDIA Corporation Built on Wed_Apr__9_19:24:57_PDT_2025 Cuda compilation tools, release 12.9, V12.9.41 Build cuda_12.9.r12.9/compiler.35813241_0 root@9cf31f6d831c:/# (cd /tmp/DeepGEMM/tests && python3 test_fp8.py) /usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.) cpu = _conversion_method_template(device=torch.device("cpu")) Library path: > ['/usr/local/lib/python3.12/dist-packages/deep_gemm'] Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 2064 us | 1009 TFLOPS | 325 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 1101 us | 963 TFLOPS | 593 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1917 us | 992 TFLOPS | 386 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 895 us | 1025 TFLOPS | 705 GB/s ```
bsakaguchi commented 2026-10-26 21:12:23 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 20, 2025):

@zhonghui-J btw, can you reproduce the correct performance using v2 code on your machines? If so, could you please try my docker reproduction above to see whether it is slow on your machine, and/or show your envs. Thanks!

EDIT: or could you please provide reproduction steps to your fast performance scenario

<!-- gh-comment-id:3203943786 --> @fzyzcjy commented on GitHub (Aug 20, 2025): @zhonghui-J btw, can you reproduce the correct performance using v2 code on your machines? If so, could you please try my docker reproduction above to see whether it is slow on your machine, and/or show your envs. Thanks! EDIT: or could you please provide reproduction steps to your fast performance scenario
bsakaguchi commented 2026-10-26 21:12:23 +00:00
Author
Owner
Copy link

@coconut902 commented on GitHub (Aug 20, 2025):

@fzyzcjy I just reproduced the bad performance.
And I found it's mainly caused by get_num_1d_blocks_per_group() in deep_gemm/include/deep_gemm/common/scheduler.cuh.

After I commented some lines like the following, the performance gets better.

template <GemmType kGemmType, uint32_t BLOCK_M, uint32_t BLOCK_N, uint32_t kNumSMs, bool kIsMulticastOnA>
static constexpr uint32_t get_num_1d_blocks_per_group() {
    // Select the best from candidates
    uint32_t num_best_blocks = 0, min_usage = cute::numeric_limits<uint32_t>::max();
    // if constexpr (kGemmType == GemmType::MGroupedContiguous or
    //               kGemmType == GemmType::MGroupedMasked) {
    //     // For grouped GEMMs, let weights always stay in the L2 cache and read activations by once
    //     num_best_blocks = kNumSMs;
    // } else {
        for (const auto& candidate: {8u, 16u}) {
            const auto& usage = kIsMulticastOnA ?
                        candidate * BLOCK_N + constexpr_ceil_div(kNumSMs, candidate) * BLOCK_M: // Grouping on N
                        candidate * BLOCK_M + constexpr_ceil_div(kNumSMs, candidate) * BLOCK_N; // Grouping on M
            if (usage < min_usage)
                min_usage = usage, num_best_blocks = candidate;
        }
    // }
    return num_best_blocks;
}

On my H800,
before:

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 2061 us | 1010 TFLOPS |  325 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 1099 us |  964 TFLOPS |  594 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1921 us |  990 TFLOPS |  385 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  920 us |  996 TFLOPS |  686 GB/s

after:

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1490 us | 1397 TFLOPS |  450 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT):  801 us | 1323 TFLOPS |  815 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1370 us | 1388 TFLOPS |  540 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  699 us | 1312 TFLOPS |  903 GB/s

@RayWang96 Could you please tell us what's the point of For grouped GEMMs, let weights always stay in the L2 cache and read activations by once?

<!-- gh-comment-id:3204142418 --> @coconut902 commented on GitHub (Aug 20, 2025): @fzyzcjy I just reproduced the bad performance. And I found it's mainly caused by `get_num_1d_blocks_per_group()` in `deep_gemm/include/deep_gemm/common/scheduler.cuh`. After I commented some lines like the following, the performance gets better. ``` template <GemmType kGemmType, uint32_t BLOCK_M, uint32_t BLOCK_N, uint32_t kNumSMs, bool kIsMulticastOnA> static constexpr uint32_t get_num_1d_blocks_per_group() { // Select the best from candidates uint32_t num_best_blocks = 0, min_usage = cute::numeric_limits<uint32_t>::max(); // if constexpr (kGemmType == GemmType::MGroupedContiguous or // kGemmType == GemmType::MGroupedMasked) { // // For grouped GEMMs, let weights always stay in the L2 cache and read activations by once // num_best_blocks = kNumSMs; // } else { for (const auto& candidate: {8u, 16u}) { const auto& usage = kIsMulticastOnA ? candidate * BLOCK_N + constexpr_ceil_div(kNumSMs, candidate) * BLOCK_M: // Grouping on N candidate * BLOCK_M + constexpr_ceil_div(kNumSMs, candidate) * BLOCK_N; // Grouping on M if (usage < min_usage) min_usage = usage, num_best_blocks = candidate; } // } return num_best_blocks; } ``` On my H800, before: ``` Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 2061 us | 1010 TFLOPS | 325 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 1099 us | 964 TFLOPS | 594 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1921 us | 990 TFLOPS | 385 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 920 us | 996 TFLOPS | 686 GB/s ``` after: ``` Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1490 us | 1397 TFLOPS | 450 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 801 us | 1323 TFLOPS | 815 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1370 us | 1388 TFLOPS | 540 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 699 us | 1312 TFLOPS | 903 GB/s ``` @RayWang96 Could you please tell us what's the point of `For grouped GEMMs, let weights always stay in the L2 cache and read activations by once`?
bsakaguchi commented 2026-10-26 21:12:23 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 20, 2025):

Interesting, DeepSeek said they could reproduce the good performance on hopper using latest main (i.e. I think without code changes). iirc in their blog they say also use H800

<!-- gh-comment-id:3204145779 --> @fzyzcjy commented on GitHub (Aug 20, 2025): Interesting, DeepSeek said they could reproduce the good performance on hopper using latest main (i.e. I think without code changes). iirc in their blog they say also use H800
bsakaguchi commented 2026-10-26 21:12:23 +00:00
Author
Owner
Copy link

@coconut902 commented on GitHub (Aug 20, 2025):

Interesting, DeepSeek said they could reproduce the good performance on hopper using latest main (i.e. I think without code changes). iirc in their blog they say also use H800

Did you mean commit f85ec649d7 (0815)?
Performance of commit 3254b758e2 (0814) is still good, though slightly worse than v1. But commit f85ec649d7 (0815) is much worse.

<!-- gh-comment-id:3204235372 --> @coconut902 commented on GitHub (Aug 20, 2025): > Interesting, DeepSeek said they could reproduce the good performance on hopper using latest main (i.e. I think without code changes). iirc in their blog they say also use H800 Did you mean commit f85ec649d76846552cfd637b6c99fb9c985fd9eb (0815)? Performance of commit 3254b758e27a5b2f2ae68279314acb1adcb6c1bc (0814) is still good, though slightly worse than v1. But commit f85ec649d76846552cfd637b6c99fb9c985fd9eb (0815) is much worse.
bsakaguchi commented 2026-10-26 21:12:24 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 20, 2025):

Yes I test via latest commit. Let me test 0814 commit later.

<!-- gh-comment-id:3204276174 --> @fzyzcjy commented on GitHub (Aug 20, 2025): Yes I test via latest commit. Let me test 0814 commit later.
bsakaguchi commented 2026-10-26 21:12:24 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 20, 2025):

By changing my reproduction above to 0814 commit

(cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && git checkout 3254b758e27a5b2f2ae68279314acb1adcb6c1bc && ./install.sh)

The speed is better now, though it is still worse than v1.

root@db3dc888a553:/# (cd /tmp/DeepGEMM/tests && python3 test_fp8.py)
/usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.)
  cpu = _conversion_method_template(device=torch.device("cpu"))
Library path:
 > ['/usr/local/lib/python3.12/dist-packages/deep_gemm']

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1692 us | 1230 TFLOPS |  396 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT):  893 us | 1187 TFLOPS |  731 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1568 us | 1213 TFLOPS |  472 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  781 us | 1173 TFLOPS |  808 GB/s
<!-- gh-comment-id:3204295423 --> @fzyzcjy commented on GitHub (Aug 20, 2025): By changing my reproduction above to 0814 commit ``` (cd /tmp && rm -rf DeepGEMM && git clone --recursive https://github.com/deepseek-ai/DeepGEMM && cd DeepGEMM && git checkout 3254b758e27a5b2f2ae68279314acb1adcb6c1bc && ./install.sh) ``` The speed is better now, though it is still worse than v1. ``` root@db3dc888a553:/# (cd /tmp/DeepGEMM/tests && python3 test_fp8.py) /usr/local/lib/python3.12/dist-packages/torch/_subclasses/functional_tensor.py:279: UserWarning: Failed to initialize NumPy: No module named 'numpy' (Triggered internally at /pytorch/torch/csrc/utils/tensor_numpy.cpp:81.) cpu = _conversion_method_template(device=torch.device("cpu")) Library path: > ['/usr/local/lib/python3.12/dist-packages/deep_gemm'] Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1692 us | 1230 TFLOPS | 396 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 893 us | 1187 TFLOPS | 731 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1568 us | 1213 TFLOPS | 472 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 781 us | 1173 TFLOPS | 808 GB/s ```
bsakaguchi commented 2026-10-26 21:12:24 +00:00
Author
Owner
Copy link

@coconut902 commented on GitHub (Aug 20, 2025):

The bad performance of 0814 commit is due to non-optimal configs, issue #161

On my H800,
before fixing issue #161

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1687 us | 1234 TFLOPS |  397 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT):  887 us | 1195 TFLOPS |  736 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1569 us | 1212 TFLOPS |  471 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  778 us | 1178 TFLOPS |  811 GB/s

after:

Testing m-grouped contiguous GEMM:
 > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1522 us | 1368 TFLOPS |  440 GB/s
 > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT):  830 us | 1276 TFLOPS |  786 GB/s
 > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1398 us | 1360 TFLOPS |  529 GB/s
 > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT):  728 us | 1259 TFLOPS |  867 GB/s

0815 commit fixed the issue, but introduced a new problem.

<!-- gh-comment-id:3204330238 --> @coconut902 commented on GitHub (Aug 20, 2025): The bad performance of 0814 commit is due to non-optimal configs, issue #161 On my H800, before fixing issue #161 ``` Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1687 us | 1234 TFLOPS | 397 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 887 us | 1195 TFLOPS | 736 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1569 us | 1212 TFLOPS | 471 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 778 us | 1178 TFLOPS | 811 GB/s ``` after: ``` Testing m-grouped contiguous GEMM: > Perf (num_groups=4, m=35456, n= 4096, k= 7168, 1D2D, layout=NT): 1522 us | 1368 TFLOPS | 440 GB/s > Perf (num_groups=4, m=36096, n= 7168, k= 2048, 1D2D, layout=NT): 830 us | 1276 TFLOPS | 786 GB/s > Perf (num_groups=8, m=32384, n= 4096, k= 7168, 1D2D, layout=NT): 1398 us | 1360 TFLOPS | 529 GB/s > Perf (num_groups=8, m=31232, n= 7168, k= 2048, 1D2D, layout=NT): 728 us | 1259 TFLOPS | 867 GB/s ``` 0815 commit fixed the issue, but introduced a new problem.
bsakaguchi commented 2026-10-26 21:12:25 +00:00
Author
Owner
Copy link

@fzyzcjy commented on GitHub (Aug 20, 2025):

I see, looks pretty reasonable. Maybe I will wait a bit to see whether this will be fixed soon. cc @LyricZhao

<!-- gh-comment-id:3204348078 --> @fzyzcjy commented on GitHub (Aug 20, 2025): I see, looks pretty reasonable. Maybe I will wait a bit to see whether this will be fixed soon. cc @LyricZhao
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
public-sync-20261014
0930-release
nv_dev
fix-mega-moe-fence
public-release-2609
public-release-2607
public-release-2606
mega-update
public-release-260416
ko3n1g/ci/build-wheels
nv_dev_f8e8fb5
nv_dev_950d31a
nv_dev_67fc648
nv_dev_c491439
nv_dev_4ff3f54
v2.1.1.post3
v2.1.1.post2
v2.1.1.post1
v2.1.1
v2.1.0
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepGEMM#92
No description provided.