Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 3 additions & 5 deletions benchmarks/access_overhead/access_overhead.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -117,13 +117,13 @@ 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();
for (int i = 0; i < iters; i++) {
Kokkos::parallel_for("access_overhead", policy_update_t(0, N), *this);
RemoteSpace_t().fence();
Kokkos::fence();
}
time_b = timer.seconds();
time += time_b - time_a;
Expand Down Expand Up @@ -183,7 +183,7 @@ struct Access<ViewType_t, typename std::enable_if_t<std::is_same<
time_a = timer.seconds();
for (int i = 0; i < iters; i++) {
Kokkos::parallel_for("access_overhead", policy_update_t(0, N), *this);
RemoteSpace_t().fence();
Kokkos::fence();
}
time_b = timer.seconds();
time += time_b - time_a;
Expand Down Expand Up @@ -248,8 +248,6 @@ int main(int argc, char *argv[]) {
}
} while (false);

Kokkos::fence();

Kokkos::finalize();
#ifdef KRS_ENABLE_SHMEMSPACE
shmem_finalize();
Expand Down
11 changes: 6 additions & 5 deletions benchmarks/access_overhead/access_overhead_p2p.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -24,13 +24,14 @@
#include <type_traits>
#include <string>

#define LDC_LEAGUE_SIZE 4096
#define LDC_LEAGUE_SIZE 2
#define LDC_TEAM_SIZE 1
//#define CHECK_FOR_CORRECTNESS

#define CHECK_FOR_CORRECTNESS

using RemoteSpace_t = Kokkos::Experimental::DefaultRemoteMemorySpace;
using RemoteView_t = Kokkos::View<double *, RemoteSpace_t>;
using PlainView_t = Kokkos::View<double *, Kokkos::LayoutLeft>;
using PlainView_t = Kokkos::View<double *>;
using UnmanagedView_t =
Kokkos::View<double *, Kokkos::MemoryTraits<Kokkos::Unmanaged>>;
using HostView_t = typename RemoteView_t::HostMirror;
Expand All @@ -53,8 +54,8 @@ using policy_check_t = Kokkos::RangePolicy<CheckTag, size_t>;

// Default values
#define default_Mode 0
#define default_N 134217728
#define default_Iters 3
#define default_N 128
#define default_Iters 1
#define default_RmaOp RMA_GET
#define TAG 0

Expand Down
Loading