-
Notifications
You must be signed in to change notification settings - Fork 787
[SYCL][HIP] Support of AMD matrix core instructions #11485
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
Merged
Merged
Changes from all commits
Commits
Show all changes
56 commits
Select commit
Hold shift + click to select a range
b68aead
[Matrix] syntax changes as prepraration before moving joint matrix from
yubingex007-a11y 5fbb285
clang-format
yubingex007-a11y bf6cd56
fix typo: dest->dst
yubingex007-a11y b399041
fix testcase
yubingex007-a11y dae1ec6
fix mad bug
yubingex007-a11y 4ec8360
fix cuda const joint_matrix_cuda
yubingex007-a11y a461cbb
fix const issue of jm_store_cuda
yubingex007-a11y 5ff715b
fix const
yubingex007-a11y 8ad7da9
lint
yubingex007-a11y 26ea49d
address dounia's comments and roll back all the testcase changes
yubingex007-a11y a09a778
test changes: mov D in mad
yubingex007-a11y 821fa89
testcase changes: ext_intel_layout
yubingex007-a11y a3921b5
testcase changes: wi_data=>jm_apply
yubingex007-a11y ef1bc67
lint
yubingex007-a11y f395199
Merge remote-tracking branch 'intel_llvm/sycl' into jm_syntax
yubingex007-a11y c71fee6
Merge remote-tracking branch 'intel_llvm/sycl' into jm_syntax
yubingex007-a11y 8f2f197
handle cuda testcase compfail
yubingex007-a11y 1411376
address dounia's comments
yubingex007-a11y 95df3b1
lint
yubingex007-a11y fb1afdc
rm sycl/test/matrix/query-use.cpp
yubingex007-a11y 11df531
fix x jm_mad in joint_matrix_bf16_fill_k_cache_impl.hpp
yubingex007-a11y a29e8f3
Merge remote-tracking branch 'intel_llvm/sycl' into jm_syntax
yubingex007-a11y a821107
address comments
yubingex007-a11y 3f1b575
Merge remote-tracking branch 'intel_llvm/sycl' into jm_syntax
yubingex007-a11y 1d091de
rm element_wise_irreg_sum_rows_impl.hpp
yubingex007-a11y 1e20968
small fix
yubingex007-a11y 1fe7fcd
small fix
yubingex007-a11y b5c0911
* Support one block AMD matrix core instructions
f6b2cb3
* Update matrix core support into joint_matrix documentation.
919884b
- Fix adding `gpu-amd-gfx90a` as available feature.
02bec23
Fix a missing variable name change.
3c460af
Add decoration type for call to get_multi_ptr function.
42e0c62
Update use cases of `mad` to have variables holding result of `mad` a…
be7a90f
Implement joint_matrix_copy for HIP(gfx90a) backend.
3d1237a
std::ignore unused input parameters.
ad7b8cd
std::ignore unused `sg` parameters in`joint_matrix_store`.
1d46e33
Merge branch 'sycl' into amd-joint-matrix
4231a1c
Fix AMD `joint_matrix_copy` function.
1a59580
- Fix joint_matrix_hip_copy.
fc31965
Remove curly braces for initialization of joint_matrix.
8bba0fb
- Use sycl::marray as container for jont_matrix data.
a152eda
Modify `joint_matrix_apply` test.
b44a76c
Update allow difference after using matrix random input.
66131d7
Fix `hip_matrix_copy`.
f719779
- Improve hip mfma tests to support matrices of multiple of K size.
7aa4ce3
Fix call to `copy` and `fill` for hip joint matrix.
3d13484
Fix reference mma calculation.
af5cc07
Rename `cuda_impl` and `hip_impl` member of `joint_matrix` to `matrix…
91cbab7
Merge branch 'sycl' into amd-joint-matrix
c6a3cee
Replace `cuda_impl` with `matrix_impl`
2f6885f
Rename `data` in `joint_matrix_hip` with `wi_marray`.
defd874
Improve `joint_matrix_copy` by avoiding the loop.
89c52d7
Add a missing `comma` to the test.
abfa2ab
Remove `-DSYCL_EXT_ONEAPI_MATRIX_VERSION=4` from AMD matrix compilati…
048ac39
Guard including `matrix-hip.hpp`
05d2e9d
Minor macro readability improvement.
File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
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
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
Oops, something went wrong.
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.
Uh oh!
There was an error while loading. Please reload this page.