feat(quant_grouped_gemm): add bquant with preshuffleB support to grouped_gemm example #3115
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Co-authored-by: Copilot <[email protected]>
Co-authored-by: Copilot <[email protected]>
Calculate has_hot_loop, num_loop, and tail_number on device side for each GEMM problem instead of using default values. This fixes incorrect results when different problems in the group have different K dimensions.
AviralGoelAMD
requested review from
ThomasNing,
afagaj,
andriy-ca,
aosewski,
asleepzzz,
bartekxk,
carlushuang,
cgmillette,
coderfeli,
geyyer,
illsilin,
poyenc,
qianfengz,
shumway,
tenpercent and
vidyasagar-amd
as code owners
7 tasks
There was a problem hiding this comment.
Pull Request Overview
This PR adds support for double shared memory buffer (DoubleSmemBuffer) with weight pre-shuffling (PreshuffleB) for BQuantGrouped quantization mode in grouped GEMM operations.
- Implements a new overload of the pipeline operator to support double shared memory buffers with TailNumber parameter
- Adds conditional logic to allocate and use two shared memory buffers when both DoubleSmemBuffer and BQuantGrouped are enabled
- Introduces a new configuration
GemmConfigPreshuffleB_Bquant_prefillwith PreshuffleB and DoubleSmemBuffer enabled
Reviewed Changes
Copilot reviewed 7 out of 7 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp | Adds new operator overload accepting TailNumber parameter for double buffer support |
| include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp | Implements conditional double shared memory buffer allocation and adds RunGemmWithPipelineSelection2LDS method |
| include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp | Minor formatting changes (blank lines added) |
| include/ck_tile/ops/gemm/pipeline/wp_pipeline_agmem_bgmem_creg_v2.hpp | Changes BlockHasHotloop from host-only to host-device callable |
| example/ck_tile/17_grouped_gemm/quant_run_grouped_gemm_example.inc | Adds support for BQuantGrouped stride configuration, pre-shuffling B tensor, and initialization method parameter |
| example/ck_tile/17_grouped_gemm/quant_grouped_gemm.hpp | Adds new configuration and removes pipeline-related macros, adds get_k_from_preshuffled_warp_tile function |
| example/ck_tile/17_grouped_gemm/quant_grouped_gemm.cpp | Updates pipeline selection logic and switches main to use new PreshuffleB configuration |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
ThomasNing
requested changes
Oct 30, 2025
| bq_dev_buf.push_back( | ||
| std::make_unique<ck_tile::DeviceMem>(bq_tensors[i].get_element_space_size_in_bytes())); | ||
|
|
||
| if constexpr(GemmConfig::PreshuffleB && QuantMode == ck_tile::QuantType::BQuantGrouped) |
There was a problem hiding this comment.
If it is preshuffle B, even it is not BQuant we should still shuffle the B.
There was a problem hiding this comment.
Based on my quick testing, gemm_quant_basic example does not support PreshuffleB == true when quant_mode=rowcol | tensor.
Hence, in quant_grouped_gemm, we consider it unsupported for other quant_mode(s)
ThomasNing
approved these changes
Oct 30, 2025
amd-khushbu
approved these changes
Oct 30, 2025
|
Close as the #3119 has been merged to the develop |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Proposed changes
This PR adds bquant support to grouped_gemm examples
(example/ck_tile/17_grouped_gemm/quant_grouped_gemm.cpp)and kernel(include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp).Checklist
Please put an
xinto the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.clang-formaton all changed filesDiscussion
If this is a relatively large or complex change, feel free to start a discussion by explaining why you chose the solution you did and what alternatives you considered