DeviceGemm_Wmma_CShuffleV3 with BlockGemmPipelineVersion::v3 #2096
+5,324
−568
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
They are based on files used for DeviceGemm_Xdl_CShuffleV3: * only parts needed for BlockGemmPipelineVersion::v3 are included; * instances only for one type and one layout; Co-authored-by: Anca Hamuraru <[email protected]>
* instances only for f16_mk_kn_mn are added; * only gfx11 is supported; * SplitK is disabled for now, it needs some additional work because of atomics.
gfx12 is not tested yet.
Based on e5ad48a
ex-rzr
force-pushed
the
agorenko/devicegemm_wmma_cshufflev3
branch
from
4da42af to
ed515b7
Compare
ex-rzr
changed the title
ex-rzr
requested review from
illsilin,
carlushuang,
qianfengz,
aosewski,
poyenc,
geyyer,
bartekxk,
andriy-ca,
afagaj and
asleepzzz
as code owners
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
DeviceGemm_Wmma_CShuffleV3(DeviceGemmV2)BlockwiseGemmWmmaops_pipeline_v3(BlockGemmPipelineVersion::v3) is implemented, other pipelines will be added in next PRs.HotLoopScheduleris commented as it may require changes for correct__builtin_amdgcn_sched_group_barriers due to differences between mfma and wmma (layouts etc.);BlockGemmPipelineVersion::v3is used everywhere for testing different configs, after adding other pipelines some instances (e.g. latency friendly) will use them;GemmSpecialization::Defaultis instantiated, mainly for faster building during development, specializations with padding work and will be added later.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