Skip to content

[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm#1862

Merged
ThomasNing merged 123 commits into
developfrom
update_cka8w8_uc
Feb 20, 2025
Merged

[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm#1862
ThomasNing merged 123 commits into
developfrom
update_cka8w8_uc

Conversation

@aska-0096
Copy link
Copy Markdown
Contributor

@aska-0096 aska-0096 commented Feb 5, 2025

Proposed changes

Please describe the motivation behind the pull request, whether it enables a new feature or fixes a bug. If there are associated pull requests or issues, please link them to the pull request.

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

aska-0096 and others added 30 commits August 22, 2024 07:30
@aska-0096
Copy link
Copy Markdown
Contributor Author

Hi @illsilin
CI passed.

Comment thread example/65_gemm_multiply_multiply/CMakeLists.txt Outdated
Comment thread example/65_gemm_multiply_multiply/CMakeLists.txt Outdated
Comment thread example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp8_bpreshuffle.cpp Outdated
Comment thread profiler/include/profiler/profile_gemm_multiply_multiply_impl.hpp Outdated
Comment on lines +274 to +275
// hipStream_t stream;
// hip_check_error(hipStreamCreate(&stream));
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Was this code meant to be uncommented?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

This code is a kind of explanatory note to whoever wondering why we use null hipStream in the ckProfiler kernel launch,

@ThomasNing
Copy link
Copy Markdown
Contributor

Already add the 16x16 instances to the CK Profiler for pre-shuffle.

Copy link
Copy Markdown
Collaborator

@aosewski aosewski left a comment

Choose a reason for hiding this comment

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

Are the changes to rmsnorm2d_rdquant relevant to this PR or is it a bad merge artifact?

@ThomasNing ThomasNing changed the title [A8W8 GEMM] Optimized weight-preshuffled implementation [A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm Feb 20, 2025
@ThomasNing
Copy link
Copy Markdown
Contributor

@aosewski It is a feature adding in this PR to add the quant data type for rmsnorm2d_rdquant ck_tile operator.

@ThomasNing ThomasNing self-assigned this Feb 20, 2025
@ThomasNing ThomasNing dismissed aosewski’s stale review February 20, 2025 18:16

Replied to Adam's question.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.