Refactored blas1 dot & sdsdot operators #471
Merged
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
OuadiElfarouki
requested review from
s-Nick,
muhammad-tanvir-1211 and
pgorlani
OuadiElfarouki
force-pushed
the
dot_refacto_pr
branch
2 times, most recently
from
f5887ca
to
a134cb4
Compare
s-Nick
reviewed
Oct 18, 2023
s-Nick
approved these changes
Oct 20, 2023
OuadiElfarouki
force-pushed
the
dot_refacto_pr
branch
4 times, most recently
from
6c8cb3f
to
2e77adf
Compare
OuadiElfarouki
force-pushed
the
dot_refacto_pr
branch
from
2e77adf
to
5fb556e
Compare
muhammad-tanvir-1211
approved these changes
Nov 16, 2023
muhammad-tanvir-1211
merged commit ece6336
into
codeplaysoftware:master
3 checks passed
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.
This PR follows the refactoring approach of reduction based operators (asum, nrm2) for the dot (& sdsdot), namely the reduction through the
sycl::reduce_over_groupencapsulated in the blas::WGAtomicReduction kernel.Also, per-target device tuned wg size & local memory configurations are added for the dot operator to be selected depending on tuning target (at compile time) & input size (at run-time).
The changes affect the
dotandsdsdotasync operators directly and their synchronous (with output return) versions indirectly. The async version of the two operators API expect the input/output to be initialized to zero, while the sync version handles the initialization to zero internally.