Skip to content

Ck tile/gemm blockscale opt - #3227

Merged
ThomasNing merged 5 commits into
developfrom
ck_tile/gemm_blockscale_opt
Dec 4, 2025
Merged

Ck tile/gemm blockscale opt#3227
ThomasNing merged 5 commits into
developfrom
ck_tile/gemm_blockscale_opt

Conversation

@kensclin

@kensclin kensclin commented Nov 18, 2025

Copy link
Copy Markdown
Contributor

Proposed changes

Performance profiling of the GEMM BlockScale kernel revealed a major bottleneck caused by inefficient instruction scheduling, leading to excessive pipeline bubbles/stalls and subsequent performance degradation.

Testing

./tile_example_gemm_quant -m=32768 -n=4096 -k=2048 -quant_mode=bquant -preshuffleb=true -v=1

The result improved from 769us -> 619us.
The improvement rate is approximately 20%.

Checklist

Please put an x into 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.

  • I have added tests relevant to the introduced functionality, and the unit tests are passing locally
  • I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, IF the test takes more than 30 seconds to run.
  • I have added inline documentation which enables the maintainers with understanding the motivation
  • I have removed the stale documentation which is no longer relevant after this pull request
  • (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request
  • I have run clang-format on all changed files
  • Any dependent changes have been merged

Discussion

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

@ThomasNing

Copy link
Copy Markdown
Contributor

@kensclin Please solve the compilation error and the merge conflict.

@kensclin
kensclin force-pushed the ck_tile/gemm_blockscale_opt branch 4 times, most recently from e4bf959 to 689ad32 Compare November 25, 2025 02:36
@kensclin
kensclin force-pushed the ck_tile/gemm_blockscale_opt branch from 689ad32 to 818b23f Compare November 25, 2025 03:09
@ThomasNing

Copy link
Copy Markdown
Contributor

@kensclin Plz solve the current CI errors.

Comment thread include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp Outdated
@ThomasNing
ThomasNing merged commit ffc3120 into develop Dec 4, 2025
25 checks passed
@ThomasNing
ThomasNing deleted the ck_tile/gemm_blockscale_opt branch December 4, 2025 06:07
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