This repository has been archived by the owner on Jan 13, 2025. It is now read-only.
New implementation of GEMM without using local memory #80
Conversation
This file contains 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
Ruyk
changed the title
Ruyk
approved these changes
Oct 5, 2018
include/operations/blas3_trees.hpp
Outdated
| "%ld\n", | ||
| id.get_global_id(0), tile_size, tile_id, tile_local_id, tiles_per_col, | ||
| tile_row, tile_col, wg_row, wg_col);*/ | ||
|
|
include/operations/blas3_trees.hpp
Outdated
| @@ -169,6 +170,303 @@ inline bool do_check<false>(bool) { | |||
| return true; | |||
| } | |||
|
|
|||
| /*! | |||
| * @brief NoLocalGemmFactory is a template class whose instantiations provide | |||
| * different implementations of the GEMM device function where the is no | |||
There was a problem hiding this comment.
of the GEMM device function -> of the GEMM kernel
include/operations/blas3_trees.hpp
Outdated
| * line size, some values fetched will be wasted, which can | ||
| * significantly reduce performance. It can be set to a | ||
| * multiple of the physical cache line size. In this case, it | ||
| * will significantly increase scratchpad memory usage, but |
include/operations/blas3_trees.hpp
Outdated
| auto A = _A.getData().get_pointer().get(); | ||
| auto B = _B.getData().get_pointer().get(); | ||
| auto C = _C.getData().get_pointer().get(); | ||
| const auto tile_per_row = ((m - 1) / block_rows) + 1; |
include/operations/blas3_trees.hpp
Outdated
| value_type reg_a[item_rows]; | ||
| value_type reg_b[item_cols]; | ||
|
|
||
| /*! @brief exiting from any threads outside of the m and n boundary */ |
There was a problem hiding this comment.
@brief is only for starting a function or a method. This is just a normal comment
include/operations/blas3_trees.hpp
Outdated
| * A, B and C matrices. | ||
| */ | ||
| int ma = (local_item_id_row + wg_row); | ||
| int na = (local_item_id_col + wg_col); |
| : store<true>(C, reg_res, alpha, beta, ldc, ma, na, boundary_check_c); | ||
| } | ||
|
|
||
| void bind(cl::sycl::handler &h) { |
There was a problem hiding this comment.
add a doxygen comment clarifying what bind does
|
Looks good to me |
AdamBrouwersHarries
approved these changes
Oct 5, 2018
Sign up for free
to subscribe to this conversation on GitHub.
Already have an account?
Sign in.
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.
No description provided.