From 8e84c53a64250fd51bfae636f5288b6026d1aa4d Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Thu, 30 Apr 2026 11:17:10 -0700 Subject: [PATCH 1/5] draft of forall_nd --- docs/sphinx/user_guide/feature/loop_basic.rst | 58 ++++++ examples/CMakeLists.txt | 4 + examples/forall_nd_layout.cpp | 197 ++++++++++++++++++ include/RAJA/RAJA.hpp | 1 + include/RAJA/pattern/forall_nd.hpp | 133 ++++++++++++ test/functional/util/CMakeLists.txt | 4 + test/functional/util/test-forall_nd.cpp | 68 ++++++ 7 files changed, 465 insertions(+) create mode 100644 examples/forall_nd_layout.cpp create mode 100644 include/RAJA/pattern/forall_nd.hpp create mode 100644 test/functional/util/test-forall_nd.cpp diff --git a/docs/sphinx/user_guide/feature/loop_basic.rst b/docs/sphinx/user_guide/feature/loop_basic.rst index 830f92c744..098ecf7a9f 100644 --- a/docs/sphinx/user_guide/feature/loop_basic.rst +++ b/docs/sphinx/user_guide/feature/loop_basic.rst @@ -412,3 +412,61 @@ the provided lambda with the appropriate indices. .. note:: CombiningAdapter currently only supports ``RAJA::TypedRangeSegment`` segments. + +.. _loop_elements-forall_nd-label: + +-------------------------------------------------------- +Flattened multi-dimensional loops with ``RAJA::forall_nd`` +-------------------------------------------------------- + +``RAJA::forall_nd`` is a convenience wrapper built on +``RAJA::dynamic_forall`` and ``RAJA::CombiningAdapter``. It executes a +flattened 1-D iteration space selected from a run-time policy list and +reconstructs the multi-dimensional indices before invoking the user lambda. +On CUDA, HIP, and SYCL backends this means the logical N-D iteration space is +flattened to a single 1-D device loop of length equal to the product of the +segment sizes. For example, ``cuda_exec``, ``hip_exec``, +and ``sycl_exec`` execute a 1-D kernel/work-group shape over that +flattened range, and ``RAJA::forall_nd`` performs the index reconstruction +before calling the user lambda. + +The API takes a policy list, a run-time policy index, a pack of +``RAJA::TypedRangeSegment`` objects created by ``RAJA::segments``, and a +lambda whose arguments are the recovered logical indices:: + + using policy_list = camp::list + #endif + #if defined(RAJA_ENABLE_HIP) + , RAJA::hip_exec<256> + #endif + #if defined(RAJA_ENABLE_SYCL) + , RAJA::sycl_exec<256> + #endif + >; + + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + +The layout tag controls which logical index is unit stride in the flattened +space. ``RAJA::layout_right`` makes the right-most index vary fastest: + +.. literalinclude:: ../../../../examples/forall_nd_layout.cpp + :start-after: // _forall_nd_right_start + :end-before: // _forall_nd_right_end + :language: C++ + +``RAJA::layout_left`` makes the left-most index vary fastest: + +.. literalinclude:: ../../../../examples/forall_nd_layout.cpp + :start-after: // _forall_nd_left_start + :end-before: // _forall_nd_left_end + :language: C++ + +The example source file ``RAJA/examples/forall_nd_layout.cpp`` contains a +complete runnable example. diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index e77b5dc493..19b132d878 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -61,6 +61,10 @@ raja_add_executable( NAME launch_flatten SOURCES launch_flatten.cpp) +raja_add_executable( + NAME forall_nd_layout + SOURCES forall_nd_layout.cpp) + raja_add_executable( NAME launch_reductions SOURCES launch_reductions.cpp) diff --git a/examples/forall_nd_layout.cpp b/examples/forall_nd_layout.cpp new file mode 100644 index 0000000000..ffe680c70e --- /dev/null +++ b/examples/forall_nd_layout.cpp @@ -0,0 +1,197 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) Lawrence Livermore National Security, LLC and other +// RAJA Project Developers. See top-level LICENSE and COPYRIGHT +// files for dates and other details. No copyright assignment is required +// to contribute to RAJA. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#include +#include +#include + +#include "RAJA/RAJA.hpp" + +/* + * Flattened N-D forall helper with layout-left/layout-right selection + * + * This example shows the public RAJA::forall_nd wrapper, which executes a + * flattened 1-D kernel selected from a run-time policy list and reconstructs + * the multi-dimensional indices for the user lambda. + * + * The public call shape is: + * + * RAJA::forall_nd( + * res, pol, RAJA::segments(rows, batches), + * [=] RAJA_HOST_DEVICE(int r, int b) { ... }); + * + * Supported layouts: + * - RAJA::layout_right : right-most index varies fastest + * - RAJA::layout_left : left-most index varies fastest + */ + +namespace +{ + +using policy_list = camp::list< + RAJA::seq_exec, + RAJA::simd_exec +#if defined(RAJA_ENABLE_OPENMP) + , + RAJA::omp_parallel_for_exec +#endif +#if defined(RAJA_ENABLE_CUDA) + , + RAJA::cuda_exec<256> +#endif +#if defined(RAJA_ENABLE_HIP) + , + RAJA::hip_exec<256> +#endif +#if defined(RAJA_ENABLE_SYCL) + , + RAJA::sycl_exec<256> +#endif + >; + +constexpr int host_policy_count = + 2 +#if defined(RAJA_ENABLE_OPENMP) + + 1 +#endif + ; + +RAJA::resources::Resource get_resource_for_policy(int pol) +{ + if (pol < 0) + { + RAJA_ABORT_OR_THROW("Policy value out of range for this build"); + } + + RAJA::resources::Host host_res; + if (pol < host_policy_count) + { + return RAJA::resources::Resource(host_res); + } + +#if defined(RAJA_ENABLE_CUDA) + if (pol == host_policy_count) + { + return RAJA::resources::Resource(RAJA::resources::Cuda {}); + } +#endif + +#if defined(RAJA_ENABLE_HIP) + if (pol == host_policy_count +#if defined(RAJA_ENABLE_CUDA) + + 1 +#endif + ) + { + return RAJA::resources::Resource(RAJA::resources::Hip {}); + } +#endif + +#if defined(RAJA_ENABLE_SYCL) + if (pol == host_policy_count +#if defined(RAJA_ENABLE_CUDA) + + 1 +#endif +#if defined(RAJA_ENABLE_HIP) + + 1 +#endif + ) + { + return RAJA::resources::Resource(RAJA::resources::Sycl {}); + } +#endif + + RAJA_ABORT_OR_THROW("Policy value out of range for this build"); +} + +void print_linear_order(char const* label, std::vector const& values) +{ + std::cout << " " << label << ":"; + for (auto value : values) + { + std::cout << ' ' << value; + } + std::cout << '\n'; +} + +} // namespace + +int main(int argc, char** argv) +{ + const int pol = (argc > 1) ? std::atoi(argv[1]) : 0; + + std::cout << "\nRAJA flattened N-D forall helper example...\n"; + std::cout << "Using policy #" << pol << '\n'; + + constexpr int n = 3; + constexpr int batch_size = 4; + constexpr int total_size = n * batch_size; + + auto res = get_resource_for_policy(pol); + + int* right_ptr = res.allocate(total_size); + int* left_ptr = res.allocate(total_size); + + res.memset(right_ptr, 0xff, sizeof(int) * total_size); + res.memset(left_ptr, 0xff, sizeof(int) * total_size); + + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + + // layout_right makes the right-most logical index ('b') unit stride. + // _forall_nd_right_start + RAJA::forall_nd( + res, pol, RAJA::segments(rows, batches), + [=] RAJA_HOST_DEVICE(int r, int b) { + const int flat = b + batch_size * r; + right_ptr[flat] = 100 * r + b; + }); + // _forall_nd_right_end + + // layout_left makes the left-most logical index ('r') unit stride. + // _forall_nd_left_start + RAJA::forall_nd( + res, pol, RAJA::segments(rows, batches), + [=] RAJA_HOST_DEVICE(int r, int b) { + const int flat = r + n * b; + left_ptr[flat] = 100 * r + b; + }); + // _forall_nd_left_end + + std::vector right(total_size); + std::vector left(total_size); + std::vector expect_right(total_size); + std::vector expect_left(total_size); + + res.wait(); + res.memcpy(right.data(), right_ptr, sizeof(int) * total_size); + res.memcpy(left.data(), left_ptr, sizeof(int) * total_size); + + for (int r = 0; r < n; ++r) + { + for (int b = 0; b < batch_size; ++b) + { + expect_right[b + batch_size * r] = 100 * r + b; + expect_left[r + n * b] = 100 * r + b; + } + } + + print_linear_order("layout_right", right); + print_linear_order("layout_left ", left); + + const bool ok = (right == expect_right) && (left == expect_left); + + std::cout << " result -- " << (ok ? "PASS" : "FAIL") << '\n'; + + res.deallocate(right_ptr); + res.deallocate(left_ptr); + + std::cout << "\nDONE!...\n"; + return ok ? 0 : 1; +} diff --git a/include/RAJA/RAJA.hpp b/include/RAJA/RAJA.hpp index 343d552f76..f367bd0555 100644 --- a/include/RAJA/RAJA.hpp +++ b/include/RAJA/RAJA.hpp @@ -124,6 +124,7 @@ #include "RAJA/util/StaticLayout.hpp" #include "RAJA/util/IndexLayout.hpp" #include "RAJA/util/View.hpp" +#include "RAJA/pattern/forall_nd.hpp" // diff --git a/include/RAJA/pattern/forall_nd.hpp b/include/RAJA/pattern/forall_nd.hpp new file mode 100644 index 0000000000..10997882e7 --- /dev/null +++ b/include/RAJA/pattern/forall_nd.hpp @@ -0,0 +1,133 @@ +/*! + ****************************************************************************** + * + * \file + * + * \brief Header file for flattened N-D forall helpers built on + *dynamic_forall. + * + ****************************************************************************** + */ + +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) Lawrence Livermore National Security, LLC and other +// RAJA Project Developers. See top-level LICENSE and COPYRIGHT +// files for dates and other details. No copyright assignment is required +// to contribute to RAJA. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#ifndef RAJA_pattern_forall_nd_HPP +#define RAJA_pattern_forall_nd_HPP + +#include + +#include "RAJA/pattern/forall.hpp" +#include "RAJA/util/CombiningAdapter.hpp" +#include "RAJA/util/View.hpp" + +namespace RAJA +{ + +template +struct TypedRangeSegmentPack +{ + camp::tuple...> data; +}; + +template +RAJA_INLINE auto segments(RAJA::TypedRangeSegment const&... segs) +{ + return TypedRangeSegmentPack {camp::make_tuple(segs...)}; +} + +namespace detail +{ + +template +struct reverse_idx_seq; + +template +struct reverse_idx_seq> +{ + using type = camp::idx_seq; +}; + +template +struct layout_to_permutation; + +template +struct layout_to_permutation +{ + using type = camp::make_idx_seq_t; +}; + +template +struct layout_to_permutation +{ + using type = typename reverse_idx_seq>::type; +}; + +template +RAJA_INLINE auto make_forall_nd_adapter( + Lambda&& body, + TypedRangeSegmentPack const& segs) +{ + using perm = + typename layout_to_permutation::type; + + return camp::apply( + [&](auto const&... unpacked_segs) { + return RAJA::make_PermutedCombiningAdapter( + std::forward(body), unpacked_segs...); + }, + segs.data); +} + +} // namespace detail + +template +void forall_nd(int pol, + TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + if (pol < 0) + { + RAJA_ABORT_OR_THROW("Policy value out of range"); + } + + auto adapter = detail::make_forall_nd_adapter( + std::forward(body), segs); + + RAJA::dynamic_forall(pol, adapter.getRange(), adapter); +} + +template +resources::EventProxy forall_nd( + RAJA::resources::Resource resource, + int pol, + TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + if (pol < 0) + { + RAJA_ABORT_OR_THROW("Policy value out of range"); + } + + auto adapter = detail::make_forall_nd_adapter( + std::forward(body), segs); + + return RAJA::dynamic_forall(resource, pol, adapter.getRange(), + adapter); +} + +} // namespace RAJA + +#endif /* RAJA_pattern_forall_nd_HPP */ diff --git a/test/functional/util/CMakeLists.txt b/test/functional/util/CMakeLists.txt index 10dfcba0c9..9f901605f5 100644 --- a/test/functional/util/CMakeLists.txt +++ b/test/functional/util/CMakeLists.txt @@ -30,3 +30,7 @@ raja_add_test( raja_add_test( NAME test-PermutedCombiningAdapter-3D SOURCES test-PermutedCombiningAdapter-3D.cpp) + +raja_add_test( + NAME test-forall_nd + SOURCES test-forall_nd.cpp) diff --git a/test/functional/util/test-forall_nd.cpp b/test/functional/util/test-forall_nd.cpp new file mode 100644 index 0000000000..75df6b2ad7 --- /dev/null +++ b/test/functional/util/test-forall_nd.cpp @@ -0,0 +1,68 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) Lawrence Livermore National Security, LLC and other +// RAJA Project Developers. See top-level LICENSE and COPYRIGHT +// files for dates and other details. No copyright assignment is required +// to contribute to RAJA. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#include "RAJA_test-base.hpp" + +#include + +namespace +{ + +using policy_list = camp::list; + +} // namespace + +TEST(forall_nd, layout_right_default_resource) +{ + constexpr int n = 3; + constexpr int batch_size = 4; + + std::vector values(n * batch_size, -1); + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + + RAJA::forall_nd( + 0, RAJA::segments(rows, batches), [&](int r, int b) { + values[b + batch_size * r] = 100 * r + b; + }); + + for (int r = 0; r < n; ++r) + { + for (int b = 0; b < batch_size; ++b) + { + ASSERT_EQ(values[b + batch_size * r], 100 * r + b); + } + } +} + +TEST(forall_nd, layout_left_resource) +{ + constexpr int n = 3; + constexpr int batch_size = 4; + + std::vector values(n * batch_size, -1); + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + + RAJA::resources::Host host_res; + RAJA::resources::Resource res(host_res); + + RAJA::forall_nd( + res, 0, RAJA::segments(rows, batches), [&](int r, int b) { + values[r + n * b] = 100 * r + b; + }); + + for (int r = 0; r < n; ++r) + { + for (int b = 0; b < batch_size; ++b) + { + ASSERT_EQ(values[r + n * b], 100 * r + b); + } + } +} From 7e6c77ae6573fab43202e242116b9dc6f53298f8 Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Thu, 30 Apr 2026 16:02:55 -0700 Subject: [PATCH 2/5] draft of launch_nd --- docs/sphinx/user_guide/cook_book.rst | 2 +- .../sphinx/user_guide/cook_book/launch-nd.rst | 116 +++++ docs/sphinx/user_guide/feature/loop_basic.rst | 58 --- examples/CMakeLists.txt | 4 +- examples/forall_nd_layout.cpp | 197 -------- examples/launch_nd.cpp | 175 +++++++ include/RAJA/RAJA.hpp | 2 +- include/RAJA/pattern/forall_nd.hpp | 133 ----- include/RAJA/pattern/launch_nd.hpp | 454 ++++++++++++++++++ test/functional/util/CMakeLists.txt | 4 +- test/functional/util/test-forall_nd.cpp | 68 --- test/functional/util/test-launch_nd.cpp | 102 ++++ 12 files changed, 853 insertions(+), 462 deletions(-) create mode 100644 docs/sphinx/user_guide/cook_book/launch-nd.rst delete mode 100644 examples/forall_nd_layout.cpp create mode 100644 examples/launch_nd.cpp delete mode 100644 include/RAJA/pattern/forall_nd.hpp create mode 100644 include/RAJA/pattern/launch_nd.hpp delete mode 100644 test/functional/util/test-forall_nd.cpp create mode 100644 test/functional/util/test-launch_nd.cpp diff --git a/docs/sphinx/user_guide/cook_book.rst b/docs/sphinx/user_guide/cook_book.rst index 5c7e5b8275..1034fb45a2 100644 --- a/docs/sphinx/user_guide/cook_book.rst +++ b/docs/sphinx/user_guide/cook_book.rst @@ -22,4 +22,4 @@ to provide users with complete beyond usage examples beyond what can be found in cook_book/reduction cook_book/multi-reduction - + cook_book/launch-nd diff --git a/docs/sphinx/user_guide/cook_book/launch-nd.rst b/docs/sphinx/user_guide/cook_book/launch-nd.rst new file mode 100644 index 0000000000..22711ab3dc --- /dev/null +++ b/docs/sphinx/user_guide/cook_book/launch-nd.rst @@ -0,0 +1,116 @@ +.. ## +.. ## Copyright (c) Lawrence Livermore National Security, LLC and other +.. ## RAJA Project Developers. See top-level LICENSE and COPYRIGHT +.. ## files for dates and other details. No copyright assignment is required +.. ## to contribute to RAJA. +.. ## +.. ## SPDX-License-Identifier: (BSD-3-Clause) +.. ## + +.. _cook-book-launch-nd-label: + +============================ +Cooking with RAJA::launch_nd +============================ + +``RAJA::launch_nd`` runs a logical 2-D or 3-D loop body through selectable +launch-backed mappings. It is intended for kernels where the source code should +stay written in logical multi-dimensional indices, but the best GPU mapping is +not known from the source alone. + +The interface supports two mapping policy families: + + * ``RAJA::launch_nd_flattened_policy`` maps the product + of the logical dimensions to a 1-D launch. The ``ExecPolicy`` is a regular + forall-style policy such as ``RAJA::cuda_exec<256>`` or + ``RAJA::hip_exec<256>``. RAJA derives the launch parameters and performs the + linear-to-logical index reconstruction internally. + * ``RAJA::launch_nd_grid_policy`` maps the + logical dimensions directly to a true 2-D or 3-D launch. The user supplies + the launch policy, one loop policy per segment, and the ``RAJA::LaunchParams``. + +This makes it possible to put both mappings behind one abstraction and choose +the mapping policy from problem size, backend, or measurements while preserving +one logical loop body. + +---------------------------------- +Why Compare Flat and Grid Mappings +---------------------------------- + +Many application kernels are logically 2-D or 3-D but have historically run on +a 1-D iteration space, with division and modulo operations used to recover the +logical indices. That approach can expose more parallelism when one logical +dimension is small. + +For example, a ``cells x components`` kernel with only a few components may not +fit a fixed ``16 x 16`` block well. A true 2-D grid can leave many component +threads inactive in each block. The flattened mapping instead launches over +``cells * components`` as a 1-D space, which can produce fuller blocks. The +tradeoff is the extra index reconstruction arithmetic. The true grid mapping +can still win when the dimensions fit the block shape, when the body is very +small and index reconstruction dominates, or when direct multi-dimensional +mapping improves memory access or scheduling. + +---------------- +Mapping Policies +---------------- + +The example source defines backend-specific policy aliases. CUDA uses direct +global loop policies for the true grid mapping: + +.. literalinclude:: ../../../../examples/launch_nd.cpp + :start-after: // _launch_nd_policy_aliases_start + :end-before: // _launch_nd_policy_aliases_end + :language: C++ + +The flattened policy uses a forall-style execution policy such as +``RAJA::cuda_exec`` or ``RAJA::hip_exec``. The +grid policy uses ``RAJA::LaunchPolicy`` and direct global loop policies such as +``RAJA::cuda_global_y_direct`` and ``RAJA::cuda_global_x_direct``. + +Both mappings call the same logical body through one ``RAJA::launch_nd`` call: + +.. literalinclude:: ../../../../examples/launch_nd.cpp + :start-after: // _launch_nd_call_start + :end-before: // _launch_nd_call_end + :language: C++ + +---------------------- +Runtime Policy Choice +---------------------- + +The mapping can be selected at run time by choosing which policy object is +passed to the common implementation: + +.. literalinclude:: ../../../../examples/launch_nd.cpp + :start-after: // _launch_nd_runtime_select_start + :end-before: // _launch_nd_runtime_select_end + :language: C++ + +Run the example both ways and compare timing with the profiling tool normally +used for the target backend:: + + ./launch_nd flat + ./launch_nd grid + +The example uses ``num_cells = 257`` and ``num_comp = 5`` to show a case where +the ``16 x 16`` grid mapping has a small logical component dimension. Change +``num_cells``, ``num_comp``, ``block_size_1d``, ``block_x``, and ``block_y`` in +``RAJA/examples/launch_nd.cpp`` to explore when the flattened or true-grid +mapping is better for a kernel shape. + +-------------------- +Current Capabilities +-------------------- + +``RAJA::launch_nd`` currently supports ``RAJA::TypedRangeSegment`` packs created +with ``RAJA::segments``. The grid mapping supports 2-D and 3-D loops and +requires one loop policy per segment. The flattened mapping uses +``RAJA::layout_right`` by default, or ``RAJA::layout_left`` when that layout tag +is supplied, to control which logical index is unit stride in the flattened +space. + +Use direct ``RAJA::launch`` when the kernel needs explicit shared memory, team +synchronization, multiple cooperating loops in a single launch body, or other +hierarchical launch features that do not fit the single logical body accepted by +``RAJA::launch_nd``. diff --git a/docs/sphinx/user_guide/feature/loop_basic.rst b/docs/sphinx/user_guide/feature/loop_basic.rst index 098ecf7a9f..830f92c744 100644 --- a/docs/sphinx/user_guide/feature/loop_basic.rst +++ b/docs/sphinx/user_guide/feature/loop_basic.rst @@ -412,61 +412,3 @@ the provided lambda with the appropriate indices. .. note:: CombiningAdapter currently only supports ``RAJA::TypedRangeSegment`` segments. - -.. _loop_elements-forall_nd-label: - --------------------------------------------------------- -Flattened multi-dimensional loops with ``RAJA::forall_nd`` --------------------------------------------------------- - -``RAJA::forall_nd`` is a convenience wrapper built on -``RAJA::dynamic_forall`` and ``RAJA::CombiningAdapter``. It executes a -flattened 1-D iteration space selected from a run-time policy list and -reconstructs the multi-dimensional indices before invoking the user lambda. -On CUDA, HIP, and SYCL backends this means the logical N-D iteration space is -flattened to a single 1-D device loop of length equal to the product of the -segment sizes. For example, ``cuda_exec``, ``hip_exec``, -and ``sycl_exec`` execute a 1-D kernel/work-group shape over that -flattened range, and ``RAJA::forall_nd`` performs the index reconstruction -before calling the user lambda. - -The API takes a policy list, a run-time policy index, a pack of -``RAJA::TypedRangeSegment`` objects created by ``RAJA::segments``, and a -lambda whose arguments are the recovered logical indices:: - - using policy_list = camp::list - #endif - #if defined(RAJA_ENABLE_HIP) - , RAJA::hip_exec<256> - #endif - #if defined(RAJA_ENABLE_SYCL) - , RAJA::sycl_exec<256> - #endif - >; - - auto rows = RAJA::TypedRangeSegment(0, n); - auto batches = RAJA::TypedRangeSegment(0, batch_size); - -The layout tag controls which logical index is unit stride in the flattened -space. ``RAJA::layout_right`` makes the right-most index vary fastest: - -.. literalinclude:: ../../../../examples/forall_nd_layout.cpp - :start-after: // _forall_nd_right_start - :end-before: // _forall_nd_right_end - :language: C++ - -``RAJA::layout_left`` makes the left-most index vary fastest: - -.. literalinclude:: ../../../../examples/forall_nd_layout.cpp - :start-after: // _forall_nd_left_start - :end-before: // _forall_nd_left_end - :language: C++ - -The example source file ``RAJA/examples/forall_nd_layout.cpp`` contains a -complete runnable example. diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 19b132d878..c6c590f47a 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -62,8 +62,8 @@ raja_add_executable( SOURCES launch_flatten.cpp) raja_add_executable( - NAME forall_nd_layout - SOURCES forall_nd_layout.cpp) + NAME launch_nd + SOURCES launch_nd.cpp) raja_add_executable( NAME launch_reductions diff --git a/examples/forall_nd_layout.cpp b/examples/forall_nd_layout.cpp deleted file mode 100644 index ffe680c70e..0000000000 --- a/examples/forall_nd_layout.cpp +++ /dev/null @@ -1,197 +0,0 @@ -//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// -// Copyright (c) Lawrence Livermore National Security, LLC and other -// RAJA Project Developers. See top-level LICENSE and COPYRIGHT -// files for dates and other details. No copyright assignment is required -// to contribute to RAJA. -// -// SPDX-License-Identifier: (BSD-3-Clause) -//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// - -#include -#include -#include - -#include "RAJA/RAJA.hpp" - -/* - * Flattened N-D forall helper with layout-left/layout-right selection - * - * This example shows the public RAJA::forall_nd wrapper, which executes a - * flattened 1-D kernel selected from a run-time policy list and reconstructs - * the multi-dimensional indices for the user lambda. - * - * The public call shape is: - * - * RAJA::forall_nd( - * res, pol, RAJA::segments(rows, batches), - * [=] RAJA_HOST_DEVICE(int r, int b) { ... }); - * - * Supported layouts: - * - RAJA::layout_right : right-most index varies fastest - * - RAJA::layout_left : left-most index varies fastest - */ - -namespace -{ - -using policy_list = camp::list< - RAJA::seq_exec, - RAJA::simd_exec -#if defined(RAJA_ENABLE_OPENMP) - , - RAJA::omp_parallel_for_exec -#endif -#if defined(RAJA_ENABLE_CUDA) - , - RAJA::cuda_exec<256> -#endif -#if defined(RAJA_ENABLE_HIP) - , - RAJA::hip_exec<256> -#endif -#if defined(RAJA_ENABLE_SYCL) - , - RAJA::sycl_exec<256> -#endif - >; - -constexpr int host_policy_count = - 2 -#if defined(RAJA_ENABLE_OPENMP) - + 1 -#endif - ; - -RAJA::resources::Resource get_resource_for_policy(int pol) -{ - if (pol < 0) - { - RAJA_ABORT_OR_THROW("Policy value out of range for this build"); - } - - RAJA::resources::Host host_res; - if (pol < host_policy_count) - { - return RAJA::resources::Resource(host_res); - } - -#if defined(RAJA_ENABLE_CUDA) - if (pol == host_policy_count) - { - return RAJA::resources::Resource(RAJA::resources::Cuda {}); - } -#endif - -#if defined(RAJA_ENABLE_HIP) - if (pol == host_policy_count -#if defined(RAJA_ENABLE_CUDA) - + 1 -#endif - ) - { - return RAJA::resources::Resource(RAJA::resources::Hip {}); - } -#endif - -#if defined(RAJA_ENABLE_SYCL) - if (pol == host_policy_count -#if defined(RAJA_ENABLE_CUDA) - + 1 -#endif -#if defined(RAJA_ENABLE_HIP) - + 1 -#endif - ) - { - return RAJA::resources::Resource(RAJA::resources::Sycl {}); - } -#endif - - RAJA_ABORT_OR_THROW("Policy value out of range for this build"); -} - -void print_linear_order(char const* label, std::vector const& values) -{ - std::cout << " " << label << ":"; - for (auto value : values) - { - std::cout << ' ' << value; - } - std::cout << '\n'; -} - -} // namespace - -int main(int argc, char** argv) -{ - const int pol = (argc > 1) ? std::atoi(argv[1]) : 0; - - std::cout << "\nRAJA flattened N-D forall helper example...\n"; - std::cout << "Using policy #" << pol << '\n'; - - constexpr int n = 3; - constexpr int batch_size = 4; - constexpr int total_size = n * batch_size; - - auto res = get_resource_for_policy(pol); - - int* right_ptr = res.allocate(total_size); - int* left_ptr = res.allocate(total_size); - - res.memset(right_ptr, 0xff, sizeof(int) * total_size); - res.memset(left_ptr, 0xff, sizeof(int) * total_size); - - auto rows = RAJA::TypedRangeSegment(0, n); - auto batches = RAJA::TypedRangeSegment(0, batch_size); - - // layout_right makes the right-most logical index ('b') unit stride. - // _forall_nd_right_start - RAJA::forall_nd( - res, pol, RAJA::segments(rows, batches), - [=] RAJA_HOST_DEVICE(int r, int b) { - const int flat = b + batch_size * r; - right_ptr[flat] = 100 * r + b; - }); - // _forall_nd_right_end - - // layout_left makes the left-most logical index ('r') unit stride. - // _forall_nd_left_start - RAJA::forall_nd( - res, pol, RAJA::segments(rows, batches), - [=] RAJA_HOST_DEVICE(int r, int b) { - const int flat = r + n * b; - left_ptr[flat] = 100 * r + b; - }); - // _forall_nd_left_end - - std::vector right(total_size); - std::vector left(total_size); - std::vector expect_right(total_size); - std::vector expect_left(total_size); - - res.wait(); - res.memcpy(right.data(), right_ptr, sizeof(int) * total_size); - res.memcpy(left.data(), left_ptr, sizeof(int) * total_size); - - for (int r = 0; r < n; ++r) - { - for (int b = 0; b < batch_size; ++b) - { - expect_right[b + batch_size * r] = 100 * r + b; - expect_left[r + n * b] = 100 * r + b; - } - } - - print_linear_order("layout_right", right); - print_linear_order("layout_left ", left); - - const bool ok = (right == expect_right) && (left == expect_left); - - std::cout << " result -- " << (ok ? "PASS" : "FAIL") << '\n'; - - res.deallocate(right_ptr); - res.deallocate(left_ptr); - - std::cout << "\nDONE!...\n"; - return ok ? 0 : 1; -} diff --git a/examples/launch_nd.cpp b/examples/launch_nd.cpp new file mode 100644 index 0000000000..351967857a --- /dev/null +++ b/examples/launch_nd.cpp @@ -0,0 +1,175 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) Lawrence Livermore National Security, LLC and other +// RAJA Project Developers. See top-level LICENSE and COPYRIGHT +// files for dates and other details. No copyright assignment is required +// to contribute to RAJA. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#include +#include +#include + +#include "RAJA/RAJA.hpp" + +/* + * RAJA::launch_nd mapping policies + * + * This example shows how to select a launch_nd mapping at run time while + * keeping one logical 2-D loop body. Run with: + * + * ./launch_nd flat + * ./launch_nd grid + * + * The flattened policy takes a regular RAJA forall execution policy, such as + * cuda_exec<256> or hip_exec<256>, and RAJA maps it through launch internally. + * The grid policy takes an explicit launch policy, loop policies, and + * LaunchParams for a true 2-D launch mapping. + */ + +namespace +{ + +constexpr int block_size_1d = 256; +constexpr int block_x = 16; +constexpr int block_y = 16; +constexpr int num_cells = 257; +constexpr int num_comp = 5; +constexpr int total = num_cells * num_comp; + +int ceil_div(int value, int divisor) { return (value + divisor - 1) / divisor; } + +// _launch_nd_mapping_enum_start +enum class Mapping +{ + Flat, + Grid +}; +// _launch_nd_mapping_enum_end + +// _launch_nd_policy_aliases_start +#if defined(RAJA_ENABLE_CUDA) +using launch_policy = RAJA::LaunchPolicy>; +using flat_exec = RAJA::cuda_exec; +using cell_loop = RAJA::LoopPolicy; +using comp_loop = RAJA::LoopPolicy; +using resource_type = RAJA::resources::Cuda; + +#elif defined(RAJA_ENABLE_HIP) +using launch_policy = RAJA::LaunchPolicy>; +using flat_exec = RAJA::hip_exec; +using cell_loop = RAJA::LoopPolicy; +using comp_loop = RAJA::LoopPolicy; +using resource_type = RAJA::resources::Hip; + +#else +using launch_policy = RAJA::LaunchPolicy; +using flat_exec = RAJA::seq_exec; +using cell_loop = RAJA::LoopPolicy; +using comp_loop = RAJA::LoopPolicy; +using resource_type = RAJA::resources::Host; +#endif +// _launch_nd_policy_aliases_end + +Mapping get_mapping(int argc, char** argv) +{ + if (argc < 2) + { + return Mapping::Flat; + } + + const std::string arg = argv[1]; + if (arg == "flat" || arg == "flattened") + { + return Mapping::Flat; + } + if (arg == "grid") + { + return Mapping::Grid; + } + + std::cout << "Unknown mapping '" << arg << "'. Use 'flat' or 'grid'.\n"; + return Mapping::Flat; +} + +template +int run_mapping(RAJA::resources::Resource res, + LaunchNdPolicy policy, + char const* mapping_name) +{ + int* values_ptr = res.allocate(total); + res.memset(values_ptr, 0, sizeof(int) * total); + + auto cells = RAJA::TypedRangeSegment(0, num_cells); + auto comps = RAJA::TypedRangeSegment(0, num_comp); + + // _launch_nd_call_start + RAJA::launch_nd(res, policy, RAJA::segments(cells, comps), + [=] RAJA_HOST_DEVICE(int cell, int comp) { + const int idx = comp + num_comp * cell; + values_ptr[idx] = 1000 * cell + comp; + }); + // _launch_nd_call_end + + std::vector values(total); + res.wait(); + res.memcpy(values.data(), values_ptr, sizeof(int) * total); + res.wait(); + + int errors = 0; + for (int cell = 0; cell < num_cells; ++cell) + { + for (int comp = 0; comp < num_comp; ++comp) + { + const int idx = comp + num_comp * cell; + const int expected = 1000 * cell + comp; + if (values[idx] != expected) + { + ++errors; + } + } + } + + std::cout << " mapping: " << mapping_name << '\n'; + std::cout << " logical size: " << num_cells << " x " << num_comp << '\n'; + std::cout << " result -- " << (errors == 0 ? "PASS" : "FAIL") << '\n'; + + res.deallocate(values_ptr); + return errors; +} + +} // namespace + +int main(int argc, char** argv) +{ + std::cout << "\nRAJA launch_nd mapping example...\n"; + + RAJA::resources::Resource res(resource_type {}); + const Mapping mapping = get_mapping(argc, argv); + + int errors = 0; + // _launch_nd_runtime_select_start + if (mapping == Mapping::Flat) + { + auto policy = RAJA::launch_nd_flattened_policy {}; + errors = run_mapping(res, policy, "flat"); + std::cout << " flattened launch threads per block: " << block_size_1d + << '\n'; + } + else + { + auto policy = + RAJA::launch_nd_grid_policy( + RAJA::LaunchParams(RAJA::Teams(ceil_div(num_comp, block_x), + ceil_div(num_cells, block_y)), + RAJA::Threads(block_x, block_y))); + errors = run_mapping(res, policy, "grid"); + std::cout << " grid launch block shape: " << block_x << " x " << block_y + << '\n'; + } + // _launch_nd_runtime_select_end + + std::cout << "\nDONE!...\n"; + return errors == 0 ? 0 : 1; +} diff --git a/include/RAJA/RAJA.hpp b/include/RAJA/RAJA.hpp index f367bd0555..e1675cb2ce 100644 --- a/include/RAJA/RAJA.hpp +++ b/include/RAJA/RAJA.hpp @@ -124,7 +124,7 @@ #include "RAJA/util/StaticLayout.hpp" #include "RAJA/util/IndexLayout.hpp" #include "RAJA/util/View.hpp" -#include "RAJA/pattern/forall_nd.hpp" +#include "RAJA/pattern/launch_nd.hpp" // diff --git a/include/RAJA/pattern/forall_nd.hpp b/include/RAJA/pattern/forall_nd.hpp deleted file mode 100644 index 10997882e7..0000000000 --- a/include/RAJA/pattern/forall_nd.hpp +++ /dev/null @@ -1,133 +0,0 @@ -/*! - ****************************************************************************** - * - * \file - * - * \brief Header file for flattened N-D forall helpers built on - *dynamic_forall. - * - ****************************************************************************** - */ - -//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// -// Copyright (c) Lawrence Livermore National Security, LLC and other -// RAJA Project Developers. See top-level LICENSE and COPYRIGHT -// files for dates and other details. No copyright assignment is required -// to contribute to RAJA. -// -// SPDX-License-Identifier: (BSD-3-Clause) -//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// - -#ifndef RAJA_pattern_forall_nd_HPP -#define RAJA_pattern_forall_nd_HPP - -#include - -#include "RAJA/pattern/forall.hpp" -#include "RAJA/util/CombiningAdapter.hpp" -#include "RAJA/util/View.hpp" - -namespace RAJA -{ - -template -struct TypedRangeSegmentPack -{ - camp::tuple...> data; -}; - -template -RAJA_INLINE auto segments(RAJA::TypedRangeSegment const&... segs) -{ - return TypedRangeSegmentPack {camp::make_tuple(segs...)}; -} - -namespace detail -{ - -template -struct reverse_idx_seq; - -template -struct reverse_idx_seq> -{ - using type = camp::idx_seq; -}; - -template -struct layout_to_permutation; - -template -struct layout_to_permutation -{ - using type = camp::make_idx_seq_t; -}; - -template -struct layout_to_permutation -{ - using type = typename reverse_idx_seq>::type; -}; - -template -RAJA_INLINE auto make_forall_nd_adapter( - Lambda&& body, - TypedRangeSegmentPack const& segs) -{ - using perm = - typename layout_to_permutation::type; - - return camp::apply( - [&](auto const&... unpacked_segs) { - return RAJA::make_PermutedCombiningAdapter( - std::forward(body), unpacked_segs...); - }, - segs.data); -} - -} // namespace detail - -template -void forall_nd(int pol, - TypedRangeSegmentPack const& segs, - Lambda&& body) -{ - if (pol < 0) - { - RAJA_ABORT_OR_THROW("Policy value out of range"); - } - - auto adapter = detail::make_forall_nd_adapter( - std::forward(body), segs); - - RAJA::dynamic_forall(pol, adapter.getRange(), adapter); -} - -template -resources::EventProxy forall_nd( - RAJA::resources::Resource resource, - int pol, - TypedRangeSegmentPack const& segs, - Lambda&& body) -{ - if (pol < 0) - { - RAJA_ABORT_OR_THROW("Policy value out of range"); - } - - auto adapter = detail::make_forall_nd_adapter( - std::forward(body), segs); - - return RAJA::dynamic_forall(resource, pol, adapter.getRange(), - adapter); -} - -} // namespace RAJA - -#endif /* RAJA_pattern_forall_nd_HPP */ diff --git a/include/RAJA/pattern/launch_nd.hpp b/include/RAJA/pattern/launch_nd.hpp new file mode 100644 index 0000000000..0731cc60b3 --- /dev/null +++ b/include/RAJA/pattern/launch_nd.hpp @@ -0,0 +1,454 @@ +/*! + ****************************************************************************** + * + * \file + * + * \brief Header file for N-D launch helpers. + * + ****************************************************************************** + */ + +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) Lawrence Livermore National Security, LLC and other +// RAJA Project Developers. See top-level LICENSE and COPYRIGHT +// files for dates and other details. No copyright assignment is required +// to contribute to RAJA. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#ifndef RAJA_pattern_launch_nd_HPP +#define RAJA_pattern_launch_nd_HPP + +#include +#include + +#include "RAJA/pattern/launch.hpp" +#include "RAJA/policy/sequential.hpp" +#if defined(RAJA_ENABLE_CUDA) +#include "RAJA/policy/cuda.hpp" +#endif +#if defined(RAJA_ENABLE_HIP) +#include "RAJA/policy/hip.hpp" +#endif +#if defined(RAJA_ENABLE_SYCL) +#include "RAJA/policy/sycl.hpp" +#endif +#include "RAJA/util/CombiningAdapter.hpp" +#include "RAJA/util/View.hpp" + +namespace RAJA +{ + +template +struct TypedRangeSegmentPack +{ + camp::tuple...> data; +}; + +template +RAJA_INLINE auto segments(RAJA::TypedRangeSegment const&... segs) +{ + return TypedRangeSegmentPack {camp::make_tuple(segs...)}; +} + +template +struct launch_nd_flattened_policy +{ + using exec_policy = ExecPolicy; + using layout_tag = LayoutTag; +}; + +template +struct launch_nd_grid_policy +{ + using launch_policy = LaunchPolicy; + using loop_policies = camp::list; + + LaunchParams launch_params; + + explicit launch_nd_grid_policy(LaunchParams params) : launch_params(params) + { + static_assert(sizeof...(LoopPolicies) == 2 || sizeof...(LoopPolicies) == 3, + "RAJA::launch_nd launch backend supports 2D or 3D loops"); + } +}; + +namespace detail +{ + +template +struct reverse_idx_seq; + +template +struct reverse_idx_seq> +{ + using type = camp::idx_seq; +}; + +template +struct layout_to_permutation; + +template +struct layout_to_permutation +{ + using type = camp::make_idx_seq_t; +}; + +template +struct layout_to_permutation +{ + using type = typename reverse_idx_seq>::type; +}; + +template +RAJA_INLINE auto make_launch_nd_adapter( + Lambda&& body, + TypedRangeSegmentPack const& segs) +{ + using perm = + typename layout_to_permutation::type; + + return camp::apply( + [&](auto const&... unpacked_segs) { + return RAJA::make_PermutedCombiningAdapter( + std::forward(body), unpacked_segs...); + }, + segs.data); +} + +template +struct launch_nd_flattened_launch_traits; + +template<> +struct launch_nd_flattened_launch_traits +{ + using launch_policy = RAJA::LaunchPolicy; + using loop_policy = RAJA::LoopPolicy; + + template + static LaunchParams make_launch_params(SizeT) + { + return LaunchParams {}; + } +}; + +#if defined(RAJA_CUDA_ACTIVE) +template +struct launch_nd_flattened_launch_traits< + RAJA::policy::cuda::cuda_exec_explicit> +{ + using launch_policy = RAJA::LaunchPolicy>; + using loop_policy = RAJA::LoopPolicy; + + template + static LaunchParams make_launch_params(SizeT size) + { + constexpr int block_size = IterationGetter::block_size; + static_assert(block_size > 0, + "RAJA::launch_nd flattened launch requires an execution " + "policy with a fixed block size"); + + constexpr int grid_size = IterationGetter::grid_size; + const int teams = + (grid_size == RAJA::named_usage::unspecified) + ? RAJA_DIVIDE_CEILING_INT(static_cast(size), block_size) + : grid_size; + + return LaunchParams(RAJA::Teams(teams), RAJA::Threads(block_size)); + } +}; +#endif + +#if defined(RAJA_HIP_ACTIVE) +template +struct launch_nd_flattened_launch_traits< + RAJA::policy::hip:: + hip_exec> +{ + using launch_policy = RAJA::LaunchPolicy>; + using loop_policy = RAJA::LoopPolicy; + + template + static LaunchParams make_launch_params(SizeT size) + { + constexpr int block_size = IterationGetter::block_size; + static_assert(block_size > 0, + "RAJA::launch_nd flattened launch requires an execution " + "policy with a fixed block size"); + + constexpr int grid_size = IterationGetter::grid_size; + const int teams = + (grid_size == RAJA::named_usage::unspecified) + ? RAJA_DIVIDE_CEILING_INT(static_cast(size), block_size) + : grid_size; + + return LaunchParams(RAJA::Teams(teams), RAJA::Threads(block_size)); + } +}; +#endif + +#if defined(RAJA_SYCL_ACTIVE) +template +struct launch_nd_flattened_launch_traits< + RAJA::policy::sycl::sycl_exec> +{ + using launch_policy = RAJA::LaunchPolicy>; + using loop_policy = RAJA::LoopPolicy>; + + template + static LaunchParams make_launch_params(SizeT size) + { + return LaunchParams( + RAJA::Teams(RAJA_DIVIDE_CEILING_INT(static_cast(size), + static_cast(BlockSize))), + RAJA::Threads(static_cast(BlockSize))); + } +}; +#endif + +template +void launch_nd_flattened_execute(TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + auto adapter = + make_launch_nd_adapter(std::forward(body), segs); + + using traits = launch_nd_flattened_launch_traits; + using launch_policy = typename traits::launch_policy; + using loop_policy = typename traits::loop_policy; + + RAJA::launch(traits::make_launch_params(adapter.size()), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) { + RAJA::loop(ctx, adapter.getRange(), + adapter); + }); +} + +template +resources::EventProxy launch_nd_flattened_execute( + RAJA::resources::Resource resource, + TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + auto adapter = + make_launch_nd_adapter(std::forward(body), segs); + + using traits = launch_nd_flattened_launch_traits; + using launch_policy = typename traits::launch_policy; + using loop_policy = typename traits::loop_policy; + + return RAJA::launch( + resource, traits::make_launch_params(adapter.size()), + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) { + RAJA::loop(ctx, adapter.getRange(), adapter); + }); +} + +template +void launch_nd_grid_execute(LaunchParams const& launch_params, + TypedRangeSegmentPack const& segs, + Lambda&& body, + camp::num<2>) +{ + using loop0 = typename camp::at>::type; + using loop1 = typename camp::at>::type; + + auto const seg0 = camp::get<0>(segs.data); + auto const seg1 = camp::get<1>(segs.data); + auto user_body = std::forward(body); + + RAJA::launch(launch_params, + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) { + RAJA::loop(ctx, seg0, [&](Idx0 i0) { + RAJA::loop(ctx, seg1, [&](Idx1 i1) { + user_body(i0, i1); + }); + }); + }); +} + +template +void launch_nd_grid_execute(LaunchParams const& launch_params, + TypedRangeSegmentPack const& segs, + Lambda&& body, + camp::num<3>) +{ + using loop0 = typename camp::at>::type; + using loop1 = typename camp::at>::type; + using loop2 = typename camp::at>::type; + + auto const seg0 = camp::get<0>(segs.data); + auto const seg1 = camp::get<1>(segs.data); + auto const seg2 = camp::get<2>(segs.data); + auto user_body = std::forward(body); + + RAJA::launch(launch_params, + [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) { + RAJA::loop(ctx, seg0, [&](Idx0 i0) { + RAJA::loop(ctx, seg1, [&](Idx1 i1) { + RAJA::loop(ctx, seg2, [&](Idx2 i2) { + user_body(i0, i1, i2); + }); + }); + }); + }); +} + +template +resources::EventProxy launch_nd_grid_execute( + RAJA::resources::Resource resource, + LaunchParams const& launch_params, + TypedRangeSegmentPack const& segs, + Lambda&& body, + camp::num<2>) +{ + using loop0 = typename camp::at>::type; + using loop1 = typename camp::at>::type; + + auto const seg0 = camp::get<0>(segs.data); + auto const seg1 = camp::get<1>(segs.data); + auto user_body = std::forward(body); + + return RAJA::launch( + resource, launch_params, [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) { + RAJA::loop(ctx, seg0, [&](Idx0 i0) { + RAJA::loop(ctx, seg1, [&](Idx1 i1) { + user_body(i0, i1); + }); + }); + }); +} + +template +resources::EventProxy launch_nd_grid_execute( + RAJA::resources::Resource resource, + LaunchParams const& launch_params, + TypedRangeSegmentPack const& segs, + Lambda&& body, + camp::num<3>) +{ + using loop0 = typename camp::at>::type; + using loop1 = typename camp::at>::type; + using loop2 = typename camp::at>::type; + + auto const seg0 = camp::get<0>(segs.data); + auto const seg1 = camp::get<1>(segs.data); + auto const seg2 = camp::get<2>(segs.data); + auto user_body = std::forward(body); + + return RAJA::launch( + resource, launch_params, [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) { + RAJA::loop(ctx, seg0, [&](Idx0 i0) { + RAJA::loop(ctx, seg1, [&](Idx1 i1) { + RAJA::loop(ctx, seg2, [&](Idx2 i2) { + user_body(i0, i1, i2); + }); + }); + }); + }); +} + +} // namespace detail + +template +void launch_nd(launch_nd_flattened_policy, + TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + detail::launch_nd_flattened_execute( + segs, std::forward(body)); +} + +template +resources::EventProxy launch_nd( + RAJA::resources::Resource resource, + launch_nd_flattened_policy, + TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + return detail::launch_nd_flattened_execute( + resource, segs, std::forward(body)); +} + +template +void launch_nd(launch_nd_grid_policy policy, + TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + static_assert(sizeof...(IdxTs) == sizeof...(LoopPolicies), + "RAJA::launch_nd launch backend requires one loop policy per " + "segment"); + + detail::launch_nd_grid_execute>( + policy.launch_params, segs, std::forward(body), + camp::num {}); +} + +template +resources::EventProxy launch_nd( + RAJA::resources::Resource resource, + launch_nd_grid_policy policy, + TypedRangeSegmentPack const& segs, + Lambda&& body) +{ + static_assert(sizeof...(IdxTs) == sizeof...(LoopPolicies), + "RAJA::launch_nd launch backend requires one loop policy per " + "segment"); + + return detail::launch_nd_grid_execute>( + resource, policy.launch_params, segs, std::forward(body), + camp::num {}); +} + +} // namespace RAJA + +#endif /* RAJA_pattern_launch_nd_HPP */ diff --git a/test/functional/util/CMakeLists.txt b/test/functional/util/CMakeLists.txt index 9f901605f5..d7c9d5cc84 100644 --- a/test/functional/util/CMakeLists.txt +++ b/test/functional/util/CMakeLists.txt @@ -32,5 +32,5 @@ raja_add_test( SOURCES test-PermutedCombiningAdapter-3D.cpp) raja_add_test( - NAME test-forall_nd - SOURCES test-forall_nd.cpp) + NAME test-launch_nd + SOURCES test-launch_nd.cpp) diff --git a/test/functional/util/test-forall_nd.cpp b/test/functional/util/test-forall_nd.cpp deleted file mode 100644 index 75df6b2ad7..0000000000 --- a/test/functional/util/test-forall_nd.cpp +++ /dev/null @@ -1,68 +0,0 @@ -//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// -// Copyright (c) Lawrence Livermore National Security, LLC and other -// RAJA Project Developers. See top-level LICENSE and COPYRIGHT -// files for dates and other details. No copyright assignment is required -// to contribute to RAJA. -// -// SPDX-License-Identifier: (BSD-3-Clause) -//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// - -#include "RAJA_test-base.hpp" - -#include - -namespace -{ - -using policy_list = camp::list; - -} // namespace - -TEST(forall_nd, layout_right_default_resource) -{ - constexpr int n = 3; - constexpr int batch_size = 4; - - std::vector values(n * batch_size, -1); - auto rows = RAJA::TypedRangeSegment(0, n); - auto batches = RAJA::TypedRangeSegment(0, batch_size); - - RAJA::forall_nd( - 0, RAJA::segments(rows, batches), [&](int r, int b) { - values[b + batch_size * r] = 100 * r + b; - }); - - for (int r = 0; r < n; ++r) - { - for (int b = 0; b < batch_size; ++b) - { - ASSERT_EQ(values[b + batch_size * r], 100 * r + b); - } - } -} - -TEST(forall_nd, layout_left_resource) -{ - constexpr int n = 3; - constexpr int batch_size = 4; - - std::vector values(n * batch_size, -1); - auto rows = RAJA::TypedRangeSegment(0, n); - auto batches = RAJA::TypedRangeSegment(0, batch_size); - - RAJA::resources::Host host_res; - RAJA::resources::Resource res(host_res); - - RAJA::forall_nd( - res, 0, RAJA::segments(rows, batches), [&](int r, int b) { - values[r + n * b] = 100 * r + b; - }); - - for (int r = 0; r < n; ++r) - { - for (int b = 0; b < batch_size; ++b) - { - ASSERT_EQ(values[r + n * b], 100 * r + b); - } - } -} diff --git a/test/functional/util/test-launch_nd.cpp b/test/functional/util/test-launch_nd.cpp new file mode 100644 index 0000000000..9cf1de9e83 --- /dev/null +++ b/test/functional/util/test-launch_nd.cpp @@ -0,0 +1,102 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// Copyright (c) Lawrence Livermore National Security, LLC and other +// RAJA Project Developers. See top-level LICENSE and COPYRIGHT +// files for dates and other details. No copyright assignment is required +// to contribute to RAJA. +// +// SPDX-License-Identifier: (BSD-3-Clause) +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// + +#include "RAJA_test-base.hpp" + +#include + +namespace +{ + +using launch_policy = RAJA::LaunchPolicy; +using flat_policy = RAJA::launch_nd_flattened_policy; +using flat_left_policy = + RAJA::launch_nd_flattened_policy; +using row_loop = RAJA::LoopPolicy; +using col_loop = RAJA::LoopPolicy; + +} // namespace + +TEST(launch_nd, layout_right_default_resource) +{ + constexpr int n = 3; + constexpr int batch_size = 4; + + std::vector values(n * batch_size, -1); + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + + RAJA::launch_nd(flat_policy {}, RAJA::segments(rows, batches), + [&](int r, int b) { + values[b + batch_size * r] = 100 * r + b; + }); + + for (int r = 0; r < n; ++r) + { + for (int b = 0; b < batch_size; ++b) + { + ASSERT_EQ(values[b + batch_size * r], 100 * r + b); + } + } +} + +TEST(launch_nd, layout_left_resource) +{ + constexpr int n = 3; + constexpr int batch_size = 4; + + std::vector values(n * batch_size, -1); + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + + RAJA::resources::Host host_res; + RAJA::resources::Resource res(host_res); + + RAJA::launch_nd(res, flat_left_policy {}, RAJA::segments(rows, batches), + [&](int r, int b) { + values[r + n * b] = 100 * r + b; + }); + + for (int r = 0; r < n; ++r) + { + for (int b = 0; b < batch_size; ++b) + { + ASSERT_EQ(values[r + n * b], 100 * r + b); + } + } +} + +TEST(launch_nd, grid_resource) +{ + constexpr int n = 3; + constexpr int batch_size = 4; + + std::vector values(n * batch_size, -1); + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + + RAJA::resources::Host host_res; + RAJA::resources::Resource res(host_res); + + RAJA::launch_nd( + res, + RAJA::launch_nd_grid_policy( + RAJA::LaunchParams()), + RAJA::segments(rows, batches), [&](int r, int b) { + values[b + batch_size * r] = 100 * r + b; + }); + + for (int r = 0; r < n; ++r) + { + for (int b = 0; b < batch_size; ++b) + { + ASSERT_EQ(values[b + batch_size * r], 100 * r + b); + } + } +} From 65706f2cd179299d34e751cd06926568f49749ea Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Mon, 11 May 2026 16:19:26 -0700 Subject: [PATCH 3/5] WIP benchmarking with RAJA launch_nd --- examples/launch_nd.cpp | 522 +++++++++++++++++++---- examples/launch_nd_benchmark_workflow.md | 148 +++++++ full_study.sh | 168 ++++++++ include/RAJA/pattern/launch_nd.hpp | 128 +++++- scripts/lc-builds/toss4_amdclang.sh | 21 + 5 files changed, 895 insertions(+), 92 deletions(-) create mode 100644 examples/launch_nd_benchmark_workflow.md create mode 100644 full_study.sh diff --git a/examples/launch_nd.cpp b/examples/launch_nd.cpp index 351967857a..832eb4c93d 100644 --- a/examples/launch_nd.cpp +++ b/examples/launch_nd.cpp @@ -8,24 +8,27 @@ //~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// #include +#include #include #include #include "RAJA/RAJA.hpp" /* - * RAJA::launch_nd mapping policies + * RAJA::launch_nd benchmark driver * - * This example shows how to select a launch_nd mapping at run time while - * keeping one logical 2-D loop body. Run with: + * This example extends the launch_nd mapping demo into a benchmark driver for + * flattened and device launch_nd policies. It supports single-size runs and + * multi-size studies, and uses RAJA::Name so Caliper can aggregate results by + * policy/size combination. * - * ./launch_nd flat - * ./launch_nd grid + * Typical runs: * - * The flattened policy takes a regular RAJA forall execution policy, such as - * cuda_exec<256> or hip_exec<256>, and RAJA maps it through launch internally. - * The grid policy takes an explicit launch policy, loop policies, and - * LaunchParams for a true 2-D launch mapping. + * ./launch_nd --mapping all --repetitions 50 --warmup 5 + * RAJA_CALIPER=1 CALI_CONFIG=runtime-report \ + * ./launch_nd --mapping all --sizes 65536x8,262144x8 + * RAJA_CALIPER=1 CALI_CONFIG=runtime-profile(output=launch_nd.cali,output.format=cali) \ + * ./launch_nd --mapping all --sizes 65536x8,262144x8,1048576x8 */ namespace @@ -34,95 +37,334 @@ namespace constexpr int block_size_1d = 256; constexpr int block_x = 16; constexpr int block_y = 16; -constexpr int num_cells = 257; -constexpr int num_comp = 5; -constexpr int total = num_cells * num_comp; +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) int ceil_div(int value, int divisor) { return (value + divisor - 1) / divisor; } +#endif -// _launch_nd_mapping_enum_start enum class Mapping { Flat, - Grid + Global, + Block, + ThreadLocal, + All +}; + +struct ProblemSize +{ + int cells = 0; + int components = 0; +}; + +struct Options +{ + Mapping mapping = Mapping::All; + int warmup = 5; + int repetitions = 50; + std::vector sizes = {}; +}; + +struct RunResult +{ + int errors = 0; }; -// _launch_nd_mapping_enum_end -// _launch_nd_policy_aliases_start #if defined(RAJA_ENABLE_CUDA) using launch_policy = RAJA::LaunchPolicy>; using flat_exec = RAJA::cuda_exec; -using cell_loop = RAJA::LoopPolicy; -using comp_loop = RAJA::LoopPolicy; +using global_cell_loop = RAJA::LoopPolicy; +using global_comp_loop = RAJA::LoopPolicy; +using block_cell_loop = RAJA::LoopPolicy; +using block_comp_loop = RAJA::LoopPolicy; +using thread_cell_loop = RAJA::LoopPolicy; +using thread_comp_loop = RAJA::LoopPolicy; using resource_type = RAJA::resources::Cuda; #elif defined(RAJA_ENABLE_HIP) using launch_policy = RAJA::LaunchPolicy>; using flat_exec = RAJA::hip_exec; -using cell_loop = RAJA::LoopPolicy; -using comp_loop = RAJA::LoopPolicy; +using global_cell_loop = RAJA::LoopPolicy; +using global_comp_loop = RAJA::LoopPolicy; +using block_cell_loop = RAJA::LoopPolicy; +using block_comp_loop = RAJA::LoopPolicy; +using thread_cell_loop = RAJA::LoopPolicy; +using thread_comp_loop = RAJA::LoopPolicy; using resource_type = RAJA::resources::Hip; #else using launch_policy = RAJA::LaunchPolicy; using flat_exec = RAJA::seq_exec; -using cell_loop = RAJA::LoopPolicy; -using comp_loop = RAJA::LoopPolicy; +using global_cell_loop = RAJA::LoopPolicy; +using global_comp_loop = RAJA::LoopPolicy; +using block_cell_loop = RAJA::LoopPolicy; +using block_comp_loop = RAJA::LoopPolicy; +using thread_cell_loop = RAJA::LoopPolicy; +using thread_comp_loop = RAJA::LoopPolicy; using resource_type = RAJA::resources::Host; #endif -// _launch_nd_policy_aliases_end -Mapping get_mapping(int argc, char** argv) +const char* mapping_name(Mapping mapping) { - if (argc < 2) + switch (mapping) { - return Mapping::Flat; + case Mapping::Flat: + return "flat"; + case Mapping::Global: + return "global"; + case Mapping::Block: + return "block"; + case Mapping::ThreadLocal: + return "thread_local"; + case Mapping::All: + return "all"; } - const std::string arg = argv[1]; - if (arg == "flat" || arg == "flattened") + return "unknown"; +} + +ProblemSize default_problem_size() +{ + return ProblemSize {262144, 8}; +} + +void print_usage(const char* executable) +{ + std::cout + << "Usage: " << executable + << " [flat|global|block|thread_local|all] [options]\n" + << "Options:\n" + << " --mapping \n" + << " Select policy set to benchmark.\n" + << " --sizes Problem sizes, e.g. 262144x8 or" + << " 65536x8,262144x8.\n" + << " --warmup Warmup launches per mapping/size.\n" + << " --repetitions Timed launches per mapping/size.\n" + << " --help Show this message.\n"; +} + +Mapping parse_mapping(const std::string& value) +{ + if (value == "flat" || value == "flattened") { return Mapping::Flat; } - if (arg == "grid") + if (value == "global" || value == "grid") + { + return Mapping::Global; + } + if (value == "block") + { + return Mapping::Block; + } + if (value == "thread" || value == "thread_local" || + value == "thread-local") { - return Mapping::Grid; + return Mapping::ThreadLocal; + } + if (value == "all") + { + return Mapping::All; } - std::cout << "Unknown mapping '" << arg << "'. Use 'flat' or 'grid'.\n"; - return Mapping::Flat; + throw std::runtime_error("unknown mapping '" + value + + "' (expected flat, global, block, thread_local," + " or all)"); } -template -int run_mapping(RAJA::resources::Resource res, - LaunchNdPolicy policy, - char const* mapping_name) +int parse_positive_int(const std::string& name, const std::string& value) { - int* values_ptr = res.allocate(total); - res.memset(values_ptr, 0, sizeof(int) * total); + try + { + const int parsed = std::stoi(value); + if (parsed <= 0) + { + throw std::runtime_error(name + " must be greater than zero"); + } + return parsed; + } + catch (const std::invalid_argument&) + { + throw std::runtime_error("invalid integer for " + name + ": '" + value + "'"); + } + catch (const std::out_of_range&) + { + throw std::runtime_error("integer out of range for " + name + ": '" + value + + "'"); + } +} + +ProblemSize parse_problem_size(const std::string& text) +{ + const std::size_t sep = text.find_first_of("xX"); + if (sep == std::string::npos) + { + throw std::runtime_error("invalid size '" + text + + "' (expected cellsxcomponents)"); + } + + return {parse_positive_int("cells", text.substr(0, sep)), + parse_positive_int("components", text.substr(sep + 1))}; +} + +std::vector parse_problem_sizes(const std::string& text) +{ + std::vector result; + std::size_t start = 0; + + while (start < text.size()) + { + const std::size_t end = text.find(',', start); + const std::string item = + text.substr(start, end == std::string::npos ? std::string::npos + : end - start); + if (item.empty()) + { + throw std::runtime_error("empty entry in --sizes"); + } + + result.push_back(parse_problem_size(item)); + + if (end == std::string::npos) + { + break; + } + start = end + 1; + } + + if (result.empty()) + { + throw std::runtime_error("--sizes requires at least one entry"); + } + + return result; +} + +const char* require_value(int& index, int argc, char** argv, const char* option) +{ + if (index + 1 >= argc) + { + throw std::runtime_error(std::string(option) + " requires a value"); + } + + return argv[++index]; +} + +Options parse_options(int argc, char** argv) +{ + Options options; + + for (int i = 1; i < argc; ++i) + { + const std::string arg = argv[i]; + + if (arg == "--help" || arg == "-h") + { + print_usage(argv[0]); + std::exit(0); + } + else if (arg == "--mapping") + { + options.mapping = parse_mapping(require_value(i, argc, argv, "--mapping")); + } + else if (arg == "--sizes") + { + options.sizes = parse_problem_sizes(require_value(i, argc, argv, "--sizes")); + } + else if (arg == "--warmup") + { + options.warmup = + parse_positive_int("--warmup", require_value(i, argc, argv, "--warmup")); + } + else if (arg == "--repetitions") + { + options.repetitions = parse_positive_int("--repetitions", + require_value(i, argc, argv, + "--repetitions")); + } + else if (arg.rfind("--", 0) == 0) + { + throw std::runtime_error("unknown option '" + arg + "'"); + } + else + { + options.mapping = parse_mapping(arg); + } + } + + return options; +} + +std::vector selected_sizes(const Options& options) +{ + if (!options.sizes.empty()) + { + return options.sizes; + } + + return {default_problem_size()}; +} + +std::vector selected_mappings(Mapping mapping) +{ + if (mapping == Mapping::All) + { + return {Mapping::Flat, Mapping::Global, Mapping::Block, + Mapping::ThreadLocal}; + } + + return {mapping}; +} - auto cells = RAJA::TypedRangeSegment(0, num_cells); - auto comps = RAJA::TypedRangeSegment(0, num_comp); +std::string size_label(const ProblemSize& size) +{ + return std::to_string(size.cells) + "x" + std::to_string(size.components); +} + +std::string kernel_name(Mapping mapping, const ProblemSize& size) +{ + return "launch_nd_" + std::string(mapping_name(mapping)) + "_cells" + + std::to_string(size.cells) + "_comp" + std::to_string(size.components); +} - // _launch_nd_call_start - RAJA::launch_nd(res, policy, RAJA::segments(cells, comps), +template +void launch_kernel(RAJA::resources::Resource res, + LaunchNdPolicy policy, + int* values_ptr, + const ProblemSize& size, + const std::string& name) +{ + auto cell_segment = RAJA::TypedRangeSegment(0, size.cells); + auto comp_segment = RAJA::TypedRangeSegment(0, size.components); + + RAJA::launch_nd(res, policy, RAJA::segments(cell_segment, comp_segment), + RAJA::Name(name.c_str()), [=] RAJA_HOST_DEVICE(int cell, int comp) { - const int idx = comp + num_comp * cell; - values_ptr[idx] = 1000 * cell + comp; + values_ptr[comp + size.components * cell] = + 1000 * cell + comp; }); - // _launch_nd_call_end - std::vector values(total); + // Synchronize so Caliper observes completed work for each named launch. res.wait(); +} + +int verify_result(RAJA::resources::Resource res, + int* values_ptr, + const ProblemSize& size) +{ + const std::size_t total = + static_cast(size.cells) * size.components; + std::vector values(total); + res.memcpy(values.data(), values_ptr, sizeof(int) * total); res.wait(); int errors = 0; - for (int cell = 0; cell < num_cells; ++cell) + for (int cell = 0; cell < size.cells; ++cell) { - for (int comp = 0; comp < num_comp; ++comp) + for (int comp = 0; comp < size.components; ++comp) { - const int idx = comp + num_comp * cell; + const int idx = comp + size.components * cell; const int expected = 1000 * cell + comp; if (values[idx] != expected) { @@ -131,45 +373,171 @@ int run_mapping(RAJA::resources::Resource res, } } - std::cout << " mapping: " << mapping_name << '\n'; - std::cout << " logical size: " << num_cells << " x " << num_comp << '\n'; - std::cout << " result -- " << (errors == 0 ? "PASS" : "FAIL") << '\n'; + return errors; +} + +template +RunResult benchmark_mapping(RAJA::resources::Resource res, + LaunchNdPolicy policy, + const Options& options, + Mapping mapping, + const ProblemSize& size) +{ + const std::size_t total = + static_cast(size.cells) * size.components; + const std::string name = kernel_name(mapping, size); + int* values_ptr = res.allocate(total); + + for (int step = 0; step < options.warmup; ++step) + { + launch_kernel(res, policy, values_ptr, size, name); + } + + for (int rep = 0; rep < options.repetitions; ++rep) + { + launch_kernel(res, policy, values_ptr, size, name); + } + + RunResult result; + result.errors = verify_result(res, values_ptr, size); + + std::cout << " kernel: " << name << '\n' + << " logical size: " << size.cells << " x " << size.components + << '\n' + << " warmup launches: " << options.warmup << '\n' + << " timed launches: " << options.repetitions << '\n' + << " result -- " << (result.errors == 0 ? "PASS" : "FAIL") + << '\n'; res.deallocate(values_ptr); - return errors; + return result; +} + +RAJA::LaunchParams make_global_launch_params(const ProblemSize& size) +{ +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) + return RAJA::LaunchParams( + RAJA::Teams(ceil_div(size.cells, block_x), + ceil_div(size.components, block_y)), + RAJA::Threads(block_x, block_y)); +#else + RAJA_UNUSED_VAR(size); + return RAJA::LaunchParams {}; +#endif +} + +RAJA::LaunchParams make_block_launch_params(const ProblemSize& size) +{ +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) + return RAJA::LaunchParams(RAJA::Teams(size.cells, size.components), + RAJA::Threads(1, 1)); +#else + RAJA_UNUSED_VAR(size); + return RAJA::LaunchParams {}; +#endif +} + +RAJA::LaunchParams make_thread_launch_params(const ProblemSize& size) +{ +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) + RAJA_UNUSED_VAR(size); + return RAJA::LaunchParams(RAJA::Teams(1, 1), RAJA::Threads(block_x, block_y)); +#else + RAJA_UNUSED_VAR(size); + return RAJA::LaunchParams {}; +#endif } } // namespace int main(int argc, char** argv) { - std::cout << "\nRAJA launch_nd mapping example...\n"; + try + { + const Options options = parse_options(argc, argv); + const std::vector sizes = selected_sizes(options); - RAJA::resources::Resource res(resource_type {}); - const Mapping mapping = get_mapping(argc, argv); + std::cout << "\nRAJA launch_nd benchmark driver...\n"; +#if defined(RAJA_ENABLE_CALIPER) + std::cout << " Caliper support: enabled\n"; +#else + std::cout << " Caliper support: disabled\n"; +#endif + std::cout << " study sizes:"; + for (const ProblemSize& size : sizes) + { + std::cout << ' ' << size_label(size); + } + std::cout << '\n'; - int errors = 0; - // _launch_nd_runtime_select_start - if (mapping == Mapping::Flat) - { - auto policy = RAJA::launch_nd_flattened_policy {}; - errors = run_mapping(res, policy, "flat"); - std::cout << " flattened launch threads per block: " << block_size_1d - << '\n'; + RAJA::resources::Resource res(resource_type {}); + + int total_errors = 0; + + for (const ProblemSize& size : sizes) + { + std::cout << "\nStudy size " << size_label(size) << '\n'; + + for (Mapping mapping : selected_mappings(options.mapping)) + { + if (mapping == Mapping::Flat) + { + auto policy = RAJA::launch_nd_flattened_policy {}; + const RunResult result = + benchmark_mapping(res, policy, options, mapping, size); + total_errors += result.errors; + std::cout << " flattened launch threads per block: " << block_size_1d + << "\n\n"; + } + else if (mapping == Mapping::Global) + { + auto policy = + RAJA::launch_nd_grid_policy( + make_global_launch_params(size)); + const RunResult result = + benchmark_mapping(res, policy, options, mapping, size); + total_errors += result.errors; + std::cout << " global launch block shape: " << block_x << " x " + << block_y << "\n\n"; + } + else if (mapping == Mapping::Block) + { + auto policy = + RAJA::launch_nd_grid_policy( + make_block_launch_params(size)); + const RunResult result = + benchmark_mapping(res, policy, options, mapping, size); + total_errors += result.errors; + std::cout << " block launch uses one logical iteration per team" + << "\n\n"; + } + else + { + auto policy = + RAJA::launch_nd_grid_policy( + make_thread_launch_params(size)); + const RunResult result = + benchmark_mapping(res, policy, options, mapping, size); + total_errors += result.errors; + std::cout << " thread-local launch uses a single team with thread" + << " loops of shape " << block_x << " x " << block_y + << "\n\n"; + } + } + } + + std::cout << "DONE!...\n"; + return total_errors == 0 ? 0 : 1; } - else + catch (const std::exception& ex) { - auto policy = - RAJA::launch_nd_grid_policy( - RAJA::LaunchParams(RAJA::Teams(ceil_div(num_comp, block_x), - ceil_div(num_cells, block_y)), - RAJA::Threads(block_x, block_y))); - errors = run_mapping(res, policy, "grid"); - std::cout << " grid launch block shape: " << block_x << " x " << block_y - << '\n'; + std::cerr << "Error: " << ex.what() << '\n'; + return 1; } - // _launch_nd_runtime_select_end - - std::cout << "\nDONE!...\n"; - return errors == 0 ? 0 : 1; } diff --git a/examples/launch_nd_benchmark_workflow.md b/examples/launch_nd_benchmark_workflow.md new file mode 100644 index 0000000000..db26f69e55 --- /dev/null +++ b/examples/launch_nd_benchmark_workflow.md @@ -0,0 +1,148 @@ +# `launch_nd` Benchmark Workflow + +This workflow uses the updated `examples/launch_nd.cpp` benchmark driver to +compare the `RAJA::launch_nd` mappings across a size sweep using Caliper only. +Each kernel launch is labeled with `RAJA::Name`, so Caliper reports are grouped +by mapping and problem size. + +## 1. Build Caliper + +```bash +export REPO_DIR=$(pwd) +export CALIPER_SOURCE_DIR=/path/to/caliper +export CALIPER_INSTALL_PREFIX="${REPO_DIR}/../install-caliper" + +mkdir -p "${REPO_DIR}/../build-caliper" +cd "${REPO_DIR}/../build-caliper" +/usr/tce/packages/cmake/cmake-3.29.2/bin/cmake "${CALIPER_SOURCE_DIR}" \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_PREFIX="${CALIPER_INSTALL_PREFIX}" +/usr/tce/packages/cmake/cmake-3.29.2/bin/cmake --build . --target install -j +``` + +## 2. Build RAJA with Caliper support + +```bash +cd "${REPO_DIR}" +export CALIPER_DIR="${CALIPER_INSTALL_PREFIX}/share/cmake/caliper" + +mkdir -p build-raja-clang +cd build-raja-clang +/usr/tce/packages/cmake/cmake-3.29.2/bin/cmake "${REPO_DIR}" \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_C_COMPILER=/opt/cray/pe/craype/2.7.35/bin/cc \ + -DCMAKE_CXX_COMPILER=/opt/cray/pe/craype/2.7.35/bin/CC \ + -DRAJA_ENABLE_RUNTIME_PLUGINS=ON \ + -DRAJA_ENABLE_CALIPER=ON \ + -Dcaliper_DIR="${CALIPER_DIR}" +/usr/tce/packages/cmake/cmake-3.29.2/bin/cmake --build . --target launch_nd -j +``` + +The LC helper script now enforces the same Caliper-only configuration: + +```bash +cd "${REPO_DIR}" +CALIPER_DIR="${CALIPER_DIR}" ./scripts/lc-builds/toss4_amdclang.sh 6.4.3 gfx90a +``` + +## 3. Generate a Caliper text report + +```bash +cd "${REPO_DIR}/build-raja-clang" +RAJA_CALIPER=1 CALI_CONFIG=runtime-report ./bin/launch_nd \ + --mapping all \ + --sizes 65536x8,262144x8,1048576x8 \ + --warmup 5 \ + --repetitions 50 +``` + +Look for kernels like: + +- `launch_nd_flat_cells65536_comp8` +- `launch_nd_global_cells65536_comp8` +- `launch_nd_block_cells262144_comp8` +- `launch_nd_thread_local_cells1048576_comp8` + +Those names come from `RAJA::Name` and let Caliper separate each policy/size +combination in a single run. + +## 4. Generate a `.cali` profile for offline analysis + +```bash +cd "${REPO_DIR}/build-raja-clang" +RAJA_CALIPER=1 CALI_CONFIG=runtime-profile(output=launch_nd.cali,output.format=cali) \ +./bin/launch_nd \ + --mapping all \ + --sizes 65536x8,262144x8,1048576x8 \ + --warmup 5 \ + --repetitions 50 +``` + +## 5. Convert the profile into a throughput CSV + +The command below computes average time per named kernel and converts it into +logical-iteration throughput: + +```bash +"${CALIPER_INSTALL_PREFIX}/bin/cali-query" -e launch_nd.cali | awk -F, ' +BEGIN { + print "mapping,cells,components,total_iterations,avg_seconds,throughput_iterations_per_second" +} +{ + kernel_name = "" + total_sec = "" + + for (i = 1; i <= NF; ++i) { + split($i, kv, "=") + key = kv[1] + value = kv[2] + + if ((key == "loop" || key == "region" || key == "path") && + value ~ /launch_nd_/) { + kernel_name = value + } else if (key == "sum#sum#time.duration") { + total_sec = value + 0 + } + } + + if (kernel_name ~ /launch_nd_/ && + match(kernel_name, /launch_nd_(.+)_cells([0-9]+)_comp([0-9]+)/, m)) { + mapping = m[1] + cells = m[2] + 0 + comp = m[3] + 0 + total_iterations = cells * comp + avg_seconds = total_sec / 50 + throughput = total_iterations / avg_seconds + print mapping "," cells "," comp "," total_iterations "," avg_seconds "," throughput + } +}' > launch_nd_throughput.csv +``` + +`launch_nd_throughput.csv` is the file to use for plotting throughput curves. + +## 6. Plot throughput curves with `gnuplot` + +```bash +gnuplot -e " +set datafile separator ','; +set key left top; +set logscale x 2; +set xlabel 'Total logical iterations'; +set ylabel 'Throughput (iterations/s)'; +plot \ + '< awk -F, \"NR==1 || \\$1==\\\"flat\\\"\" launch_nd_throughput.csv' using 4:6 with linespoints title 'flat', \ + '< awk -F, \"NR==1 || \\$1==\\\"global\\\"\" launch_nd_throughput.csv' using 4:6 with linespoints title 'global', \ + '< awk -F, \"NR==1 || \\$1==\\\"block\\\"\" launch_nd_throughput.csv' using 4:6 with linespoints title 'block', \ + '< awk -F, \"NR==1 || \\$1==\\\"thread_local\\\"\" launch_nd_throughput.csv' using 4:6 with linespoints title 'thread_local' +" +``` + +## 7. Suggested study design + +- Keep `components` fixed and sweep `cells` first. +- Run all mappings in the same executable invocation so the environment is + identical. +- Use at least 5 warmup launches and 30-100 timed launches. +- Save both the Caliper profile and the derived throughput CSV for each build. +- Compare `flat`, `global`, `block`, and `thread_local` at each size, then + examine how the gap changes with total problem size. diff --git a/full_study.sh b/full_study.sh new file mode 100644 index 0000000000..7f36be1f13 --- /dev/null +++ b/full_study.sh @@ -0,0 +1,168 @@ +#!/usr/bin/env bash + +set -euo pipefail + +export REPO_DIR=$(pwd)/RAJA +export CALIPER_INSTALL_PREFIX=$(pwd)/install-caliper +export CALIPER_DIR="${CALIPER_INSTALL_PREFIX}/share/cmake/caliper" + +COMPILER_VERSION="${COMPILER_VERSION:-7.2.1}" +GPU_ARCH="${GPU_ARCH:-gfx942}" +BUILD_DIR="build_lc_toss4-amdclang-${COMPILER_VERSION}-${GPU_ARCH}" +SIZES="${SIZES:-65536x8,262144x8,1048576x8}" +WARMUP="${WARMUP:-5}" +REPETITIONS="${REPETITIONS:-50}" +TERMINAL_PLOT="${TERMINAL_PLOT:-0}" +TERMINAL_PLOT_SIZE="${TERMINAL_PLOT_SIZE:-120,35}" +PLOT_FORMAT="${PLOT_FORMAT:-svg}" +LAUNCH_ND_EXE="${REPO_DIR}/${BUILD_DIR}/bin/launch_nd" +LAUNCH_ND_SOURCE="${REPO_DIR}/examples/launch_nd.cpp" + +echo "Running full study with ROCm ${COMPILER_VERSION} on ${GPU_ARCH}" + +cd "${REPO_DIR}" +if [ ! -x "${LAUNCH_ND_EXE}" ]; then + CALIPER_DIR="${CALIPER_DIR}" ./scripts/lc-builds/toss4_amdclang.sh "${COMPILER_VERSION}" "${GPU_ARCH}" + + cmake --build "${BUILD_DIR}" --target launch_nd -j +elif [ "${LAUNCH_ND_SOURCE}" -nt "${LAUNCH_ND_EXE}" ]; then + echo "Rebuilding launch_nd because the source is newer than the executable" + cmake --build "${BUILD_DIR}" --target launch_nd -j +else + echo "Reusing existing build at ${BUILD_DIR}" +fi + +cd "${BUILD_DIR}" +RAJA_CALIPER=1 \ +CALI_CONFIG='runtime-profile(output=launch_nd.cali,output.format=cali)' \ +./bin/launch_nd \ + --mapping all \ + --sizes "${SIZES}" \ + --warmup "${WARMUP}" \ + --repetitions "${REPETITIONS}" + +"${CALIPER_INSTALL_PREFIX}/bin/cali-query" \ + -e \ + launch_nd.cali \ +| awk -F, -v repetitions="${REPETITIONS}" ' +BEGIN { + print "mapping,cells,components,total_iterations,avg_seconds,throughput_iterations_per_second" +} +{ + kernel_name = "" + total_sec = "" + + for (i = 1; i <= NF; ++i) { + split($i, kv, "=") + key = kv[1] + value = kv[2] + + if (key == "loop" || key == "region" || key == "path") { + if (value ~ /launch_nd_/) { + kernel_name = value + } + } else if (key == "sum#sum#time.duration") { + total_sec = value + 0 + } + } + + if (kernel_name ~ /launch_nd_/ && + match(kernel_name, /launch_nd_(.+)_cells([0-9]+)_comp([0-9]+)/, m)) { + mapping = m[1] + cells = m[2] + 0 + comp = m[3] + 0 + total_iterations = cells * comp + avg_seconds = total_sec / repetitions + throughput = total_iterations / avg_seconds + print mapping "," cells "," comp "," total_iterations "," avg_seconds "," throughput + } +}' > launch_nd_throughput.csv + +awk -F, 'NR == 1 || $1 == "flat"' launch_nd_throughput.csv > launch_nd_throughput_flat.csv +awk -F, 'NR == 1 || $1 == "global"' launch_nd_throughput.csv > launch_nd_throughput_global.csv +awk -F, 'NR == 1 || $1 == "block"' launch_nd_throughput.csv > launch_nd_throughput_block.csv +awk -F, 'NR == 1 || $1 == "thread_local"' launch_nd_throughput.csv > launch_nd_throughput_thread_local.csv + +if [ "${PLOT_FORMAT}" = "png" ]; then +gnuplot -e " +set datafile separator ','; +set terminal pngcairo size 1400,900 enhanced font ',14'; +set output 'launch_nd_throughput.png'; +set object 1 rectangle from screen 0,0 to screen 1,1 fillcolor rgb '#ffffff' behind; +set border lw 2 lc rgb '#000000'; +set key outside right top; +set logscale x 2; +set logscale y 10; +set grid xtics ytics lc rgb '#bfbfbf' lw 1.5; +set tics nomirror; +set tics textcolor rgb '#000000'; +set xlabel 'Total logical iterations'; +set ylabel 'Throughput (iterations/s)'; +set format x '2^{%L}'; +set format y '10^{%L}'; +set style line 1 lt 1 lw 5 pt 7 ps 1.8 lc rgb '#0072B2'; +set style line 2 lt 1 lw 5 pt 5 ps 1.8 lc rgb '#009E73'; +set style line 3 lt 1 lw 5 pt 9 ps 1.8 lc rgb '#D55E00'; +set style line 4 lt 1 lw 5 pt 13 ps 1.8 lc rgb '#CC79A7'; +plot \ + 'launch_nd_throughput_flat.csv' using 4:6 with linespoints ls 1 title 'flat', \ + 'launch_nd_throughput_global.csv' using 4:6 with linespoints ls 2 title 'global', \ + 'launch_nd_throughput_block.csv' using 4:6 with linespoints ls 3 title 'block', \ + 'launch_nd_throughput_thread_local.csv' using 4:6 with linespoints ls 4 title 'thread_local' +" +elif [ "${PLOT_FORMAT}" = "svg" ]; then +gnuplot -e " +set datafile separator ','; +set terminal svg size 1400,900 dynamic; +set output 'launch_nd_throughput.svg'; +set object 1 rectangle from screen 0,0 to screen 1,1 fillcolor rgb '#ffffff' behind; +set border lw 2 lc rgb '#000000'; +set key outside right top; +set logscale x 2; +set logscale y 10; +set grid xtics ytics lc rgb '#bfbfbf' lw 1.5; +set tics nomirror; +set tics textcolor rgb '#000000'; +set xlabel 'Total logical iterations'; +set ylabel 'Throughput (iterations/s)'; +set format x '2^{%L}'; +set format y '10^{%L}'; +set style line 1 lt 1 lw 5 pt 7 ps 1.8 lc rgb '#0072B2'; +set style line 2 lt 1 lw 5 pt 5 ps 1.8 lc rgb '#009E73'; +set style line 3 lt 1 lw 5 pt 9 ps 1.8 lc rgb '#D55E00'; +set style line 4 lt 1 lw 5 pt 13 ps 1.8 lc rgb '#CC79A7'; +plot \ + 'launch_nd_throughput_flat.csv' using 4:6 with linespoints ls 1 title 'flat', \ + 'launch_nd_throughput_global.csv' using 4:6 with linespoints ls 2 title 'global', \ + 'launch_nd_throughput_block.csv' using 4:6 with linespoints ls 3 title 'block', \ + 'launch_nd_throughput_thread_local.csv' using 4:6 with linespoints ls 4 title 'thread_local' +" +elif [ "${PLOT_FORMAT}" != "none" ]; then + echo "Unsupported PLOT_FORMAT=${PLOT_FORMAT} (expected png, svg, or none)" >&2 + exit 1 +fi + +if [ "${TERMINAL_PLOT}" = "1" ]; then + gnuplot -e " +set datafile separator ','; +set terminal dumb size ${TERMINAL_PLOT_SIZE}; +set key outside; +set logscale x 2; +set logscale y 10; +set xlabel 'Total logical iterations'; +set ylabel 'Throughput (iterations/s)'; +plot \ + 'launch_nd_throughput_flat.csv' using 4:6 with linespoints title 'flat', \ + 'launch_nd_throughput_global.csv' using 4:6 with linespoints title 'global', \ + 'launch_nd_throughput_block.csv' using 4:6 with linespoints title 'block', \ + 'launch_nd_throughput_thread_local.csv' using 4:6 with linespoints title 'thread_local' +" +fi + +echo "Wrote $(pwd)/launch_nd.cali" +echo "Wrote $(pwd)/launch_nd_throughput.csv" +if [ "${PLOT_FORMAT}" = "png" ]; then + echo "Wrote $(pwd)/launch_nd_throughput.png" +elif [ "${PLOT_FORMAT}" = "svg" ]; then + echo "Wrote $(pwd)/launch_nd_throughput.svg" +fi diff --git a/include/RAJA/pattern/launch_nd.hpp b/include/RAJA/pattern/launch_nd.hpp index 0731cc60b3..7bfc367e8f 100644 --- a/include/RAJA/pattern/launch_nd.hpp +++ b/include/RAJA/pattern/launch_nd.hpp @@ -77,6 +77,25 @@ struct launch_nd_grid_policy namespace detail { +template +RAJA_INLINE auto make_launch_nd_context(std::string&& kernel_name) +{ + return util::make_context(std::move(kernel_name)); +} + +#if defined(RAJA_GPU_ACTIVE) +template +RAJA_INLINE auto make_launch_nd_context(RAJA::resources::Resource resource, + std::string&& kernel_name) +{ + return resource.get_platform() == RAJA::Platform::host + ? util::make_context( + std::move(kernel_name)) + : util::make_context( + std::move(kernel_name)); +} +#endif + template struct reverse_idx_seq; @@ -389,64 +408,143 @@ resources::EventProxy launch_nd_grid_execute( template + typename... Params> void launch_nd(launch_nd_flattened_policy, TypedRangeSegmentPack const& segs, - Lambda&& body) + Params&&... params) { + auto f_params = expt::make_forall_param_pack(std::forward(params)...); + std::string kernel_name = + expt::get_kernel_name(std::forward(params)...); + auto&& loop_body = expt::get_lambda(std::forward(params)...); + expt::check_forall_optional_args(loop_body, f_params); + + util::PluginContext context { + detail::make_launch_nd_context(std::move(kernel_name))}; + util::callPreCapturePlugins(context); + + using RAJA::util::trigger_updates_before; + auto body = trigger_updates_before(loop_body); + + util::callPostCapturePlugins(context); + util::callPreLaunchPlugins(context); + detail::launch_nd_flattened_execute( - segs, std::forward(body)); + segs, std::move(body)); + + util::callPostLaunchPlugins(context); } template + typename... Params> resources::EventProxy launch_nd( RAJA::resources::Resource resource, launch_nd_flattened_policy, TypedRangeSegmentPack const& segs, - Lambda&& body) + Params&&... params) { - return detail::launch_nd_flattened_execute( - resource, segs, std::forward(body)); + auto f_params = expt::make_forall_param_pack(std::forward(params)...); + std::string kernel_name = + expt::get_kernel_name(std::forward(params)...); + auto&& loop_body = expt::get_lambda(std::forward(params)...); + expt::check_forall_optional_args(loop_body, f_params); + + util::PluginContext context { + detail::make_launch_nd_context(std::move(kernel_name))}; + util::callPreCapturePlugins(context); + + using RAJA::util::trigger_updates_before; + auto body = trigger_updates_before(loop_body); + + util::callPostCapturePlugins(context); + util::callPreLaunchPlugins(context); + + auto event = detail::launch_nd_flattened_execute( + resource, segs, std::move(body)); + + util::callPostLaunchPlugins(context); + return event; } template + typename... Params> void launch_nd(launch_nd_grid_policy policy, TypedRangeSegmentPack const& segs, - Lambda&& body) + Params&&... params) { static_assert(sizeof...(IdxTs) == sizeof...(LoopPolicies), "RAJA::launch_nd launch backend requires one loop policy per " "segment"); + auto f_params = expt::make_forall_param_pack(std::forward(params)...); + std::string kernel_name = + expt::get_kernel_name(std::forward(params)...); + auto&& loop_body = expt::get_lambda(std::forward(params)...); + expt::check_forall_optional_args(loop_body, f_params); + + util::PluginContext context {detail::make_launch_nd_context< + typename LaunchPolicy::host_policy_t>(std::move(kernel_name))}; + util::callPreCapturePlugins(context); + + using RAJA::util::trigger_updates_before; + auto body = trigger_updates_before(loop_body); + + util::callPostCapturePlugins(context); + util::callPreLaunchPlugins(context); + detail::launch_nd_grid_execute>( - policy.launch_params, segs, std::forward(body), + policy.launch_params, segs, std::move(body), camp::num {}); + + util::callPostLaunchPlugins(context); } template + typename... Params> resources::EventProxy launch_nd( RAJA::resources::Resource resource, launch_nd_grid_policy policy, TypedRangeSegmentPack const& segs, - Lambda&& body) + Params&&... params) { static_assert(sizeof...(IdxTs) == sizeof...(LoopPolicies), "RAJA::launch_nd launch backend requires one loop policy per " "segment"); - return detail::launch_nd_grid_execute>( - resource, policy.launch_params, segs, std::forward(body), + auto f_params = expt::make_forall_param_pack(std::forward(params)...); + std::string kernel_name = + expt::get_kernel_name(std::forward(params)...); + auto&& loop_body = expt::get_lambda(std::forward(params)...); + expt::check_forall_optional_args(loop_body, f_params); + +#if defined(RAJA_GPU_ACTIVE) + util::PluginContext context {detail::make_launch_nd_context( + resource, std::move(kernel_name))}; +#else + util::PluginContext context {detail::make_launch_nd_context< + typename LaunchPolicy::host_policy_t>(std::move(kernel_name))}; +#endif + util::callPreCapturePlugins(context); + + using RAJA::util::trigger_updates_before; + auto body = trigger_updates_before(loop_body); + + util::callPostCapturePlugins(context); + util::callPreLaunchPlugins(context); + + auto event = detail::launch_nd_grid_execute>( + resource, policy.launch_params, segs, std::move(body), camp::num {}); + + util::callPostLaunchPlugins(context); + return event; } } // namespace RAJA diff --git a/scripts/lc-builds/toss4_amdclang.sh b/scripts/lc-builds/toss4_amdclang.sh index f61a8fe307..0bff17c4b2 100755 --- a/scripts/lc-builds/toss4_amdclang.sh +++ b/scripts/lc-builds/toss4_amdclang.sh @@ -11,6 +11,7 @@ # Default CMake version if not provided DEFAULT_CMAKE_VER=3.24.2 +DEFAULT_CALIPER_DIR="$(pwd)/../install-caliper/share/cmake/caliper" if [[ $# -lt 2 ]]; then echo @@ -22,6 +23,7 @@ if [[ $# -lt 2 ]]; then echo "For example: " echo " toss4_amdclang.sh 4.1.0 gfx906 [3.27.4]" echo "If no CMake version is provided, version ${DEFAULT_CMAKE_VER} will be used." + echo "Set CALIPER_DIR to override the default Caliper package location." exit 1 fi @@ -52,10 +54,26 @@ else fi BUILD_SUFFIX=lc_toss4-amdclang-${COMP_VER}-${COMP_ARCH} +CRAY_RPATH="${CRAY_LD_LIBRARY_PATH//:/;}" + +CALIPER_DIR_TO_USE="${CALIPER_DIR:-${DEFAULT_CALIPER_DIR}}" +if [[ ! -d "${CALIPER_DIR_TO_USE}" ]]; then + echo "Caliper is required for this build." + echo "Set CALIPER_DIR to a valid Caliper CMake package directory." + echo "Expected to find: ${CALIPER_DIR_TO_USE}" + exit 1 +fi + +CALIPER_ARGS=( + -DRAJA_ENABLE_RUNTIME_PLUGINS=ON + -DRAJA_ENABLE_CALIPER=ON + -Dcaliper_DIR="${CALIPER_DIR_TO_USE}" +) echo echo "Creating build directory build_${BUILD_SUFFIX} and generating configuration in it" echo "Using CMake version: ${CMAKE_VER}" +echo "Caliper support: required (${CALIPER_DIR_TO_USE})" echo "Configuration extra arguments:" echo " $@" echo @@ -86,6 +104,8 @@ cmake \ -DCMAKE_HIP_ARCHITECTURES="${COMP_ARCH}" \ -DGPU_TARGETS="${COMP_ARCH}" \ -DAMDGPU_TARGETS="${COMP_ARCH}" \ + -DCMAKE_BUILD_RPATH="${CRAY_RPATH}" \ + -DCMAKE_INSTALL_RPATH="${CRAY_RPATH}" \ -DBLT_CXX_STD=c++20 \ -C "../host-configs/lc-builds/toss4/${HOSTCONFIG}.cmake" \ -DENABLE_HIP=ON \ @@ -93,6 +113,7 @@ cmake \ -DENABLE_CUDA=OFF \ -DENABLE_BENCHMARKS=On \ -DCMAKE_INSTALL_PREFIX=../install_${BUILD_SUFFIX} \ + "${CALIPER_ARGS[@]}" \ "$@" \ .. From 6349a4ef4193ad8d43fd308a2186e785d5c666f8 Mon Sep 17 00:00:00 2001 From: Arturo Vargas Date: Tue, 9 Jun 2026 13:09:55 -0700 Subject: [PATCH 4/5] minor changes for application usage --- examples/launch_nd.cpp | 143 +++++++++++++++++---- examples/launch_nd_benchmark_workflow.md | 4 + include/RAJA/pattern/launch_nd.hpp | 154 +++++++++++++++++++++++ 3 files changed, 275 insertions(+), 26 deletions(-) diff --git a/examples/launch_nd.cpp b/examples/launch_nd.cpp index 832eb4c93d..a4e7aa1e56 100644 --- a/examples/launch_nd.cpp +++ b/examples/launch_nd.cpp @@ -22,6 +22,10 @@ * multi-size studies, and uses RAJA::Name so Caliper can aggregate results by * policy/size combination. * + * It also demonstrates runtime backend selection using RAJA::ExecPlace and the + * extended RAJA::launch_nd overloads that accept an ExecPlace plus a + * host/device policy pair. + * * Typical runs: * * ./launch_nd --mapping all --repetitions 50 --warmup 5 @@ -60,6 +64,12 @@ struct ProblemSize struct Options { Mapping mapping = Mapping::All; + RAJA::ExecPlace exec_place = +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) + RAJA::ExecPlace::DEVICE; +#else + RAJA::ExecPlace::HOST; +#endif int warmup = 5; int repetitions = 50; std::vector sizes = {}; @@ -128,6 +138,18 @@ ProblemSize default_problem_size() return ProblemSize {262144, 8}; } +const char* exec_place_name(RAJA::ExecPlace place) +{ + switch (place) + { + case RAJA::ExecPlace::HOST: + return "host"; + case RAJA::ExecPlace::DEVICE: + return "device"; + } + return "unknown"; +} + void print_usage(const char* executable) { std::cout @@ -136,6 +158,9 @@ void print_usage(const char* executable) << "Options:\n" << " --mapping \n" << " Select policy set to benchmark.\n" +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) + << " --exec-place Select execution backend at runtime.\n" +#endif << " --sizes Problem sizes, e.g. 262144x8 or" << " 65536x8,262144x8.\n" << " --warmup Warmup launches per mapping/size.\n" @@ -172,6 +197,21 @@ Mapping parse_mapping(const std::string& value) " or all)"); } +RAJA::ExecPlace parse_exec_place(const std::string& value) +{ + if (value == "host" || value == "cpu") + { + return RAJA::ExecPlace::HOST; + } + if (value == "device" || value == "gpu") + { + return RAJA::ExecPlace::DEVICE; + } + + throw std::runtime_error("unknown exec-place '" + value + + "' (expected host or device)"); +} + int parse_positive_int(const std::string& name, const std::string& value) { try @@ -267,6 +307,13 @@ Options parse_options(int argc, char** argv) { options.mapping = parse_mapping(require_value(i, argc, argv, "--mapping")); } +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) + else if (arg == "--exec-place") + { + options.exec_place = + parse_exec_place(require_value(i, argc, argv, "--exec-place")); + } +#endif else if (arg == "--sizes") { options.sizes = parse_problem_sizes(require_value(i, argc, argv, "--sizes")); @@ -329,7 +376,8 @@ std::string kernel_name(Mapping mapping, const ProblemSize& size) template void launch_kernel(RAJA::resources::Resource res, - LaunchNdPolicy policy, + RAJA::ExecPlace exec_place, + LaunchNdPolicy const& policy, int* values_ptr, const ProblemSize& size, const std::string& name) @@ -337,7 +385,7 @@ void launch_kernel(RAJA::resources::Resource res, auto cell_segment = RAJA::TypedRangeSegment(0, size.cells); auto comp_segment = RAJA::TypedRangeSegment(0, size.components); - RAJA::launch_nd(res, policy, RAJA::segments(cell_segment, comp_segment), + RAJA::launch_nd(res, exec_place, policy, RAJA::segments(cell_segment, comp_segment), RAJA::Name(name.c_str()), [=] RAJA_HOST_DEVICE(int cell, int comp) { values_ptr[comp + size.components * cell] = @@ -376,9 +424,11 @@ int verify_result(RAJA::resources::Resource res, return errors; } -template +template RunResult benchmark_mapping(RAJA::resources::Resource res, - LaunchNdPolicy policy, + RAJA::ExecPlace exec_place, + HostPolicy host_policy, + DevicePolicy device_policy, const Options& options, Mapping mapping, const ProblemSize& size) @@ -388,14 +438,18 @@ RunResult benchmark_mapping(RAJA::resources::Resource res, const std::string name = kernel_name(mapping, size); int* values_ptr = res.allocate(total); + const auto policy = + RAJA::make_launch_nd_place_policy(std::move(host_policy), + std::move(device_policy)); + for (int step = 0; step < options.warmup; ++step) { - launch_kernel(res, policy, values_ptr, size, name); + launch_kernel(res, exec_place, policy, values_ptr, size, name); } for (int rep = 0; rep < options.repetitions; ++rep) { - launch_kernel(res, policy, values_ptr, size, name); + launch_kernel(res, exec_place, policy, values_ptr, size, name); } RunResult result; @@ -471,9 +525,23 @@ int main(int argc, char** argv) std::cout << '\n'; RAJA::resources::Resource res(resource_type {}); + const RAJA::ExecPlace exec_place = options.exec_place; + +#if !defined(RAJA_ENABLE_CUDA) && !defined(RAJA_ENABLE_HIP) + if (exec_place == RAJA::ExecPlace::DEVICE) + { + throw std::runtime_error( + "--exec-place device requested but no GPU backend is enabled"); + } +#endif + + std::cout << " exec place: " << exec_place_name(exec_place) << '\n'; int total_errors = 0; + using host_launch_policy = RAJA::LaunchPolicy; + using host_loop_policy = RAJA::LoopPolicy; + for (const ProblemSize& size : sizes) { std::cout << "\nStudy size " << size_label(size) << '\n'; @@ -482,48 +550,71 @@ int main(int argc, char** argv) { if (mapping == Mapping::Flat) { - auto policy = RAJA::launch_nd_flattened_policy {}; const RunResult result = - benchmark_mapping(res, policy, options, mapping, size); + benchmark_mapping(res, + exec_place, + RAJA::launch_nd_flattened_policy {}, + RAJA::launch_nd_flattened_policy {}, + options, + mapping, + size); total_errors += result.errors; std::cout << " flattened launch threads per block: " << block_size_1d << "\n\n"; } else if (mapping == Mapping::Global) { - auto policy = - RAJA::launch_nd_grid_policy( - make_global_launch_params(size)); const RunResult result = - benchmark_mapping(res, policy, options, mapping, size); + benchmark_mapping( + res, + exec_place, + RAJA::launch_nd_grid_policy(RAJA::LaunchParams {}), + RAJA::launch_nd_grid_policy(make_global_launch_params(size)), + options, + mapping, + size); total_errors += result.errors; std::cout << " global launch block shape: " << block_x << " x " << block_y << "\n\n"; } else if (mapping == Mapping::Block) { - auto policy = - RAJA::launch_nd_grid_policy( - make_block_launch_params(size)); const RunResult result = - benchmark_mapping(res, policy, options, mapping, size); + benchmark_mapping( + res, + exec_place, + RAJA::launch_nd_grid_policy(RAJA::LaunchParams {}), + RAJA::launch_nd_grid_policy(make_block_launch_params(size)), + options, + mapping, + size); total_errors += result.errors; std::cout << " block launch uses one logical iteration per team" << "\n\n"; } else { - auto policy = - RAJA::launch_nd_grid_policy( - make_thread_launch_params(size)); const RunResult result = - benchmark_mapping(res, policy, options, mapping, size); + benchmark_mapping( + res, + exec_place, + RAJA::launch_nd_grid_policy(RAJA::LaunchParams {}), + RAJA::launch_nd_grid_policy(make_thread_launch_params(size)), + options, + mapping, + size); total_errors += result.errors; std::cout << " thread-local launch uses a single team with thread" << " loops of shape " << block_x << " x " << block_y diff --git a/examples/launch_nd_benchmark_workflow.md b/examples/launch_nd_benchmark_workflow.md index db26f69e55..8148bef206 100644 --- a/examples/launch_nd_benchmark_workflow.md +++ b/examples/launch_nd_benchmark_workflow.md @@ -51,6 +51,7 @@ CALIPER_DIR="${CALIPER_DIR}" ./scripts/lc-builds/toss4_amdclang.sh 6.4.3 gfx90a cd "${REPO_DIR}/build-raja-clang" RAJA_CALIPER=1 CALI_CONFIG=runtime-report ./bin/launch_nd \ --mapping all \ + --exec-place device \ --sizes 65536x8,262144x8,1048576x8 \ --warmup 5 \ --repetitions 50 @@ -73,6 +74,7 @@ cd "${REPO_DIR}/build-raja-clang" RAJA_CALIPER=1 CALI_CONFIG=runtime-profile(output=launch_nd.cali,output.format=cali) \ ./bin/launch_nd \ --mapping all \ + --exec-place device \ --sizes 65536x8,262144x8,1048576x8 \ --warmup 5 \ --repetitions 50 @@ -146,3 +148,5 @@ plot \ - Save both the Caliper profile and the derived throughput CSV for each build. - Compare `flat`, `global`, `block`, and `thread_local` at each size, then examine how the gap changes with total problem size. +- When building with CUDA/HIP enabled, use `--exec-place host` to run the + same benchmark on the host backend for an apples-to-apples comparison. diff --git a/include/RAJA/pattern/launch_nd.hpp b/include/RAJA/pattern/launch_nd.hpp index 7bfc367e8f..512587f866 100644 --- a/include/RAJA/pattern/launch_nd.hpp +++ b/include/RAJA/pattern/launch_nd.hpp @@ -59,6 +59,44 @@ struct launch_nd_flattened_policy using layout_tag = LayoutTag; }; +/*! + * A flattened launch_nd policy that supports selecting between a host and device + * exec policy at runtime via RAJA::ExecPlace. + * + * This is useful when the caller wants a single launch site but needs to choose + * between e.g. seq_exec and cuda_exec/hip_exec based on runtime conditions. + */ +template +struct launch_nd_flattened_place_policy +{ + using host_exec_policy = HostExecPolicy; + using device_exec_policy = DeviceExecPolicy; + using layout_tag = LayoutTag; +}; + +/*! + * General runtime host/device launch_nd policy container. + * + * This allows callers to select different launch_nd policy *kinds* for host and + * device (e.g., a grid policy on host for nested loops and a flattened policy + * on device for a 1D mapping). + */ +template +struct launch_nd_place_policy +{ + HostPolicy host; + DevicePolicy device; +}; + +template +RAJA_INLINE launch_nd_place_policy +make_launch_nd_place_policy(HostPolicy host, DevicePolicy device) +{ + return {std::move(host), std::move(device)}; +} + template struct launch_nd_grid_policy { @@ -547,6 +585,122 @@ resources::EventProxy launch_nd( return event; } +template +void launch_nd(ExecPlace place, + launch_nd_flattened_place_policy, + TypedRangeSegmentPack const& segs, + Params&&... params) +{ + switch (place) + { + case ExecPlace::HOST: + RAJA::launch_nd(launch_nd_flattened_policy{}, + segs, + std::forward(params)...); + break; +#if defined(RAJA_GPU_ACTIVE) + case ExecPlace::DEVICE: + RAJA::launch_nd(launch_nd_flattened_policy{}, + segs, + std::forward(params)...); + break; +#endif + default: + RAJA_ABORT_OR_THROW("Unknown launch place or device is not enabled"); + } +} + +template +resources::EventProxy launch_nd( + RAJA::resources::Resource resource, + ExecPlace place, + launch_nd_flattened_place_policy, + TypedRangeSegmentPack const& segs, + Params&&... params) +{ + switch (place) + { + case ExecPlace::HOST: + return RAJA::launch_nd(resource, + launch_nd_flattened_policy{}, + segs, + std::forward(params)...); +#if defined(RAJA_GPU_ACTIVE) + case ExecPlace::DEVICE: + return RAJA::launch_nd(resource, + launch_nd_flattened_policy{}, + segs, + std::forward(params)...); +#endif + default: + RAJA_ABORT_OR_THROW("Unknown launch place or device is not enabled"); + } + return resources::EventProxy(resource); +} + +template +void launch_nd(ExecPlace place, + launch_nd_place_policy const& policy, + TypedRangeSegmentPack const& segs, + Params&&... params) +{ + switch (place) + { + case ExecPlace::HOST: + RAJA::launch_nd(policy.host, segs, std::forward(params)...); + break; +#if defined(RAJA_GPU_ACTIVE) + case ExecPlace::DEVICE: + RAJA::launch_nd(policy.device, segs, std::forward(params)...); + break; +#endif + default: + RAJA_ABORT_OR_THROW("Unknown launch place or device is not enabled"); + } +} + +template +resources::EventProxy launch_nd( + RAJA::resources::Resource resource, + ExecPlace place, + launch_nd_place_policy const& policy, + TypedRangeSegmentPack const& segs, + Params&&... params) +{ + switch (place) + { + case ExecPlace::HOST: + return RAJA::launch_nd(resource, + policy.host, + segs, + std::forward(params)...); +#if defined(RAJA_GPU_ACTIVE) + case ExecPlace::DEVICE: + return RAJA::launch_nd(resource, + policy.device, + segs, + std::forward(params)...); +#endif + default: + RAJA_ABORT_OR_THROW("Unknown launch place or device is not enabled"); + } + return resources::EventProxy(resource); +} + } // namespace RAJA #endif /* RAJA_pattern_launch_nd_HPP */ From 9329f4f1e29a80033d53253c9ea29a03e038acf3 Mon Sep 17 00:00:00 2001 From: Rich Hornung Date: Wed, 8 Jul 2026 08:35:07 -0700 Subject: [PATCH 5/5] clang format --- include/RAJA/pattern/launch_nd.hpp | 67 +++++++++++++++--------------- 1 file changed, 33 insertions(+), 34 deletions(-) diff --git a/include/RAJA/pattern/launch_nd.hpp b/include/RAJA/pattern/launch_nd.hpp index 512587f866..9d05626a26 100644 --- a/include/RAJA/pattern/launch_nd.hpp +++ b/include/RAJA/pattern/launch_nd.hpp @@ -60,8 +60,8 @@ struct launch_nd_flattened_policy }; /*! - * A flattened launch_nd policy that supports selecting between a host and device - * exec policy at runtime via RAJA::ExecPlace. + * A flattened launch_nd policy that supports selecting between a host and + * device exec policy at runtime via RAJA::ExecPlace. * * This is useful when the caller wants a single launch site but needs to choose * between e.g. seq_exec and cuda_exec/hip_exec based on runtime conditions. @@ -467,8 +467,8 @@ void launch_nd(launch_nd_flattened_policy, util::callPostCapturePlugins(context); util::callPreLaunchPlugins(context); - detail::launch_nd_flattened_execute( - segs, std::move(body)); + detail::launch_nd_flattened_execute(segs, + std::move(body)); util::callPostLaunchPlugins(context); } @@ -524,8 +524,9 @@ void launch_nd(launch_nd_grid_policy policy, auto&& loop_body = expt::get_lambda(std::forward(params)...); expt::check_forall_optional_args(loop_body, f_params); - util::PluginContext context {detail::make_launch_nd_context< - typename LaunchPolicy::host_policy_t>(std::move(kernel_name))}; + util::PluginContext context { + detail::make_launch_nd_context( + std::move(kernel_name))}; util::callPreCapturePlugins(context); using RAJA::util::trigger_updates_before; @@ -565,8 +566,9 @@ resources::EventProxy launch_nd( util::PluginContext context {detail::make_launch_nd_context( resource, std::move(kernel_name))}; #else - util::PluginContext context {detail::make_launch_nd_context< - typename LaunchPolicy::host_policy_t>(std::move(kernel_name))}; + util::PluginContext context { + detail::make_launch_nd_context( + std::move(kernel_name))}; #endif util::callPreCapturePlugins(context); @@ -576,10 +578,10 @@ resources::EventProxy launch_nd( util::callPostCapturePlugins(context); util::callPreLaunchPlugins(context); - auto event = detail::launch_nd_grid_execute>( - resource, policy.launch_params, segs, std::move(body), - camp::num {}); + auto event = + detail::launch_nd_grid_execute>( + resource, policy.launch_params, segs, std::move(body), + camp::num {}); util::callPostLaunchPlugins(context); return event; @@ -591,22 +593,23 @@ template void launch_nd(ExecPlace place, - launch_nd_flattened_place_policy, + launch_nd_flattened_place_policy, TypedRangeSegmentPack const& segs, Params&&... params) { switch (place) { case ExecPlace::HOST: - RAJA::launch_nd(launch_nd_flattened_policy{}, - segs, - std::forward(params)...); + RAJA::launch_nd(launch_nd_flattened_policy {}, + segs, std::forward(params)...); break; #if defined(RAJA_GPU_ACTIVE) case ExecPlace::DEVICE: - RAJA::launch_nd(launch_nd_flattened_policy{}, - segs, - std::forward(params)...); + RAJA::launch_nd( + launch_nd_flattened_policy {}, segs, + std::forward(params)...); break; #endif default: @@ -622,23 +625,23 @@ template launch_nd( RAJA::resources::Resource resource, ExecPlace place, - launch_nd_flattened_place_policy, + launch_nd_flattened_place_policy, TypedRangeSegmentPack const& segs, Params&&... params) { switch (place) { case ExecPlace::HOST: - return RAJA::launch_nd(resource, - launch_nd_flattened_policy{}, - segs, - std::forward(params)...); + return RAJA::launch_nd( + resource, launch_nd_flattened_policy {}, + segs, std::forward(params)...); #if defined(RAJA_GPU_ACTIVE) case ExecPlace::DEVICE: - return RAJA::launch_nd(resource, - launch_nd_flattened_policy{}, - segs, - std::forward(params)...); + return RAJA::launch_nd( + resource, launch_nd_flattened_policy {}, + segs, std::forward(params)...); #endif default: RAJA_ABORT_OR_THROW("Unknown launch place or device is not enabled"); @@ -684,15 +687,11 @@ resources::EventProxy launch_nd( switch (place) { case ExecPlace::HOST: - return RAJA::launch_nd(resource, - policy.host, - segs, + return RAJA::launch_nd(resource, policy.host, segs, std::forward(params)...); #if defined(RAJA_GPU_ACTIVE) case ExecPlace::DEVICE: - return RAJA::launch_nd(resource, - policy.device, - segs, + return RAJA::launch_nd(resource, policy.device, segs, std::forward(params)...); #endif default: