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..c790bb1158 --- /dev/null +++ b/docs/sphinx/user_guide/cook_book/launch-nd.rst @@ -0,0 +1,120 @@ +.. ## +.. ## 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::nd_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``. + +When using overloads that take both a ``RAJA::resources::Resource`` and an +explicit ``RAJA::ExecPlace``, the place must match the resource platform (host +vs device); otherwise RAJA aborts or throws. diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 84f07735f4..5035751e6a 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -77,6 +77,10 @@ raja_add_executable( NAME launch_flatten SOURCES launch_flatten.cpp) +raja_add_executable( + NAME launch_nd + SOURCES launch_nd.cpp) + raja_add_executable( NAME launch_reductions SOURCES launch_reductions.cpp) diff --git a/examples/launch_nd.cpp b/examples/launch_nd.cpp new file mode 100644 index 0000000000..3dba2a50e0 --- /dev/null +++ b/examples/launch_nd.cpp @@ -0,0 +1,634 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// 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 + +#include "RAJA/RAJA.hpp" + +/* + * RAJA::launch_nd benchmark driver + * + * 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. + * + * 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 + * 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 +{ + +constexpr int block_size_1d = 256; +constexpr int block_x = 16; +constexpr int block_y = 16; + +#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP) +int ceil_div(int value, int divisor) { return (value + divisor - 1) / divisor; } +#endif + +enum class Mapping +{ + Flat, + Global, + Block, + ThreadLocal, + All +}; + +struct ProblemSize +{ + int cells = 0; + int components = 0; +}; + +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 = {}; +}; + +struct RunResult +{ + int errors = 0; +}; + +#if defined(RAJA_ENABLE_CUDA) +using launch_policy = RAJA::LaunchPolicy>; +using flat_exec = RAJA::cuda_exec; +using global_cell_grid_mapping = RAJA::LoopPolicy; +using global_comp_grid_mapping = RAJA::LoopPolicy; +using block_cell_grid_mapping = RAJA::LoopPolicy; +using block_comp_grid_mapping = RAJA::LoopPolicy; +using thread_cell_in_kernel_loops = RAJA::LoopPolicy; +using thread_comp_in_kernel_loops = 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 global_cell_grid_mapping = RAJA::LoopPolicy; +using global_comp_grid_mapping = RAJA::LoopPolicy; +using block_cell_grid_mapping = RAJA::LoopPolicy; +using block_comp_grid_mapping = RAJA::LoopPolicy; +using thread_cell_in_kernel_loops = RAJA::LoopPolicy; +using thread_comp_in_kernel_loops = RAJA::LoopPolicy; +using resource_type = RAJA::resources::Hip; + +#else +using launch_policy = RAJA::LaunchPolicy; +using flat_exec = RAJA::seq_exec; +using global_cell_grid_mapping = RAJA::LoopPolicy; +using global_comp_grid_mapping = RAJA::LoopPolicy; +using block_cell_grid_mapping = RAJA::LoopPolicy; +using block_comp_grid_mapping = RAJA::LoopPolicy; +using thread_cell_in_kernel_loops = RAJA::LoopPolicy; +using thread_comp_in_kernel_loops = RAJA::LoopPolicy; +using resource_type = RAJA::resources::Host; +#endif + +const char* mapping_name(Mapping mapping) +{ + switch (mapping) + { + 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"; + } + + return "unknown"; +} + +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 + << "Usage: " << executable + << " [flat|global|block|thread_local|all] [options]\n" + << "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" + << " --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 (value == "global" || value == "grid") + { + return Mapping::Global; + } + if (value == "block") + { + return Mapping::Block; + } + if (value == "thread" || value == "thread_local" || + value == "thread-local") + { + return Mapping::ThreadLocal; + } + if (value == "all") + { + return Mapping::All; + } + + throw std::runtime_error("unknown mapping '" + value + + "' (expected flat, global, block, thread_local," + " 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 + { + 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")); + } +#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")); + } + 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}; +} + +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); +} + +template +void launch_kernel(RAJA::resources::Resource res, + RAJA::ExecPlace exec_place, + LaunchNdPolicy const& 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, exec_place, policy, RAJA::nd_segments(cell_segment, comp_segment), + RAJA::Name(name.c_str()), + [=] RAJA_HOST_DEVICE(int cell, int comp) { + values_ptr[comp + size.components * cell] = + 1000 * cell + comp; + }); + + // 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 < size.cells; ++cell) + { + for (int comp = 0; comp < size.components; ++comp) + { + const int idx = comp + size.components * cell; + const int expected = 1000 * cell + comp; + if (values[idx] != expected) + { + ++errors; + } + } + } + + return errors; +} + +template +RunResult benchmark_mapping(RAJA::resources::Resource res, + RAJA::ExecPlace exec_place, + HostPolicy host_policy, + DevicePolicy device_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); + + 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, exec_place, policy, values_ptr, size, name); + } + + for (int rep = 0; rep < options.repetitions; ++rep) + { + launch_kernel(res, exec_place, 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 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) +{ + try + { + const Options options = parse_options(argc, argv); + const std::vector sizes = selected_sizes(options); + + 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'; + + 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'; + + for (Mapping mapping : selected_mappings(options.mapping)) + { + if (mapping == Mapping::Flat) + { + const RunResult result = + 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) + { + const RunResult result = + 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) + { + const RunResult result = + 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 + { + const RunResult result = + 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 + << "\n\n"; + } + } + } + + std::cout << "DONE!...\n"; + return total_errors == 0 ? 0 : 1; + } + catch (const std::exception& ex) + { + std::cerr << "Error: " << ex.what() << '\n'; + return 1; + } +} diff --git a/examples/launch_nd_benchmark_workflow.md b/examples/launch_nd_benchmark_workflow.md new file mode 100644 index 0000000000..8148bef206 --- /dev/null +++ b/examples/launch_nd_benchmark_workflow.md @@ -0,0 +1,152 @@ +# `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 \ + --exec-place device \ + --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 \ + --exec-place device \ + --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. +- 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/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/RAJA.hpp b/include/RAJA/RAJA.hpp index 911bec021d..a086b46ebf 100644 --- a/include/RAJA/RAJA.hpp +++ b/include/RAJA/RAJA.hpp @@ -126,6 +126,7 @@ #include "RAJA/util/StaticLayout.hpp" #include "RAJA/util/IndexLayout.hpp" #include "RAJA/util/View.hpp" +#include "RAJA/pattern/launch_nd.hpp" // diff --git a/include/RAJA/pattern/launch_nd.hpp b/include/RAJA/pattern/launch_nd.hpp new file mode 100644 index 0000000000..fdce5e2b2b --- /dev/null +++ b/include/RAJA/pattern/launch_nd.hpp @@ -0,0 +1,656 @@ +/*! + ****************************************************************************** + * + * \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 + +#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 +{ + +namespace detail +{ + +template +struct TypedRangeSegmentPack +{ + camp::tuple...> data; +}; + +template +struct is_typed_range_segment_pack : std::false_type +{}; + +template +struct is_typed_range_segment_pack> + : std::true_type +{}; + +template +inline constexpr bool is_typed_range_segment_pack_v = + is_typed_range_segment_pack::value; + +template +struct typed_range_segment_pack_rank; + +template +struct typed_range_segment_pack_rank> + : std::integral_constant +{}; + +template +inline constexpr camp::idx_t typed_range_segment_pack_rank_v = + typed_range_segment_pack_rank::value; + +template +concept typed_range_segment_pack = + is_typed_range_segment_pack_v>; + +} // namespace detail + +/*! + * Create a segment pack for use with ``RAJA::launch_nd``. + * + * Currently supports packs of ``RAJA::TypedRangeSegment`` only. + */ +template +RAJA_INLINE auto nd_segments(RAJA::TypedRangeSegment const&... segs) +{ + return detail::TypedRangeSegmentPack {camp::make_tuple(segs...)}; +} + +template +struct launch_nd_flattened_policy +{ + using exec_policy = ExecPolicy; + 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; +}; + +/*! + * Convenience alias for selecting between host/device exec policies for the + * flattened launch_nd mapping at runtime via RAJA::ExecPlace. + */ +template +using launch_nd_flattened_place_policy = launch_nd_place_policy< + launch_nd_flattened_policy, + launch_nd_flattened_policy>; + +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 +{ + 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 +RAJA_INLINE auto make_launch_nd_context(std::string&& kernel_name) +{ + return util::make_context(std::move(kernel_name)); +} + +#if defined(RAJA_GPU_ACTIVE) +RAJA_INLINE void validate_launch_nd_place_matches_resource( + RAJA::resources::Resource const& resource, + ExecPlace place) +{ + const bool resource_is_host = + (resource.get_platform() == RAJA::Platform::host); + + if (place == ExecPlace::HOST && !resource_is_host) + { + RAJA_ABORT_OR_THROW( + "RAJA::launch_nd: ExecPlace::HOST requires a host resource"); + } + else if (place == ExecPlace::DEVICE && resource_is_host) + { + RAJA_ABORT_OR_THROW( + "RAJA::launch_nd: ExecPlace::DEVICE requires a device resource"); + } +} +#else +RAJA_INLINE void validate_launch_nd_place_matches_resource( + RAJA::resources::Resource const& resource, + ExecPlace place) +{ + RAJA_UNUSED_VAR(resource); + if (place == ExecPlace::DEVICE) + { + RAJA_ABORT_OR_THROW("RAJA::launch_nd: ExecPlace::DEVICE requested but " + "device is not enabled"); + } +} +#endif + +#if defined(RAJA_GPU_ACTIVE) +template +RAJA_INLINE auto make_launch_nd_context(RAJA::resources::Resource resource, + std::string&& kernel_name) +{ + if (resource.get_platform() == RAJA::Platform::host) + { + return util::make_context( + std::move(kernel_name)); + } + return util::make_context( + std::move(kernel_name)); +} +#endif + +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, SegmentPack const& segs) +{ + constexpr camp::idx_t Rank = + typed_range_segment_pack_rank_v>; + + 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 +RAJA_INLINE auto launch_nd_with_plugins(util::PluginContext context, + Execute&& execute, + 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); + + context.kernel_name = 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); + + using result_type = decltype(std::forward(execute)(std::move(body))); + if constexpr (std::is_void::value) + { + std::forward(execute)(std::move(body)); + util::callPostLaunchPlugins(context); + } + else + { + auto event = std::forward(execute)(std::move(body)); + util::callPostLaunchPlugins(context); + return event; + } +} + +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 +struct launch_nd_flattened_body +{ + Adapter adapter; + + RAJA_HOST_DEVICE RAJA_INLINE void operator()(RAJA::LaunchContext ctx) const + { + RAJA::loop(ctx, adapter.getRange(), adapter); + } +}; + +template +void launch_nd_flattened_execute(SegmentPack 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()), + launch_nd_flattened_body> { + std::move(adapter)}); +} + +template +resources::EventProxy launch_nd_flattened_execute( + RAJA::resources::Resource resource, + SegmentPack 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()), + launch_nd_flattened_body> { + std::move(adapter)}); +} + +template +struct launch_nd_grid_body +{ + SegmentPack segs; + Body body; + + RAJA_HOST_DEVICE RAJA_INLINE void operator()(RAJA::LaunchContext ctx) const + { + exec_dim<0>(ctx); + } + +private: + static constexpr camp::idx_t rank = + typed_range_segment_pack_rank_v>; + + template + RAJA_HOST_DEVICE RAJA_INLINE void exec_dim(RAJA::LaunchContext ctx, + IdxTs... indices) const + { + if constexpr (Dim == rank) + { + body(indices...); + } + else + { + using loop_policy = + typename camp::at>::type; + auto const seg = camp::get(segs.data); + + RAJA::loop(ctx, seg, [&](auto idx) { + exec_dim(ctx, indices..., idx); + }); + } + } +}; + +template +void launch_nd_grid_execute(LaunchParams const& launch_params, + SegmentPack const& segs, + Lambda&& body) +{ + RAJA::launch( + launch_params, + launch_nd_grid_body, SegmentPack> { + segs, std::forward(body)}); +} + +template +resources::EventProxy launch_nd_grid_execute( + RAJA::resources::Resource resource, + LaunchParams const& launch_params, + SegmentPack const& segs, + Lambda&& body) +{ + return RAJA::launch( + resource, launch_params, + launch_nd_grid_body, SegmentPack> { + segs, std::forward(body)}); +} + +} // namespace detail + +template +void launch_nd(launch_nd_flattened_policy, + SegmentPack const& segs, + Params&&... params) +{ + detail::launch_nd_with_plugins( + detail::make_launch_nd_context(std::string {}), + [&](auto&& body) { + detail::launch_nd_flattened_execute( + segs, std::forward(body)); + }, + std::forward(params)...); +} + +template +resources::EventProxy launch_nd( + RAJA::resources::Resource resource, + launch_nd_flattened_policy, + SegmentPack const& segs, + Params&&... params) +{ + return detail::launch_nd_with_plugins( + detail::make_launch_nd_context(std::string {}), + [&](auto&& body) { + return detail::launch_nd_flattened_execute( + resource, segs, std::forward(body)); + }, + std::forward(params)...); +} + +template +void launch_nd(launch_nd_grid_policy policy, + SegmentPack const& segs, + Params&&... params) +{ + static_assert( + detail::typed_range_segment_pack_rank_v> == + static_cast(sizeof...(LoopPolicies)), + "RAJA::launch_nd launch backend requires one loop policy per " + "segment"); + + detail::launch_nd_with_plugins( + detail::make_launch_nd_context( + std::string {}), + [&](auto&& body) { + detail::launch_nd_grid_execute>( + policy.launch_params, segs, std::forward(body)); + }, + std::forward(params)...); +} + +template +resources::EventProxy launch_nd( + RAJA::resources::Resource resource, + launch_nd_grid_policy policy, + SegmentPack const& segs, + Params&&... params) +{ + static_assert( + detail::typed_range_segment_pack_rank_v> == + static_cast(sizeof...(LoopPolicies)), + "RAJA::launch_nd launch backend requires one loop policy per " + "segment"); + +#if defined(RAJA_GPU_ACTIVE) + util::PluginContext context { + detail::make_launch_nd_context(resource, std::string {})}; +#else + util::PluginContext context { + detail::make_launch_nd_context( + std::string {})}; +#endif + + return detail::launch_nd_with_plugins( + std::move(context), + [&](auto&& body) { + return detail::launch_nd_grid_execute>( + resource, policy.launch_params, segs, + std::forward(body)); + }, + std::forward(params)...); +} + +template +void launch_nd(ExecPlace place, + launch_nd_place_policy const& policy, + SegmentPack 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, + SegmentPack const& segs, + Params&&... params) +{ + detail::validate_launch_nd_place_matches_resource(resource, place); + 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 */ diff --git a/test/functional/util/CMakeLists.txt b/test/functional/util/CMakeLists.txt index 10dfcba0c9..d7c9d5cc84 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-launch_nd + SOURCES test-launch_nd.cpp) diff --git a/test/functional/util/test-launch_nd.cpp b/test/functional/util/test-launch_nd.cpp new file mode 100644 index 0000000000..57bbe60a25 --- /dev/null +++ b/test/functional/util/test-launch_nd.cpp @@ -0,0 +1,166 @@ +//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// +// 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; +using depth_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::nd_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::nd_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::nd_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, grid_3d_resource) +{ + constexpr int n = 3; + constexpr int batch_size = 4; + constexpr int depth = 2; + + std::vector values(n * batch_size * depth, -1); + auto rows = RAJA::TypedRangeSegment(0, n); + auto batches = RAJA::TypedRangeSegment(0, batch_size); + auto depths = RAJA::TypedRangeSegment(0, depth); + + RAJA::resources::Host host_res; + RAJA::resources::Resource res(host_res); + + RAJA::launch_nd( + res, + RAJA::launch_nd_grid_policy(RAJA::LaunchParams()), + RAJA::nd_segments(rows, batches, depths), [&](int r, int b, int d) { + values[d + depth * (b + batch_size * r)] = 100 * r + 10 * b + d; + }); + + for (int r = 0; r < n; ++r) + { + for (int b = 0; b < batch_size; ++b) + { + for (int d = 0; d < depth; ++d) + { + ASSERT_EQ(values[d + depth * (b + batch_size * r)], + 100 * r + 10 * b + d); + } + } + } +} + +TEST(launch_nd, resource_place_mismatch_throws) +{ + RAJA::resources::Host host_res; + RAJA::resources::Resource res(host_res); + + auto rows = RAJA::TypedRangeSegment(0, 1); + auto batches = RAJA::TypedRangeSegment(0, 1); + + // Need gtest death test to avoid complete failure due to eventual seg fault +#if defined(RAJA_ENABLE_TARGET_OPENMP) + EXPECT_DEATH_IF_SUPPORTED( + (RAJA::launch_nd(res, RAJA::ExecPlace::DEVICE, + RAJA::make_launch_nd_place_policy(flat_policy {}, + flat_policy {}), + RAJA::nd_segments(rows, batches), [](int, int) {})), + ""); +#else + EXPECT_THROW((RAJA::launch_nd(res, RAJA::ExecPlace::DEVICE, + RAJA::make_launch_nd_place_policy(flat_policy {}, + flat_policy {}), + RAJA::nd_segments(rows, batches), + [](int, int) {})), + std::runtime_error); +#endif +}