Skip to content
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

Add test for TeamPolicy array reduction #6296

Draft
wants to merge 5 commits into
base: develop
Choose a base branch
from

Conversation

masterleinad
Copy link
Contributor

In reference to #6293 (comment). It seems that we don't test array reductions for TeamPolicy in the CI. This pull request copies a very similar test to use array reductions instead of a C-style array of fixed size.

@masterleinad masterleinad force-pushed the test_team_policy_array_reduction branch from 2f5cdc9 to b4965c5 Compare July 19, 2023 20:46
@fnrizzi
Copy link
Contributor

fnrizzi commented Jul 25, 2023

why is this PR also touching impl for CUDA and HIP too? rather than just being about a new test?

@masterleinad
Copy link
Contributor Author

why is this PR also touching impl for CUDA and HIP too? rather than just being about a new test?

Because the code path used for array reductions did not compile for CUDA and HIP.

@masterleinad masterleinad added the Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos) label Jul 26, 2023
core/unit_test/TestTeam.hpp Show resolved Hide resolved
@masterleinad
Copy link
Contributor Author

Retest this please.

1 similar comment
@masterleinad
Copy link
Contributor Author

Retest this please.

using value_type = ScalarType[];
size_type value_count = 3;

const size_type nwork;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why const?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

All this is copied from ReduceTeamFunctor and TeamFunctor (array reduction for RangePolicy).

Comment on lines 308 to 310
KOKKOS_INLINE_FUNCTION
ArrayReduceTeamFunctor(const ArrayReduceTeamFunctor &rhs)
: nwork(rhs.nwork) {}
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why define the copy constructor at all?

Comment on lines 314 to 316
dst[0] = 0;
dst[1] = 0;
dst[2] = 0;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

for (int i = 0; i < value_count; ++i) ...

kokkos_impl_cuda_shared_memory<word_size_type>() +
threadIdx.y * word_count.value);
reference_type value =
m_functor_reducer.get_reducer().init(reinterpret_cast<pointer_type>(
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

So it is not so much a "adding test", the code was broken wasn't it?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Well, yeah, it was broken for Cuda and HIP but not for the other backends AFAIK which means we didn't have a test.

}

KOKKOS_INLINE_FUNCTION
void operator()(const typename policy_type::member_type ind,
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not blocking but I am really not a fan of calling it "ind".
Also why are we not passing by ref?

Comment on lines 423 to 424
const uint64_t correct = 0 == j % 3 ? nw : nsum;
ASSERT_EQ((ScalarType)correct, result[i][j]);
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This feels silly, just inline the assertion and inject more information with a << <helpful message>;

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Backend - CUDA Backend - HIP Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants