Skip to content

feat(cuda): add ldmatrix integer dequantization - #139

Merged
jinzhen-lin merged 9 commits into
mainfrom
feat/flat-ldmatrix-integer-dequant
Oct 13, 2026
Merged

jinzhen-lin merged 9 commits into
mainfrom
feat/flat-ldmatrix-integer-dequant

Conversation

@jinzhen-lin

Copy link
Copy Markdown
Collaborator

Stacked on #137.

  1. CUDA 13.4 adds an ldmatrix instruction that dequantizes INT4 to INT8 while loading data from shared memory. This PR adds full support for this operation in W4A8 GEMM. Thanks to @LilySu for pointing it out in Ptx94 ldmatrix s8 s4 dense w4a8 #85.

  2. Analysis and hardware testing revealed broader support than the PTX documentation describes: all architectures supporting ldmatrix (SM75+) can execute the conversion through patched SASS, and the instruction also supports INT2-to-INT4 conversion. This PR reuses patch_cubin, compiling ldmatrix.sync.aligned.m8n8.{x2,x4}.trans.shared.b16 placeholders and then patching the generated SASS. This enables W4A8 and W2A4 conversion on supported hardware, including with CUDA versions before 13.4. Native PTX is used for INT4-to-INT8 on SM90+ with CUDA 13.4+.

  3. Adds HUMMING_DISABLE_PATCH_CUBIN=1 to disable patched instructions and undocumented hardware formats, including Blackwell E3M4/E0M3 support. Integer dequantization falls back to software when native PTX is unavailable; documented native instructions remain enabled.

N = K = 8192, per-token activation × per-channel weight, default heuristics. Values are speedups over software dequantization, measured with alternating CUDA graph runs.

W4A8

M SM75 SM80 SM86 SM89 SM90 H20 SM90 H200 SM120 SM121
1 1.01x 1.16x 1.00x 1.07x 1.28x 1.11x 1.08x 1.00x
16 1.01x 1.13x 1.00x 1.04x 1.06x 1.14x 1.02x 1.00x
64 1.00x 1.11x 1.07x 1.07x 1.02x 1.15x 1.06x 1.00x
256 1.05x 1.13x 1.01x 1.04x 1.00x 1.10x 1.02x 0.97x
1024 1.05x 1.18x 1.06x 1.14x 1.00x 1.06x 1.01x 1.01x
4096 1.07x 1.10x 1.02x 1.09x 1.00x 1.06x 1.01x 1.00x

W2A4

M SM75 SM80 SM86 SM89
1 1.04x 1.16x 1.01x 1.06x
16 1.03x 1.15x 1.01x 1.04x
64 0.98x 1.09x 1.03x 1.07x
256 1.06x 1.07x 1.24x 1.05x
1024 1.08x 1.28x 1.14x 1.16x
4096 1.08x 1.12x 1.10x 1.07x

jinzhen-lin and others added 8 commits October 12, 2026 21:03
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
@jinzhen-lin
jinzhen-lin added this pull request to stack #140 October 12, 2026 17:42
Base automatically changed from feat/flat-native-weights to main October 13, 2026 07:26
Co-authored-by: Codex (GPT-6) <codex@openai.com>
Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com>
@jinzhen-lin
jinzhen-lin marked this pull request as ready for review October 13, 2026 09:12
@jinzhen-lin
jinzhen-lin requested a review from mgoin as a code owner October 13, 2026 09:12
@jinzhen-lin
jinzhen-lin merged commit f35c5ff into main Oct 13, 2026
5 checks passed
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.

1 participant