deepseek-ai/DeepGEMM

Fork 0

mirror of https://github.com/deepseek-ai/DeepGEMM.git synced 2026-11-05 00:47:26 +00:00

[PR #272] [CLOSED] Add FP8 per-tensor scaling to cuBLASLt for fair comparison #334

New issue

Closed

opened 2026-10-26 21:16:24 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 21:16:24 +00:00

Owner

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepGEMM/pull/272
Author: @yurekami
Created: 1/1/2026
Status: ❌ Closed

Base: main ← Head: fix/cublaslt-fp8-blockwise-scaling

📝 Commits (1)

5753b21 Add FP8 per-tensor scaling support to cuBLASLt for fair comparison

📊 Changes

4 files changed (+59 additions, -5 deletions)

View changed files

📝 csrc/apis/gemm.hpp (+35 -0)
📝 csrc/jit_kernels/impls/smxx_cublaslt.hpp (+18 -4)
📝 deep_gemm/__init__.py (+2 -0)
📝 tests/test_fp8.py (+4 -1)

📄 Description

Summary

This PR addresses #199 where the cuBLASLt comparison in test_fp8.py uses tensorwise scaling (no scaling) while DeepGEMM uses blockwise scaling, making the performance comparison misleading.

Changes

Add scale_a and scale_b optional parameters to call_cublaslt_api()
Set CUBLASLT_MATMUL_DESC_A_SCALE_POINTER and B_SCALE_POINTER when provided
Add new cublaslt_fp8_gemm_nt() function that accepts scale tensors
Export the new function via pybind11 and Python __init__.py
Update test_fp8.py to use per-tensor scaling (max of blockwise scales)

How it works

The test now computes the maximum value from the blockwise scales and uses it as a single per-tensor scale for cuBLASLt:

scale_a_max = a[1].max().reshape(1).to(torch.float32).contiguous()
scale_b_max = b[1].max().reshape(1).to(torch.float32).contiguous()
cublas_t, split_k_t = bench_kineto(
    lambda: deep_gemm.cublaslt_fp8_gemm_nt(a[0], b[0], d, c=c, scale_a=scale_a_max, scale_b=scale_b_max),
    ('nvjet', 'reduce'), suppress_kineto_output=True
)

This provides a fairer comparison where both implementations apply some form of FP8 dequantization.

Limitations

This is an approximation - cuBLASLt uses per-tensor scaling while DeepGEMM uses more precise blockwise (128x128) scaling. For a fully aligned comparison, cuBLASLt would need blockwise scaling support using CUBLASLT_MATMUL_SCALE_MODE_BLOCK_2D (available in CUDA 12.0+). This could be a future enhancement.

Test Plan

Verify test_fp8.py runs without errors
Verify cuBLASLt now uses scaling in the benchmark comparison
Compare performance numbers before and after

Fixes #199

🤖 Generated with Claude Code

_{🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.}

## 📋 Pull Request Information **Original PR:** https://github.com/deepseek-ai/DeepGEMM/pull/272 **Author:** [@yurekami](https://github.com/yurekami) **Created:** 1/1/2026 **Status:** ❌ Closed **Base:** `main` ← **Head:** `fix/cublaslt-fp8-blockwise-scaling` --- ### 📝 Commits (1) - [`5753b21`](https://github.com/deepseek-ai/DeepGEMM/commit/5753b21f4937e2d05502a6cce88e5ceb49628dcd) Add FP8 per-tensor scaling support to cuBLASLt for fair comparison ### 📊 Changes **4 files changed** (+59 additions, -5 deletions) <details> <summary>View changed files</summary> 📝 `csrc/apis/gemm.hpp` (+35 -0) 📝 `csrc/jit_kernels/impls/smxx_cublaslt.hpp` (+18 -4) 📝 `deep_gemm/__init__.py` (+2 -0) 📝 `tests/test_fp8.py` (+4 -1) </details> ### 📄 Description ## Summary This PR addresses #199 where the cuBLASLt comparison in `test_fp8.py` uses tensorwise scaling (no scaling) while DeepGEMM uses blockwise scaling, making the performance comparison misleading. ## Changes - Add `scale_a` and `scale_b` optional parameters to `call_cublaslt_api()` - Set `CUBLASLT_MATMUL_DESC_A_SCALE_POINTER` and `B_SCALE_POINTER` when provided - Add new `cublaslt_fp8_gemm_nt()` function that accepts scale tensors - Export the new function via pybind11 and Python `__init__.py` - Update `test_fp8.py` to use per-tensor scaling (max of blockwise scales) ## How it works The test now computes the maximum value from the blockwise scales and uses it as a single per-tensor scale for cuBLASLt: ```python scale_a_max = a[1].max().reshape(1).to(torch.float32).contiguous() scale_b_max = b[1].max().reshape(1).to(torch.float32).contiguous() cublas_t, split_k_t = bench_kineto( lambda: deep_gemm.cublaslt_fp8_gemm_nt(a[0], b[0], d, c=c, scale_a=scale_a_max, scale_b=scale_b_max), ('nvjet', 'reduce'), suppress_kineto_output=True ) ``` This provides a fairer comparison where both implementations apply some form of FP8 dequantization. ## Limitations This is an approximation - cuBLASLt uses per-tensor scaling while DeepGEMM uses more precise blockwise (128x128) scaling. For a fully aligned comparison, cuBLASLt would need blockwise scaling support using `CUBLASLT_MATMUL_SCALE_MODE_BLOCK_2D` (available in CUDA 12.0+). This could be a future enhancement. ## Test Plan - [ ] Verify `test_fp8.py` runs without errors - [ ] Verify cuBLASLt now uses scaling in the benchmark comparison - [ ] Compare performance numbers before and after Fixes #199 🤖 Generated with [Claude Code](https://claude.com/claude-code) --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi

2026-10-26 21:16:24 +00:00

closed this issue
added the
pull-request
label

No labels

pull-request

No milestone

No project

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#334

No description provided.

Rows
Columns