Skip to content

Commit

Permalink
Adds Test_ViewInit and Test_MDRangePolicy
Browse files Browse the repository at this point in the history 
Applies clang-format to more files
  • Loading branch information
@janciesko
janciesko committed Nov 6, 2024
1 parent c0b4e92 commit c595c7b
Show file tree
Hide file tree
Showing 9 changed files with 278 additions and 86 deletions.
2 changes: 1 addition & 1 deletion benchmarks/access_overhead/access_overhead.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -117,7 +117,7 @@ struct Access<ViewType_t, typename std::enable_if_t<!std::is_same<
Kokkos::parallel_for("access_overhead-init", policy_init_t(0, N), *this);
Kokkos::fence();
#ifdef KRS_ENABLE_NVSHMEMSPACE
//nvshmem_barrier_all(); // Not sure why this impacts perf
// nvshmem_barrier_all(); // Not sure why this impacts perf
#endif

time_a = timer.seconds();
Expand Down
132 changes: 64 additions & 68 deletions benchmarks/access_overhead/access_overhead_p2p_kernelconf.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -76,8 +76,8 @@ struct Args_t {
int N = default_N;
int iters = default_Iters;
int rma_op = default_RmaOp;
int ts = default_ts;
int ls = default_ls;
int ts = default_ts;
int ls = default_ls;
};

void print_help() {
Expand Down Expand Up @@ -141,22 +141,22 @@ struct Access<ViewType_t, typename std::enable_if_t<
MPI_Comm_size(MPI_COMM_WORLD, &num_ranks);
other_rank = my_rank ^ 1;
assert(num_ranks == 2);
iters_per_team = args.N / LEAGUE_SIZE;
iters_per_team = args.N / LEAGUE_SIZE;
iters_per_team_mod = args.N % LEAGUE_SIZE;
};

KOKKOS_FUNCTION
void operator()(const InitTag &, const size_t i) const { v(i) = my_rank + 1; }

KOKKOS_FUNCTION
void operator()(const UpdateTag &, typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for( Kokkos::TeamThreadRange(team, start, end + mod), [&](const int i){
v(i) += v_tmp(i);
});
void operator()(const UpdateTag &,
typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for(Kokkos::TeamThreadRange(team, start, end + mod),
[&](const int i) { v(i) += v_tmp(i); });
}

KOKKOS_FUNCTION
Expand Down Expand Up @@ -189,9 +189,9 @@ struct Access<ViewType_t, typename std::enable_if_t<
MPI_Recv(v_tmp_host.data(), N, MPI_DOUBLE, other_rank, TAG,
MPI_COMM_WORLD, MPI_STATUS_IGNORE);
Kokkos::deep_copy(v_tmp, v_tmp_host);
Kokkos::parallel_for(
"access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE), *this);
Kokkos::parallel_for("access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE),
*this);
Kokkos::fence();
time_b = timer.seconds();
time += time_b - time_a;
Expand Down Expand Up @@ -244,22 +244,22 @@ struct Access_CudaAware<
MPI_Comm_size(MPI_COMM_WORLD, &num_ranks);
other_rank = my_rank ^ 1;
assert(num_ranks == 2);
iters_per_team = args.N / LEAGUE_SIZE;
iters_per_team = args.N / LEAGUE_SIZE;
iters_per_team_mod = args.N % LEAGUE_SIZE;
};

KOKKOS_FUNCTION
void operator()(const InitTag &, const size_t i) const { v(i) = my_rank + 1; }

KOKKOS_FUNCTION
void operator()(const UpdateTag &, typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for( Kokkos::TeamThreadRange(team, start, end + mod), [&](const int i){
v(i) += v_tmp(i);
});
void operator()(const UpdateTag &,
typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for(Kokkos::TeamThreadRange(team, start, end + mod),
[&](const int i) { v(i) += v_tmp(i); });
}

KOKKOS_FUNCTION
Expand Down Expand Up @@ -287,9 +287,9 @@ struct Access_CudaAware<
} else {
MPI_Recv(v_tmp.data(), N, MPI_DOUBLE, other_rank, TAG, MPI_COMM_WORLD,
MPI_STATUS_IGNORE);
Kokkos::parallel_for(
"access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE), *this);
Kokkos::parallel_for("access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE),
*this);
Kokkos::fence();
time_b = timer.seconds();
time += time_b - time_a;
Expand Down Expand Up @@ -337,7 +337,7 @@ struct Access<ViewType_t, typename std::enable_if_t<
v = ViewType_t(std::string(typeid(v).name()), num_ranks * args.N);
auto local_range =
Kokkos::Experimental::get_local_range(num_ranks * args.N);
iters_per_team = (local_range.second - local_range.first) / LEAGUE_SIZE;
iters_per_team = (local_range.second - local_range.first) / LEAGUE_SIZE;
iters_per_team_mod = (local_range.second - local_range.first) % LEAGUE_SIZE;
};

Expand All @@ -349,31 +349,28 @@ struct Access<ViewType_t, typename std::enable_if_t<
v(i) += v(other_rank * N + i);
}

KOKKOS_FUNCTION
void operator()(const UpdateTag_get &, typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for( Kokkos::TeamThreadRange(team, start, end + mod), [&](const int i){
v(i) += v(other_rank * N + i);
});
KOKKOS_FUNCTION
void operator()(const UpdateTag_get &,
typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for(Kokkos::TeamThreadRange(team, start, end + mod),
[&](const int i) { v(i) += v(other_rank * N + i); });
}



KOKKOS_FUNCTION
void operator()(const UpdateTag_put &, typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for( Kokkos::TeamThreadRange(team, start, end + mod), [&](const int i){
v(other_rank * N + i) = v(i);
});
KOKKOS_FUNCTION
void operator()(const UpdateTag_put &,
typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for(Kokkos::TeamThreadRange(team, start, end + mod),
[&](const int i) { v(other_rank * N + i) = v(i); });
}


KOKKOS_FUNCTION
void operator()(const CheckTag &, const size_t i) const {
assert(v(i) == typename ViewType_t::traits::value_type(
Expand All @@ -398,9 +395,9 @@ struct Access<ViewType_t, typename std::enable_if_t<
for (int i = 0; i < iters; i++) {
if (my_rank == 0) {
time_a = timer.seconds();
Kokkos::parallel_for(
"access_overhead",
team_policy_get_update_t(LEAGUE_SIZE, TEAM_SIZE), *this);
Kokkos::parallel_for("access_overhead",
team_policy_get_update_t(LEAGUE_SIZE, TEAM_SIZE),
*this);
Kokkos::fence();
RemoteSpace_t().fence();
time_b = timer.seconds();
Expand Down Expand Up @@ -487,7 +484,6 @@ struct Access_LDC<
int iters_per_team;
int iters_per_team_mod;


Access_LDC(Args_t args)
: N(args.N), iters(args.iters), mode(args.mode), rma_op(args.rma_op) {
MPI_Comm_rank(MPI_COMM_WORLD, &my_rank);
Expand All @@ -498,7 +494,7 @@ struct Access_LDC<
v_tmp = ViewType_t(std::string(typeid(v).name()), num_ranks * args.N);
auto local_range =
Kokkos::Experimental::get_local_range(num_ranks * args.N);
iters_per_team = (local_range.second - local_range.first) / LEAGUE_SIZE;
iters_per_team = (local_range.second - local_range.first) / LEAGUE_SIZE;
iters_per_team_mod = (local_range.second - local_range.first) % LEAGUE_SIZE;
};

Expand All @@ -518,15 +514,15 @@ struct Access_LDC<
iters * (other_rank + 1) + (my_rank + 1)));
}

KOKKOS_FUNCTION
void operator()(const UpdateTag &, typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for( Kokkos::TeamThreadRange(team, start, end + mod), [&](const int i){
v(i) += v_tmp(i);
});
KOKKOS_FUNCTION
void operator()(const UpdateTag &,
typename team_policy_update_t::member_type team) const {
int team_id = team.league_rank();
int start = team_id * iters_per_team;
int end = start + iters_per_team;
int mod = (team_id == LEAGUE_SIZE - 1) ? iters_per_team_mod : 0;
Kokkos::parallel_for(Kokkos::TeamThreadRange(team, start, end + mod),
[&](const int i) { v(i) += v_tmp(i); });
}

KOKKOS_FUNCTION
Expand Down Expand Up @@ -585,9 +581,9 @@ struct Access_LDC<
*this);
Kokkos::fence();
#endif
Kokkos::parallel_for(
"access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE), *this);
Kokkos::parallel_for("access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE),
*this);
Kokkos::fence();
RemoteSpace_t().fence();
time_b = timer.seconds();
Expand All @@ -612,9 +608,9 @@ struct Access_LDC<
Kokkos::RangePolicy(local_range.first, local_range.second),
*this);
#endif
Kokkos::parallel_for(
"access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE), *this);
Kokkos::parallel_for("access_overhead",
team_policy_update_t(LEAGUE_SIZE, TEAM_SIZE),
*this);
Kokkos::fence();
RemoteSpace_t().fence();
time_b = timer.seconds();
Expand Down
4 changes: 2 additions & 2 deletions src/core/Kokkos_RemoteSpaces_LocalDeepCopy.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -119,8 +119,8 @@ void KOKKOS_INLINE_FUNCTION local_deep_copy_contiguous(
auto team_range = Kokkos::pair(size_type(start_offset),
size_type(start_offset + team_block));

if(team_range.first == team_range.second) return; //nothing to be done
if (team_range.first == team_range.second) return; // nothing to be done

// Construct per-team subviews
auto src_subview = Kokkos::Impl::get_local_subview(src, team_range);
auto dst_subview = Kokkos::Impl::get_local_subview(dst, team_range);
Expand Down
6 changes: 2 additions & 4 deletions src/core/Kokkos_RemoteSpaces_ViewMapping.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -344,7 +344,7 @@ class ViewMapping<
typename view_type::size_type offset;
offset = switch_to_local_indexing ? total_offset : local_offset;
dst.remote_view_props.total_offset = total_offset;

#ifdef KRS_ENABLE_MPISPACE
// Subviews propagate MPI_Window of the original view
dst.m_handle = ViewDataHandle<DstTraits>::assign(
Expand Down Expand Up @@ -396,9 +396,7 @@ class ViewMapping<Traits, Kokkos::Experimental::RemoteSpaceSpecializeTag> {
int get_PE() const { return remote_view_props.my_PE; }

KOKKOS_INLINE_FUNCTION
auto get_ptr() const {
return handle().ptr + remote_view_props.total_offset;
}
auto get_ptr() const { return handle().ptr + remote_view_props.total_offset; }

template <typename T = Traits>
KOKKOS_INLINE_FUNCTION int get_logical_PE(ENABLE_IF_GLOBAL_LAYOUT(T)) const {
Expand Down
8 changes: 1 addition & 7 deletions unit_tests/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,13 +12,7 @@ if(NOT googletest_POPULATED)
endif()

SET(NAME KokkosRemoteSpaces_TestAll)

#if (KRS_ENABLE_MPISPACE)
# FILE(GLOB TEST_SRCS *.cpp)
# list(FILTER TEST_SRCS EXCLUDE REGEX ".*Test_Atomic\\.cpp$")
#else()
FILE(GLOB TEST_SRCS *.cpp)
#endif()
FILE(GLOB TEST_SRCS *.cpp)

add_executable(${NAME} ${TEST_SRCS})

Expand Down
7 changes: 5 additions & 2 deletions unit_tests/Test_Empty.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,9 @@

using RemoteSpace_t = Kokkos::Experimental::DefaultRemoteMemorySpace;

void test_empty() { RemoteSpace_t::fence(); }
void test_empty() {}

TEST(TEST_CATEGORY, test_empty) { test_empty(); }
TEST(TEST_CATEGORY, test_empty) {
test_empty();
RemoteSpace_t::fence();
}
4 changes: 2 additions & 2 deletions unit_tests/Test_LocalDeepCopy.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -988,10 +988,10 @@ void test_localdeepcopy_withSubview(
test_localdeepcopy_withSubview<int, Kokkos::HostSpace, RemoteSpace_t, TEAM, \
get_op, IS_RANGES, LAYOUT>(12, 15); \
/* 2D with Subviews (put block transfer) */ \
test_localdeepcopy_withSubview<int, Kokkos::HostSpace, RemoteSpace_t, TEAM, \
test_localdeepcopy_withSubview<int, Kokkos::HostSpace, RemoteSpace_t, TEAM, \
put_op, IS_RANGES, LAYOUT>(5, 16); \
/* 2D with Subviews (get block transfer)*/ \
test_localdeepcopy_withSubview<int64_t, Kokkos::HostSpace, RemoteSpace_t, \
test_localdeepcopy_withSubview<int64_t, Kokkos::HostSpace, RemoteSpace_t, \
TEAM, get_op, IS_RANGES, LAYOUT>(12, 15); \
/* 2D with Subviews (put block transfer)*/ \
test_localdeepcopy_withSubview<int64_t, Kokkos::HostSpace, RemoteSpace_t, \
Expand Down
Loading

0 comments on commit c595c7b

Please sign in to comment.