Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

CK Tile Gemm API and heuristics changes #1809

Open
wants to merge 4 commits into
base: develop
Choose a base branch
from

Conversation

jakpiase
Copy link
Contributor

Proposed changes

Add api for CK TILE Gemm similar to FMHA's one and add instances

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 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

Add api for CK Tile universal gemm and some instances for fp16 and bf16 cases. I'm not sure how the heuristic of choosing between comp and mem instances should be made so for now it is trivial and will be upgraded in the near future

@jakpiase jakpiase changed the title Jakpiase/ck tile gemm api CK Tile Gemm API and heuristics changes Jan 14, 2025
Copy link
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.

Very good start!

example/ck_tile/03_gemm/CMakeLists.txt Outdated Show resolved Hide resolved
example/ck_tile/03_gemm/CMakeLists.txt Outdated Show resolved Hide resolved
example/ck_tile/03_gemm/CMakeLists.txt Show resolved Hide resolved
example/ck_tile/03_gemm/gemm_basic.cpp Outdated Show resolved Hide resolved
example/ck_tile/03_gemm/gemm_basic.hpp Outdated Show resolved Hide resolved
example/ck_tile/03_gemm/instances/gemm_api.cpp Outdated Show resolved Hide resolved
example/ck_tile/03_gemm/run_gemm_example.inc Show resolved Hide resolved
example/ck_tile/03_gemm/run_gemm_example.inc Show resolved Hide resolved
example/ck_tile/03_gemm/universal_gemm.cpp Outdated Show resolved Hide resolved
Comment on lines +229 to +236
else if(a_layout == "C" && b_layout == "C")
{
return run_gemm_example_with_layouts(argc, argv, Col{}, Col{}, Row{});
}
else if(a_layout == "C" && b_layout == "R")
{
return run_gemm_example_with_layouts(argc, argv, Col{}, Row{}, Row{});
}
Copy link
Collaborator

Choose a reason for hiding this comment

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

Does it work now?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

CCR and CRR don't work with the same tiles as RRR and RCR do, but with they seem to work with some other tile shapes, I am not sure if that's the desired behavior or if something is broken inside of them

bool is_a_rowmajor;
bool is_b_rowmajor;
bool is_c_rowmajor;
std::string data_type; /** Tensors datatype, can be set to either fp16 or bf16*/
Copy link
Collaborator

Choose a reason for hiding this comment

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

fp8 should also be supported.

@@ -2,6 +2,29 @@
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
Copy link
Collaborator

Choose a reason for hiding this comment

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

Update copyright year.

Comment on lines +21 to +39
return gemm_<gemm_traits_<FP16,
FP16,
FP32,
FP16,
Row,
Row,
Row,
256,
256,
32,
2,
2,
1,
32,
32,
16,
false,
false,
false>>(a, s);
Copy link
Collaborator

Choose a reason for hiding this comment

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

Please turn off clang format also here. It would be beneficial to have // comment, as we used to have in old ck instances, describing what specific value referes to.

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