-
Notifications
You must be signed in to change notification settings - Fork 752
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
[SYCL][E2E] Update get_coordinate tests to verify both row and columns #15880
[SYCL][E2E] Update get_coordinate tests to verify both row and columns #15880
Conversation
This PR is still in progress. There might be some code clean up needed to be done. |
@YixingZhang007 I've updated the PR description to remove the JIRA link. The policy does not allow using internal JIRA links in open-source, only JIRA ticket numbers are allowed. |
f613e96
to
fc048c9
Compare
b034c02
to
be1187a
Compare
2f361b6
to
fc18698
Compare
|
||
// This condition check can be removed once the IGC PR resolving the Matrix B | ||
// row coordinate bug is pull downed to the driver. | ||
if (Use != use::b) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
this is not needed anymore
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
Ping @intel/llvm-reviewers-runtime |
@intel/llvm-gatekeepers This PR has been approved and all CI tests passed. Could you please help merge this PR? Thanks! |
for (int i = 0; i < NUM_ROWS; i++) { | ||
sum_local_rows[i] = | ||
reduce_over_group(sg, sum_local_rows[i], sycl::plus<>()); | ||
// only Groups leader perform the global reduction | ||
if (global_idy % sg_size == 0) { | ||
sycl::atomic_ref<TResult, sycl::memory_order::relaxed, | ||
sycl::memory_scope::device> | ||
aref(v_rows[i]); | ||
aref.fetch_add(sum_local_rows[i]); | ||
} | ||
} | ||
|
||
for (int i = 0; i < NUM_COLS; i++) { | ||
sum_local_cols[i] = | ||
reduce_over_group(sg, sum_local_cols[i], sycl::plus<>()); | ||
// only Groups leader perform the global reduction | ||
if (global_idy % sg_size == 0) { | ||
sycl::atomic_ref<TResult, sycl::memory_order::relaxed, | ||
sycl::memory_scope::device> | ||
aref(v_cols[i]); | ||
aref.fetch_add(sum_local_cols[i]); | ||
} | ||
} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
@YixingZhang007 this code is duplicated. Would be great to outline it to function and reuse.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks for the suggestion! I have created a PR for fixing this. The link to the PR is #16429
for (int i = 0; i < Rows; i++) { | ||
for (int j = 0; j < Cols; j++) { | ||
M[i][j] = i + j; | ||
} | ||
} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
@YixingZhang007 , we have matrix_fill
in common.hpp
to not duplicate such loops in each test.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks! I have changed this.
{1, 1 * sg_size}), | ||
[=](nd_item<2> spmd_item) | ||
#ifdef SG_SZ | ||
[[intel::reqd_sub_group_size(SG_SZ)]] |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
attribute 'intel::reqd_sub_group_size' is deprecated
sycl::reqd_sub_group_size should be used
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Sorry for the mistake! I have fix this in the other PR.
Improving Test Coverage and Cleanliness:
get_coordinate
function to check both row and column indices for all matrix types inget_coordinate_ops.cpp
.get_coord_int8_matA.cpp
,get_coord_int8_matB.cpp
, andget_coord_float_matC.cpp
. These tests have been consolidated and are now included inget_coordinate_ops.cpp
.