Skip to content
Snippets Groups Projects
Commit b251ba7c authored by Harshada Balasubramanian's avatar Harshada Balasubramanian
Browse files

general cleanup

parent 3949a1db
No related branches found
No related tags found
No related merge requests found
Hide whitespace changes
Inline Side-by-side
Showing
with 165 additions and 158 deletions
main.cpp
+ 165
− 158
View file @ b251ba7c
......@@ -7,8 +7,6 @@
#include "Kokkos_Timer.hpp"
using space_t = Kokkos::DefaultExecutionSpace::memory_space;
typedef Kokkos::MDRangePolicy<Kokkos::DefaultExecutionSpace, Kokkos::IndexType<int>, Kokkos::Rank<3>> md_range_policy;
Kokkos::Timer timer;
// constexpr int nblocks = 2;
......@@ -17,6 +15,128 @@ Kokkos::Timer timer;
// #define ENABLE_CHECK_BOUNDS
static void validate(double* array, int nblocks, int nlev, int nproma);
KOKKOS_INLINE_FUNCTION void check_bounds(int i1, int i2, int i3, int n1, int n2, int n3);
void scenario_1(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_2(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_2b(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_3(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_4(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_4b(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_5(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_6(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_7(double* array, int nblocks, int nlev, int nproma, bool print);
void scenario_8(double* array, int nblocks, int nlev, int nproma);
void scenario_1b(double* array, int nblocks, int nlev, int nproma) {
std::cout << "scenario 1b: Default layout right; view(array, nblocks, nlev, nproma); d_view(jb, jk, jc) ----- "
<< std::endl;
Kokkos::View<double***, Kokkos::LayoutRight, Kokkos::MemoryUnmanaged> d_view(array, nblocks, nlev, nproma);
timer.reset();
for (int jb = 0; jb < nblocks; ++jb)
Kokkos::parallel_for(
"", Kokkos::RangePolicy<>(0, nproma), KOKKOS_LAMBDA(const int jc) {
for (int jk = 0; jk < nlev; ++jk) {
int p = jb * nlev * nproma + jk * nproma + jc;
check_bounds(jb, jk, jc, d_view.extent(0), d_view.extent(1), d_view.extent(2));
d_view(jb, jk, jc) = p;
}
});
Kokkos::fence();
printf("Time = %f ms\n\n", timer.seconds() * 1000);
}
void scenario_7b(double* array, int nblocks, int nlev, int nproma) {
std::cout << "scenario 7b: Default layout right; view(array, nblocks, nlev, nproma); d_view(jb, jk, jc) ----- " << std::endl;
Kokkos::View<double***, Kokkos::MemoryUnmanaged> d_view(array, nblocks, nlev, nproma);
timer.reset();
Kokkos::parallel_for("", Kokkos::RangePolicy<>(0, nblocks), KOKKOS_LAMBDA (const int jb) {
//for (int jb = 0 ; jb < nblocks; ++jb) {
for (int jc = 0 ; jc < nproma; ++jc) {
for (int jk = 0; jk < nlev; ++jk) {
int p = jb * nlev * nproma + jk * nproma + jc;
check_bounds(jb,jk,jc, d_view.extent(0), d_view.extent(1), d_view.extent(2));
d_view(jb, jk, jc) = p;
}
}
});
Kokkos::fence();
printf("Time = %f ms\n\n", timer.seconds() * 1000);
}
void openacc_calls(double* array, int nblocks, int nlev, int nproma,
std::function<void(double*, int, int, int)> func){
#if defined(DEMO_DEVICE)
#pragma acc enter data copyin(array[0:nblocks*nlev*nproma])
{
#pragma acc host_data use_device(array)
{
func(array, nblocks, nlev, nproma);
}
}
#pragma acc exit data copyout(array[0:nblocks*nlev*nproma])
#else
func(array, nblocks, nlev, nproma);
#endif
validate(array, nblocks, nlev, nproma);
}
int main() {
int ncells = atoi(std::getenv("NCELLS"));
int nlev = atoi(std::getenv("NLEV"));
int nproma = atoi(std::getenv("NPROMA"));
int nblocks = (ncells - 1) / nproma + 1;
std::cout << "nblocks=" << nblocks << ", nlev=" << nlev << ", nproma=" << nproma << std::endl;
double array[nblocks * nlev * nproma];
Kokkos::initialize();
{
scenario_1(array, nblocks, nlev, nproma, false);
//scenario_1(array, nblocks, nlev, nproma);
memset(array, 0.0, sizeof(array));
std::function<void(double*, int, int, int)> s_1b = scenario_1b;
openacc_calls(array, nblocks, nlev, nproma, s_1b);
memset(array, 0.0, sizeof(array));
std::function<void(double*, int, int, int)> s_7b = scenario_7b;
openacc_calls(array, nblocks, nlev, nproma, s_7b);
memset(array, 0.0, sizeof(array));
std::function<void(double*, int, int, int)> s_8 = scenario_8;
openacc_calls(array, nblocks, nlev, nproma, s_8);
/*
scenario_2(array, nblocks, nlev, nproma);
scenario_2b(array, nblocks, nlev, nproma);
scenario_3(array, nblocks, nlev, nproma);
scenario_4(array, nblocks, nlev, nproma);
scenario_4b(array, nblocks, nlev, nproma);
scenario_5(array, nblocks, nlev, nproma);
scenario_6(array, nblocks, nlev, nproma);
scenario_7(array, nblocks, nlev, nproma);
*/
}
Kokkos::finalize();
return 0;
}
static void validate(double* array, int nblocks, int nlev, int nproma) {
for (int i = 0; i < nblocks * nlev * nproma; ++i) {
assert(array[i] == static_cast<double>(i));
......@@ -24,9 +144,9 @@ static void validate(double* array, int nblocks, int nlev, int nproma) {
}
KOKKOS_INLINE_FUNCTION void check_bounds(int i1, int i2, int i3, int n1, int n2, int n3) {
#ifdef ENABLE_CHECK_BOUNDS
assert(i1 >= 0 && i2 >= 0 && i3 >= 0 && i1 < n1 && i2 < n2 && i3 < n3);
#endif
#ifdef ENABLE_CHECK_BOUNDS
assert(i1 >= 0 && i2 >= 0 && i3 >= 0 && i1 < n1 && i2 < n2 && i3 < n3);
#endif
}
void scenario_1(double* array, int nblocks, int nlev, int nproma, bool print = true) {
......@@ -59,29 +179,6 @@ void scenario_1(double* array, int nblocks, int nlev, int nproma, bool print = t
validate(array, nblocks, nlev, nproma);
}
void scenario_1b(double* array, int nblocks, int nlev, int nproma) {
std::cout << "scenario 1b: Default layout right; view(array, nblocks, nlev, nproma); d_view(jb, jk, jc) ----- "
<< std::endl;
Kokkos::View<double***, Kokkos::LayoutRight, Kokkos::MemoryUnmanaged> d_view(array, nblocks, nlev, nproma);
timer.reset();
for (int jb = 0; jb < nblocks; ++jb)
Kokkos::parallel_for(
"", Kokkos::RangePolicy<>(0, nproma), KOKKOS_LAMBDA(const int jc) {
for (int jk = 0; jk < nlev; ++jk) {
int p = jb * nlev * nproma + jk * nproma + jc;
check_bounds(jb, jk, jc, d_view.extent(0), d_view.extent(1), d_view.extent(2));
d_view(jb, jk, jc) = p;
}
});
Kokkos::fence();
printf("Time = %f ms\n\n", timer.seconds() * 1000);
}
void scenario_2(double* array, int nblocks, int nlev, int nproma, bool print = true) {
if (print)
......@@ -329,152 +426,62 @@ void scenario_6(double* array, int nblocks, int nlev, int nproma, bool print = t
void scenario_7(double* array, int nblocks, int nlev, int nproma, bool print=true) {
if(print)
std::cout << "scenario 7: Default layout; view(array, nblocks, nlev, nproma); d_view(jb, jk, jc) ----- " << std::endl;
Kokkos::View<double***, Kokkos::HostSpace, Kokkos::MemoryUnmanaged> view(array, nblocks, nlev, nproma);
using space_t = Kokkos::DefaultExecutionSpace::memory_space;
auto d_view = Kokkos::create_mirror_view_and_copy(space_t(), view);
timer.reset();
Kokkos::parallel_for("", Kokkos::RangePolicy<>(0, nblocks), KOKKOS_LAMBDA (const int jb) {
//for (int jb = 0 ; jb < nblocks; ++jb) {
for (int jc = 0 ; jc < nproma; ++jc) {
for (int jk = 0; jk < nlev; ++jk) {
int p = jb * nlev * nproma + jk * nproma + jc;
check_bounds(jb,jk,jc, d_view.extent(0), d_view.extent(1), d_view.extent(2));
d_view(jb, jk, jc) = p;
// printf("%f ", d_view(jb, jk, jc));
}
}
});
if(print)
std::cout << "scenario 7: Default layout; view(array, nblocks, nlev, nproma); d_view(jb, jk, jc) ----- " << std::endl;
Kokkos::fence();
if(print)
printf("Time = %f ms\n\n", timer.seconds() * 1000);
Kokkos::deep_copy(view, d_view);
validate(array, nblocks, nlev, nproma);
}
void scenario_7b(double* array, int nblocks, int nlev, int nproma) {
std::cout << "scenario 7b: +ACC Default layout; view(array, nblocks, nlev, nproma); d_view(jb, jk, jc) ----- " << std::endl;
Kokkos::View<double***, Kokkos::MemoryUnmanaged> d_view(array, nblocks, nlev, nproma);
timer.reset();
Kokkos::parallel_for("", Kokkos::RangePolicy<>(0, nblocks), KOKKOS_LAMBDA (const int jb) {
//for (int jb = 0 ; jb < nblocks; ++jb) {
for (int jc = 0 ; jc < nproma; ++jc) {
for (int jk = 0; jk < nlev; ++jk) {
#if defined(DEMO_DEVICE)
int p = jc * nlev * nblocks + jk * nblocks + jb;
#else
int p = jb * nlev * nproma + jk * nproma + jc;
Kokkos::View<double***, Kokkos::HostSpace, Kokkos::MemoryUnmanaged> view(array, nblocks, nlev, nproma);
#endif
check_bounds(jb,jk,jc, d_view.extent(0), d_view.extent(1), d_view.extent(2));
d_view(jb, jk, jc) = p;
// printf("%f ", d_view(jb, jk, jc));
}
using space_t = Kokkos::DefaultExecutionSpace::memory_space;
auto d_view = Kokkos::create_mirror_view_and_copy(space_t(), view);
timer.reset();
Kokkos::parallel_for("", Kokkos::RangePolicy<>(0, nblocks), KOKKOS_LAMBDA (const int jb) {
//for (int jb = 0 ; jb < nblocks; ++jb) {
for (int jc = 0 ; jc < nproma; ++jc) {
for (int jk = 0; jk < nlev; ++jk) {
int p = jb * nlev * nproma + jk * nproma + jc;
check_bounds(jb,jk,jc, d_view.extent(0), d_view.extent(1), d_view.extent(2));
d_view(jb, jk, jc) = p;
// printf("%f ", d_view(jb, jk, jc));
}
});
}
});
Kokkos::fence();
printf("Time = %f ms\n\n", timer.seconds() * 1000);
Kokkos::fence();
if(print)
printf("Time = %f ms\n\n", timer.seconds() * 1000);
Kokkos::deep_copy(view, d_view);
validate(array, nblocks, nlev, nproma);
}
void scenario_8(double* array, int nblocks, int nlev, int nproma) {
std::cout << "scenario 8: Default Layout as ICON pointer view(array, nproma, nlev, nblocks); d_view(jb, jk, jc) ----- " << std::endl;
std::cout << "scenario 8: Default Layout as ICON pointer view(array, nproma, nlev, nblocks); d_view(jb, jk, jc) ----- " << std::endl;
using space_t = Kokkos::DefaultExecutionSpace::memory_space;
Kokkos::View<double***, Kokkos::MemoryUnmanaged> d_view(array, nproma, nlev, nblocks);
using space_t = Kokkos::DefaultExecutionSpace::memory_space;
Kokkos::View<double***, Kokkos::MemoryUnmanaged> d_view(array, nproma, nlev, nblocks);
timer.reset();
for (int jb = 0; jb < nblocks; ++jb)
Kokkos::parallel_for(
"", Kokkos::RangePolicy<>(0, nproma), KOKKOS_LAMBDA(const int jc) {
for (int jk = 0; jk < nlev; ++jk) {
timer.reset();
for (int jb = 0; jb < nblocks; ++jb)
Kokkos::parallel_for(
"", Kokkos::RangePolicy<>(0, nproma), KOKKOS_LAMBDA(const int jc) {
for (int jk = 0; jk < nlev; ++jk) {
#if defined(DEMO_DEVICE)
int p = jb * nlev * nproma + jk * nproma + jc;
int p = jb * nlev * nproma + jk * nproma + jc;
#else
int p = jc * nlev * nblocks + jk * nblocks + jb;
int p = jc * nlev * nblocks + jk * nblocks + jb;
#endif
d_view(jc, jk, jb) = p;
d_view(jc, jk, jb) = p;
// printf("%f ", d_view(jb, jk, jc));
}
});
Kokkos::fence();
// printf("%f ", d_view(jb, jk, jc));
}
});
Kokkos::fence();
printf("Time = %f ms\n\n", timer.seconds() * 1000);
printf("Time = %f ms\n\n", timer.seconds() * 1000);
}
void openacc_calls(double* array, int nblocks, int nlev, int nproma,
std::function<void(double*, int, int, int)> func){
#if defined(DEMO_DEVICE)
#pragma acc enter data copyin(array[0:nblocks*nlev*nproma])
{
#pragma acc host_data use_device(array)
{
func(array, nblocks, nlev, nproma);
}
}
#pragma acc exit data copyout(array[0:nblocks*nlev*nproma])
#else
func(array, nblocks, nlev, nproma);
#endif
validate(array, nblocks, nlev, nproma);
}
int main() {
int ncells = atoi(std::getenv("NCELLS"));
int nlev = atoi(std::getenv("NLEV"));
int nproma = atoi(std::getenv("NPROMA"));
int nblocks = (ncells - 1) / nproma + 1;
std::cout << "nblocks=" << nblocks << ", nlev=" << nlev << ", nproma=" << nproma << std::endl;
double array[nblocks * nlev * nproma];
Kokkos::initialize();
{
scenario_1(array, nblocks, nlev, nproma, false);
//scenario_1(array, nblocks, nlev, nproma);
memset(array, 0.0, sizeof(array));
std::function<void(double*, int, int, int)> s_1b = scenario_1b;
openacc_calls(array, nblocks, nlev, nproma, s_1b);
memset(array, 0.0, sizeof(array));
std::function<void(double*, int, int, int)> s_7b = scenario_7b;
openacc_calls(array, nblocks, nlev, nproma, s_7b);
memset(array, 0.0, sizeof(array));
std::function<void(double*, int, int, int)> s_8 = scenario_8;
openacc_calls(array, nblocks, nlev, nproma, s_8);
/*
scenario_2(array, nblocks, nlev, nproma);
scenario_2b(array, nblocks, nlev, nproma);
scenario_3(array, nblocks, nlev, nproma);
scenario_4(array, nblocks, nlev, nproma);
scenario_4b(array, nblocks, nlev, nproma);
scenario_5(array, nblocks, nlev, nproma);
scenario_6(array, nblocks, nlev, nproma);
scenario_7(array, nblocks, nlev, nproma);
*/
}
Kokkos::finalize();
return 0;
}
\ No newline at end of file
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment

Imprint | Privacy Policy