Skip to content

[CUDA] fp and int4 quants for qmm_sm80#3268

Merged
zcbenz merged 1 commit intoml-explore:mainfrom
zcbenz:qmm-sm80-update
Mar 19, 2026
Merged

[CUDA] fp and int4 quants for qmm_sm80#3268
zcbenz merged 1 commit intoml-explore:mainfrom
zcbenz:qmm-sm80-update

Conversation

@zcbenz
Copy link
Copy Markdown
Collaborator

@zcbenz zcbenz commented Mar 17, 2026

Followup to #3255 finishing the missing quants, was blocked by NVIDIA/cutlass#3111.

We can't support int2/3/5/6 in this version unfortunately since the lengths supported by cp.async instruction are limited.

Copy link
Copy Markdown
Member

@angeloskath angeloskath left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks great!

How 's the initial perf looking?

@zcbenz
Copy link
Copy Markdown
Collaborator Author

zcbenz commented Mar 18, 2026

The memory bandwidth of int4 is about 50% of int8 on A100:

M N K QMM (GiB/s) CUBLAS (GiB/s) QMM (TF/s) CUBLAS (TF/s) Speedup (x)
16 4096 4096 174.3 1397.5 9.65 22.19 0.43

I think it means that the kernel is bottlenecked at loading quantized weights.

@zcbenz zcbenz merged commit dbfbc0f into ml-explore:main Mar 19, 2026
16 checks passed
@zcbenz zcbenz deleted the qmm-sm80-update branch March 19, 2026 00:38
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants