From f6bdbe8212d72dc2a5c76a229b63c9ca3263b9d7 Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Fri, 18 Nov 2022 10:04:53 -0800 Subject: [PATCH 1/3] raja launch + chai integration test --- tests/integration/CMakeLists.txt | 23 +- tests/integration/raja-chai-launch.cpp | 409 +++++++++++++++++++++++++ 2 files changed, 429 insertions(+), 3 deletions(-) create mode 100644 tests/integration/raja-chai-launch.cpp diff --git a/tests/integration/CMakeLists.txt b/tests/integration/CMakeLists.txt index b29a5046..e551cf21 100644 --- a/tests/integration/CMakeLists.txt +++ b/tests/integration/CMakeLists.txt @@ -39,14 +39,14 @@ if (CHAI_ENABLE_MANAGED_PTR) NAME managed_ptr_test COMMAND managed_ptr_tests) endif () - + if (CHAI_ENABLE_RAJA_PLUGIN) set(raja_test_depends ${chai_integration_test_depends} RAJA) blt_add_executable( - NAME raja-chai-tests + NAME raja-chai-tests SOURCES raja-chai-tests.cpp DEPENDS_ON ${raja_test_depends}) @@ -59,8 +59,9 @@ if (CHAI_ENABLE_RAJA_PLUGIN) PUBLIC ${PROJECT_BINARY_DIR}/include) if (CHAI_ENABLE_RAJA_NESTED_TEST) +#raja kernel tests blt_add_executable( - NAME raja-chai-nested-tests + NAME raja-chai-nested-tests SOURCES raja-chai-nested.cpp DEPENDS_ON ${raja_test_depends}) @@ -71,5 +72,21 @@ if (CHAI_ENABLE_RAJA_PLUGIN) target_include_directories( raja-chai-nested-tests PUBLIC ${PROJECT_BINARY_DIR}/include) + +#raja launch tests + blt_add_executable( + NAME raja-chai-launch-tests + SOURCES raja-chai-launch.cpp + DEPENDS_ON ${raja_test_depends}) + + blt_add_test( + NAME raja-chai-launch-tests + COMMAND raja-chai-launch-tests) + + target_include_directories( + raja-chai-launch-tests + PUBLIC ${PROJECT_BINARY_DIR}/include) + endif () + endif () diff --git a/tests/integration/raja-chai-launch.cpp b/tests/integration/raja-chai-launch.cpp new file mode 100644 index 00000000..13dfcf77 --- /dev/null +++ b/tests/integration/raja-chai-launch.cpp @@ -0,0 +1,409 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the COPYRIGHT file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// +/// +/// Source file containing tests for CHAI in RAJA nested loops. +/// +/// +#include +#include +#include + +#include +#include +#include + +#include "RAJA/RAJA.hpp" + +using namespace RAJA; +using namespace std; + +#include "chai/ArrayManager.hpp" +#include "chai/ManagedArrayView.hpp" +#include "chai/ManagedArray.hpp" + +#include "gtest/gtest.h" + +// TODO: add hip policy for these tests. +#if defined(RAJA_ENABLE_CUDA) +#define PARALLEL_RAJA_DEVICE __device__ +#elif defined(RAJA_ENABLE_OPENMP) +#define PARALLEL_RAJA_DEVICE +#else +#define PARALLEL_RAJA_DEVICE +#endif + +#define CUDA_TEST(X, Y) \ + static void cuda_test_##X##_##Y(); \ + TEST(X, Y) { cuda_test_##X##_##Y(); } \ + static void cuda_test_##X##_##Y() + +/* + * Simple tests using nested::forall and View + */ +CUDA_TEST(Chai, LaunchSimple) +{ + + using LAUNCH_POLICY = RAJA::LaunchPolicy; + using LOOP_POLICY_0 = RAJA::LoopPolicy; + using LOOP_POLICY_1 = RAJA::LoopPolicy; + +#if defined(RAJA_ENABLE_CUDA) + + const bool async = false; + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy>; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + +#elif defined(RAJA_ENABLE_OPENMP) + + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + +#else + using PARALLEL_LAUNCH_POLICY = LAUNCH_POLICY; + using PARALLEL_LOOP_POLICY_0 = LOOP_POLICY_0; + using PARALLEL_LOOP_POLICY_1 = LOOP_POLICY_1; +#endif + + const int X = 16; + const int Y = 16; + + chai::ManagedArray v1(X * Y); + chai::ManagedArray v2(X * Y); + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + int index = j * X + i; + v1[index] = index; + }); + }); + }); + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + int index = j * X + i; + v2[index] = v1[index] * 2.0f; + }); + }); + }); + + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + int index = j * X + i; + ASSERT_FLOAT_EQ(v1[index], index * 1.0f); + ASSERT_FLOAT_EQ(v2[index], index * 2.0f); + }); + }); + }); + + v1.free(); + v2.free(); +} + +CUDA_TEST(Chai, LaunchView) +{ + + using LAUNCH_POLICY = RAJA::LaunchPolicy; + using LOOP_POLICY_0 = RAJA::LoopPolicy; + using LOOP_POLICY_1 = RAJA::LoopPolicy; + +#if defined(RAJA_ENABLE_CUDA) + + const bool async = false; + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy>; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + +#elif defined(RAJA_ENABLE_OPENMP) + + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + +#else + + using PARALLEL_LAUNCH_POLICY = LAUNCH_POLICY; + using PARALLEL_LOOP_POLICY_0 = LOOP_POLICY_0; + using PARALLEL_LOOP_POLICY_1 = LOOP_POLICY_1; +#endif + + const int X = 16; + const int Y = 16; + + chai::ManagedArray v1_array(X * Y); + chai::ManagedArray v2_array(X * Y); + + using view = chai::ManagedArrayView>; + + view v1(v1_array, X, Y); + view v2(v2_array, X, Y); + + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + v1(i, j) = (i + (j * X)) * 1.0f; + }); + }); + }); + + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + v2(i, j) = v1(i, j) * 2.0f; + }); + }); + }); + + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + ASSERT_FLOAT_EQ(v2(i, j), v1(i, j) * 2.0f); + }); + }); + }); + + v1_array.free(); + v2_array.free(); +} + +CUDA_TEST(Chai, LaunchMultiView) +{ + + using LAUNCH_POLICY = RAJA::LaunchPolicy; + using LOOP_POLICY_0 = RAJA::LoopPolicy; + using LOOP_POLICY_1 = RAJA::LoopPolicy; + +#if defined(RAJA_ENABLE_CUDA) + + const bool async = false; + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy>; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + +#elif defined(RAJA_ENABLE_OPENMP) + + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; +#else + + using PARALLEL_LAUNCH_POLICY = LAUNCH_POLICY; + using PARALLEL_LOOP_POLICY_0 = LOOP_POLICY_0; + using PARALLEL_LOOP_POLICY_1 = LOOP_POLICY_1; +#endif + + const int X = 16; + const int Y = 16; + + chai::ManagedArray v1_array(X * Y); + chai::ManagedArray v2_array(X * Y); + + chai::ManagedArray all_arrays[2]; + all_arrays[0] = v1_array; + all_arrays[1] = v2_array; + + // default MultiView + using view = chai::ManagedArrayMultiView>; + view mview(all_arrays, RAJA::Layout<2>(X, Y)); + + // MultiView with index in 1st position + using view1p = chai::ManagedArrayMultiView, 1>; + view1p mview1p(all_arrays, RAJA::Layout<2>(X, Y)); + + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + mview(0, i, j) = (i + (j * X)) * 1.0f; + }); + }); + }); + + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + // use both MultiViews + mview(1, i, j) = mview1p(i, 0, j) * 2.0f; + }); + }); + }); + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::RangeSegment(0, X), [&] (int i) { + RAJA::loop(ctx, RAJA::RangeSegment(0, Y), [&] (int j) { + ASSERT_FLOAT_EQ(mview(1, i, j), mview(0, i, j) * 2.0f); + }); + }); + }); + + + v1_array.free(); + v2_array.free(); +} + +/////////////////////////////////////////////////////////////////////////// +// +// Example LTimes kernel test routines +// +// Demonstrates a 4-nested loop, the use of complex nested policies and +// the use of strongly-typed indices +// +// This routine computes phi(m, g, z) = SUM_d { ell(m, d)*psi(d,g,z) } +// +/////////////////////////////////////////////////////////////////////////// + +RAJA_INDEX_VALUE_T(IM, int, "IM"); +RAJA_INDEX_VALUE_T(ID, int, "ID"); +RAJA_INDEX_VALUE_T(IG, int, "IG"); +RAJA_INDEX_VALUE_T(IZ, int, "IZ"); + +void runLTimesTests(Index_type num_moments, + Index_type num_directions, + Index_type num_groups, + Index_type num_zones) +{ + // allocate data + // phi is initialized to all zeros, the others are randomized + chai::ManagedArray L_data(num_moments * num_directions); + chai::ManagedArray psi_data(num_directions * num_groups * num_zones); + chai::ManagedArray phi_data(num_moments * num_groups * num_zones); + + RAJA::forall( + RAJA::RangeSegment(0, (num_moments * num_directions)), + [=](int i) { + L_data[i] = i+2; + }); + + RAJA::forall( + RAJA::RangeSegment(0, (num_directions * num_groups * num_zones)), + [=](int i) { psi_data[i] = 2*i+1; }); + + RAJA::forall( + RAJA::RangeSegment(0, (num_moments * num_groups * num_zones)), + [=](int i) { phi_data[i] = 0.0; }); + + using LView = chai::TypedManagedArrayView, IM, ID>; + + // psi(d, g, z) : 2 -> z is stride-1 dimension + using PsiView = chai::TypedManagedArrayView, ID, IG, IZ>; + + // phi(m, g, z) : 2 -> z is stride-1 dimension + using PhiView = chai::TypedManagedArrayView, IM, IG, IZ>; + + std::array L_perm {{0, 1}}; + LView L(L_data, + RAJA::make_permuted_layout({{num_moments, num_directions}}, L_perm)); + + std::array psi_perm {{0, 1, 2}}; + PsiView psi(psi_data, + RAJA::make_permuted_layout({{num_directions, num_groups, num_zones}}, psi_perm)); + + std::array phi_perm {{0, 1, 2}}; + PhiView phi(phi_data, + RAJA::make_permuted_layout({{num_moments, num_groups, num_zones}}, phi_perm)); + +#if defined(RAJA_ENABLE_CUDA) + const bool async = false; + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy>; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_2 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_3 = RAJA::LoopPolicy; + +#elif defined(RAJA_ENABLE_OPENMP) + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_2 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_3 = RAJA::LoopPolicy; +#else + using PARALLEL_LAUNCH_POLICY = RAJA::LaunchPolicy; + using PARALLEL_LOOP_POLICY_0 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_1 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_2 = RAJA::LoopPolicy; + using PARALLEL_LOOP_POLICY_3 = RAJA::LoopPolicy; +#endif + + RAJA::launch + (RAJA::LaunchParams(RAJA::Teams(1), RAJA::Threads(10)), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) + { + RAJA::loop(ctx, RAJA::TypedRangeSegment(0, num_moments), [&] (IM m) { + RAJA::loop(ctx, RAJA::TypedRangeSegment(0, num_groups), [&] (IG g) { + RAJA::loop(ctx, RAJA::TypedRangeSegment(0, num_zones ), [&] (IZ z) { + RAJA::loop(ctx, RAJA::TypedRangeSegment(0, num_directions ), [&] (ID d) { + phi(m, g, z) += L(m, d) * psi(d, g, z); + }); + }); + }); + }); + + }); + + RAJA::forall( + RAJA::TypedRangeSegment(0, num_moments), [=] (IM m) { + for (IG g(0); g < num_groups; ++g) { + for (IZ z(0); z < num_zones; ++z) { + double total = 0.0; + for (ID d(0); d < num_directions; ++d) { + double val = L(m, d) * psi(d, g, z); + total += val; + } + ASSERT_FLOAT_EQ(total, phi(m, g, z)); + } + } + }); + + L_data.free(); + psi_data.free(); + phi_data.free(); +} + +TEST(Chai, LaunchLTimes) +{ + // runLTimesTests(2, 0, 7, 3); + runLTimesTests(2, 3, 7, 3); + runLTimesTests(2, 3, 32, 4); + runLTimesTests(25, 96, 8, 32); + runLTimesTests(100, 15, 7, 13); +} From a2ec1a50d72bf39a484b53c9c6b051dbe2bb67db Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Fri, 18 Nov 2022 10:07:12 -0800 Subject: [PATCH 2/3] bump raja up to develop --- src/tpl/raja | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/tpl/raja b/src/tpl/raja index 4351fe6a..1664d4e9 160000 --- a/src/tpl/raja +++ b/src/tpl/raja @@ -1 +1 @@ -Subproject commit 4351fe6a50bd579511a625b017c9e054885e7fd2 +Subproject commit 1664d4e95bd21ccecc66aebdcc58052df19a3636 From 0107b28a5542c83d00727bb8b8a564c86840056b Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Thu, 15 Dec 2022 09:48:02 -0800 Subject: [PATCH 3/3] update raja tpl to v2022.10.4 --- src/tpl/raja | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/tpl/raja b/src/tpl/raja index 1664d4e9..c2a6b174 160000 --- a/src/tpl/raja +++ b/src/tpl/raja @@ -1 +1 @@ -Subproject commit 1664d4e95bd21ccecc66aebdcc58052df19a3636 +Subproject commit c2a6b1740759ae3ae7c85b35e20dbffbe235355d