From c14d7592615d1cb81be2c6abf2136051fb554a9f Mon Sep 17 00:00:00 2001 From: Sander Bos Date: Mon, 7 Sep 2026 14:32:51 +0000 Subject: [PATCH 1/3] fix(primbench): `m_last_bytes_per_second` was always zero, ruining autotuning I followed the instructions in projects/rocprim/tuning/README.md to confirm that the generated default config now matches the fastest specialization. The reason this problem was hidden is because the autotuning system always picked the existing default config, which was already the fastest specialization. By passing `--exclude-default-config` to `run_tuning.py` the autotuning system before this fix commit would output terribly slow configs. --- shared/primbench/primbench.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/shared/primbench/primbench.hpp b/shared/primbench/primbench.hpp index 323c8cc677a4..84c34cbf32dd 100644 --- a/shared/primbench/primbench.hpp +++ b/shared/primbench/primbench.hpp @@ -2716,8 +2716,9 @@ class state elapsed_gpu_secs += batch_gpu_secs; - double bytes_per_batch = m_read_write_bytes * m_kernels_per_batch; - double bytes_per_sec = bytes_per_batch / batch_gpu_secs; + double bytes_per_batch = m_read_write_bytes * m_kernels_per_batch; + double bytes_per_sec = bytes_per_batch / batch_gpu_secs; + m_last_bytes_per_second = bytes_per_sec; double items_per_batch = m_items * m_kernels_per_batch; double items_per_sec = items_per_batch / batch_gpu_secs; From fb3fd0f8fe620ed302dea0634936a265082e9442 Mon Sep 17 00:00:00 2001 From: Sander Bos Date: Tue, 8 Sep 2026 10:24:16 +0000 Subject: [PATCH 2/3] fix(primbench): can't call `state.run()` twice In rocPRIM there are several benchmarks which call `state.run()` multiple times during tuning, like `projects/rocprim/benchmark/benchmark_device_segmented_reduce.hpp` calling it for 5 different segment sizes. This caused the bytes/sec that was returned to Kernel Tuner to be inaccurate, and for the same specialization to be printed 5x in the terminal, which was very confusing. --- shared/primbench/primbench.hpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/shared/primbench/primbench.hpp b/shared/primbench/primbench.hpp index 84c34cbf32dd..792b5239c2e5 100644 --- a/shared/primbench/primbench.hpp +++ b/shared/primbench/primbench.hpp @@ -2588,6 +2588,11 @@ class state std::cerr << "Error: Can't call run() before calling set_items()\n"; exit(EXIT_FAILURE); } + if(m_has_run) + { + std::cerr << "Error: Can't call run() twice\n"; + exit(EXIT_FAILURE); + } m_has_run = true; std::string name = m_meta.serialize_name(); From 5955278fbced8d775b38a63b188f3077810ff2ac Mon Sep 17 00:00:00 2001 From: Nick Breed Date: Tue, 8 Sep 2026 12:48:53 +0000 Subject: [PATCH 3/3] fix(rocprim): benchmarks updated to only call run ones for autotuning --- .../benchmark/benchmark_device_histogram.hpp | 20 +- ...hmark_device_segmented_radix_sort_keys.hpp | 229 +++++++++------ ...mark_device_segmented_radix_sort_pairs.hpp | 273 +++++++++++------- .../benchmark_device_segmented_reduce.hpp | 135 +++++---- .../benchmark_device_segmented_scan.hpp | 132 +++++---- 5 files changed, 475 insertions(+), 314 deletions(-) diff --git a/projects/rocprim/benchmark/benchmark_device_histogram.hpp b/projects/rocprim/benchmark/benchmark_device_histogram.hpp index 15ebb0a3231f..99ca403dc456 100644 --- a/projects/rocprim/benchmark/benchmark_device_histogram.hpp +++ b/projects/rocprim/benchmark/benchmark_device_histogram.hpp @@ -466,14 +466,14 @@ struct device_multi_histogram_even_benchmark : public primbench::benchmark_inter } state.set_items(items); - - for(auto& data : cases_data) - { - T* d_input = data.get_d_input(bytes); - - state.run( - [&] + state.add_reads(items * Channels * cases_data.size()); + state.run( + [&] + { + for(auto& data : cases_data) { + T* d_input = data.get_d_input(bytes); + HIP_CHECK((rocprim::multi_histogram_even( d_temporary_storage.get(), temporary_storage_bytes, @@ -485,10 +485,8 @@ struct device_multi_histogram_even_benchmark : public primbench::benchmark_inter data.upper_level, stream, false))); - }); - - state.add_reads(items * Channels); - } + } + }); for(unsigned int channel = 0; channel < ActiveChannels; ++channel) { diff --git a/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_keys.hpp b/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_keys.hpp index 9405774334ee..14a0c61e1bf3 100644 --- a/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_keys.hpp +++ b/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_keys.hpp @@ -112,7 +112,7 @@ struct device_segmented_radix_sort_keys_benchmark : public primbench::benchmark_ return j; } - void run_benchmark(primbench::state&& state, size_t num_segments, size_t mean_segment_length) + void run(primbench::state& state) override { const auto& stream = state.stream; const auto& seed = state.seed; @@ -120,117 +120,168 @@ struct device_segmented_radix_sort_keys_benchmark : public primbench::benchmark_ using offset_type = int; using key_type = Key; - primbench::log("Creating offsets"); - std::vector offsets; - offsets.push_back(0); + // Collect the (num_segments, mean_segment_length) pairs to actually run + std::vector> combos; + + if(m_segment_counts.size() == 1) + { + combos.emplace_back(m_segment_counts[0], m_segment_lengths[0]); + } + else + { + constexpr size_t min_size = 300000; + constexpr size_t max_size = 33554432; + + for(const auto segment_count : m_segment_counts) + { + for(const auto segment_length : m_segment_lengths) + { + const auto number_of_elements = segment_count * segment_length; + if(number_of_elements < min_size || number_of_elements > max_size) + { + continue; + } + combos.emplace_back(segment_count, segment_length); + } + } + } + + const int num_input_arrays = combos.size(); - primbench::log("Creating gen"); static constexpr int iseed = 716; engine_type gen(iseed); - primbench::log("Generating segment_length_dis"); - std::normal_distribution segment_length_dis( - static_cast(mean_segment_length), - 0.1 * mean_segment_length); + // Build offsets + keys for every combo up front + std::vector> offsets_arrays(num_input_arrays); + std::vector> keys_input_arrays(num_input_arrays); + std::vector items_per_run(num_input_arrays); + std::vector segments_counts(num_input_arrays); + size_t max_items = 0; - primbench::log("Calculating offsets"); - size_t offset = 0; - for(size_t segment_index = 0; segment_index < num_segments;) + for(int i = 0; i < num_input_arrays; ++i) { - const double segment_length_candidate = std::round(segment_length_dis(gen)); - if(segment_length_candidate < 0) + const size_t num_segments = combos[i].first; + const size_t mean_segment_length = combos[i].second; + + primbench::log("Creating offsets"); + std::vector& offsets = offsets_arrays[i]; + offsets.push_back(0); + + primbench::log("Generating segment_length_dis"); + std::normal_distribution segment_length_dis( + static_cast(mean_segment_length), + 0.1 * mean_segment_length); + + primbench::log("Calculating offsets"); + size_t offset = 0; + for(size_t segment_index = 0; segment_index < num_segments;) { - continue; + const double segment_length_candidate = std::round(segment_length_dis(gen)); + if(segment_length_candidate < 0) + { + continue; + } + const offset_type segment_length + = static_cast(segment_length_candidate); + offset += segment_length; + offsets.push_back(offset); + ++segment_index; } - const offset_type segment_length = static_cast(segment_length_candidate); - offset += segment_length; - offsets.push_back(offset); - ++segment_index; + const size_t items = offset; + + items_per_run[i] = items; + segments_counts[i] = offsets.size() - 1; + max_items = std::max(max_items, items); + + primbench::log("Generating keys_input"); + keys_input_arrays[i] + = get_random_data(items, + common::generate_limits::min(), + common::generate_limits::max(), + seed); } - const size_t items = offset; - const size_t segments_count = offsets.size() - 1; - primbench::log("Generating keys_input"); - std::vector keys_input - = get_random_data(items, - common::generate_limits::min(), - common::generate_limits::max(), - seed); + primbench::log("Creating d_offsets_arrays"); + std::vector> d_offsets_arrays(num_input_arrays); + for(int i = 0; i < num_input_arrays; ++i) + { + d_offsets_arrays[i].store(offsets_arrays[i]); + } - primbench::log("Creating d_offsets"); - common::device_ptr d_offsets(offsets); + primbench::log("Creating d_keys_input_arrays"); + std::vector> d_keys_input_arrays(num_input_arrays); + for(int i = 0; i < num_input_arrays; ++i) + { + d_keys_input_arrays[i].store(keys_input_arrays[i]); + } - primbench::log("Creating d_keys_input"); - common::device_ptr d_keys_input(keys_input); + // Shared, reused across all runs (sized to the largest) primbench::log("Creating d_keys_output"); - common::device_ptr d_keys_output(items); - + common::device_ptr d_keys_output(max_items); + + // Single call to segmented_radix_sort_keys for one combo, with its own storage size. + const auto dispatch_input = [&](void* d_temp_storage, + size_t& temp_storage_size_bytes, + key_type* d_keys_input, + size_t items, + size_t segments_count, + offset_type* d_offsets) + { + HIP_CHECK(rocprim::segmented_radix_sort_keys(d_temp_storage, + temp_storage_size_bytes, + d_keys_input, + d_keys_output.get(), + items, + segments_count, + d_offsets, + d_offsets + 1, + 0, + sizeof(key_type) * 8, + stream, + false)); + }; + + // Size each combo independently and track the max required storage. primbench::log("Calculating d_temporary_storage size"); - size_t temporary_storage_bytes = 0; - HIP_CHECK(rocprim::segmented_radix_sort_keys(nullptr, - temporary_storage_bytes, - d_keys_input.get(), - d_keys_output.get(), - items, - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - 0, - sizeof(key_type) * 8, - stream, - false)); + std::vector temp_storage_bytes_per_run(num_input_arrays); + size_t max_temporary_storage_bytes = 0; + + for(int i = 0; i < num_input_arrays; ++i) + { + size_t temp_storage_size_bytes = 0; + dispatch_input(nullptr, + temp_storage_size_bytes, + d_keys_input_arrays[i].get(), + items_per_run[i], + segments_counts[i], + d_offsets_arrays[i].get()); + + temp_storage_bytes_per_run[i] = temp_storage_size_bytes; + max_temporary_storage_bytes + = std::max(max_temporary_storage_bytes, temp_storage_size_bytes); + } primbench::log("Resizing d_temporary_storage"); - common::device_ptr d_temporary_storage(temporary_storage_bytes); + common::device_ptr d_temporary_storage(max_temporary_storage_bytes); - state.set_items(items); - state.add_reads(items); + const size_t total_items + = std::accumulate(items_per_run.begin(), items_per_run.end(), size_t{0}); + state.set_items(total_items); + state.add_reads(total_items); state.run( [&] { - HIP_CHECK(rocprim::segmented_radix_sort_keys(d_temporary_storage.get(), - temporary_storage_bytes, - d_keys_input.get(), - d_keys_output.get(), - items, - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - 0, - sizeof(key_type) * 8, - stream, - false)); - }); - } - - void run(primbench::state& state) override - { - if(m_segment_counts.size() == 1) - { - run_benchmark(std::forward(state), - m_segment_counts[0], - m_segment_lengths[0]); - return; - } - - constexpr size_t min_size = 300000; - constexpr size_t max_size = 33554432; - - // TODO: Replace with KernelTuner-based autotuning that generates one benchmark per segment_count+length combo. - for(const auto segment_count : m_segment_counts) - { - for(const auto segment_length : m_segment_lengths) - { - const auto number_of_elements = segment_count * segment_length; - if(number_of_elements < min_size || number_of_elements > max_size) + for(int i = 0; i < num_input_arrays; ++i) { - continue; + dispatch_input(d_temporary_storage.get(), + temp_storage_bytes_per_run[i], + d_keys_input_arrays[i].get(), + items_per_run[i], + segments_counts[i], + d_offsets_arrays[i].get()); } - - run_benchmark(std::forward(state), segment_count, segment_length); - } - } + }); } private: diff --git a/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_pairs.hpp b/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_pairs.hpp index 917226b05b7c..42ac69459b2a 100644 --- a/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_pairs.hpp +++ b/projects/rocprim/benchmark/benchmark_device_segmented_radix_sort_pairs.hpp @@ -113,7 +113,7 @@ struct device_segmented_radix_sort_pairs_benchmark : public primbench::benchmark return j; } - void run_benchmark(primbench::state&& state, size_t num_segments, size_t mean_segment_length) + void run(primbench::state& state) override { const auto& stream = state.stream; const auto& seed = state.seed; @@ -122,134 +122,191 @@ struct device_segmented_radix_sort_pairs_benchmark : public primbench::benchmark using key_type = Key; using value_type = Value; - primbench::log("Creating offsets"); - std::vector offsets; - offsets.push_back(0); + // Collect the (num_segments, mean_segment_length) pairs to actually run + std::vector> combos; + + if(m_segment_counts.size() == 1) + { + combos.emplace_back(m_segment_counts[0], m_segment_lengths[0]); + } + else + { + constexpr size_t min_size = 300000; + constexpr size_t max_size = 33554432; + + for(const auto segment_count : m_segment_counts) + { + for(const auto segment_length : m_segment_lengths) + { + const auto number_of_elements = segment_count * segment_length; + if(number_of_elements < min_size || number_of_elements > max_size) + { + continue; + } + combos.emplace_back(segment_count, segment_length); + } + } + } + + const int num_input_arrays = combos.size(); - primbench::log("Creating gen"); static constexpr int iseed = 716; engine_type gen(iseed); - primbench::log("Generating segment_length_dis"); - std::normal_distribution segment_length_dis( - static_cast(mean_segment_length), - 0.1 * mean_segment_length); + // Build offsets + keys + values for every combo up front + std::vector> offsets_arrays(num_input_arrays); + std::vector> keys_input_arrays(num_input_arrays); + std::vector> values_input_arrays(num_input_arrays); + std::vector items_per_run(num_input_arrays); + std::vector segments_counts(num_input_arrays); + size_t max_items = 0; - primbench::log("Calculating offsets"); - size_t offset = 0; - for(size_t segment_index = 0; segment_index < num_segments;) + for(int i = 0; i < num_input_arrays; ++i) { - const double segment_length_candidate = std::round(segment_length_dis(gen)); - if(segment_length_candidate < 0) + const size_t num_segments = combos[i].first; + const size_t mean_segment_length = combos[i].second; + + primbench::log("Creating offsets"); + std::vector& offsets = offsets_arrays[i]; + offsets.push_back(0); + + primbench::log("Generating segment_length_dis"); + std::normal_distribution segment_length_dis( + static_cast(mean_segment_length), + 0.1 * mean_segment_length); + + primbench::log("Calculating offsets"); + size_t offset = 0; + for(size_t segment_index = 0; segment_index < num_segments;) { - continue; + const double segment_length_candidate = std::round(segment_length_dis(gen)); + if(segment_length_candidate < 0) + { + continue; + } + const offset_type segment_length + = static_cast(segment_length_candidate); + offset += segment_length; + offsets.push_back(offset); + ++segment_index; } - const offset_type segment_length = static_cast(segment_length_candidate); - offset += segment_length; - offsets.push_back(offset); - ++segment_index; + const size_t items = offset; + + items_per_run[i] = items; + segments_counts[i] = offsets.size() - 1; + max_items = std::max(max_items, items); + + primbench::log("Generating keys_input"); + keys_input_arrays[i] + = get_random_data(items, + common::generate_limits::min(), + common::generate_limits::max(), + seed); + + primbench::log("Generating values_input"); + values_input_arrays[i] + = get_random_data(items, + common::generate_limits::min(), + common::generate_limits::max(), + seed); } - const size_t items = offset; - const size_t segments_count = offsets.size() - 1; - - primbench::log("Generating keys_input"); - std::vector keys_input - = get_random_data(items, - common::generate_limits::min(), - common::generate_limits::max(), - seed); - - primbench::log("Generating values_input"); - std::vector values_input - = get_random_data(items, - common::generate_limits::min(), - common::generate_limits::max(), - seed); - - primbench::log("Creating d_offsets"); - common::device_ptr d_offsets(offsets); - - primbench::log("Creating d_keys_input"); - common::device_ptr d_keys_input(keys_input); - primbench::log("Creating d_keys_output"); - common::device_ptr d_keys_output(items); - primbench::log("Creating d_values_input"); - common::device_ptr d_values_input(values_input); - primbench::log("Creating d_values_output"); - common::device_ptr d_values_output(items); - - primbench::log("Calculating d_temporary_storage size"); - size_t temporary_storage_bytes = 0; - HIP_CHECK(rocprim::segmented_radix_sort_pairs(nullptr, - temporary_storage_bytes, - d_keys_input.get(), - d_keys_output.get(), - d_values_input.get(), - d_values_output.get(), - items, - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - 0, - sizeof(key_type) * 8, - stream, - false)); + primbench::log("Creating d_offsets_arrays"); + std::vector> d_offsets_arrays(num_input_arrays); + for(int i = 0; i < num_input_arrays; ++i) + { + d_offsets_arrays[i].store(offsets_arrays[i]); + } - primbench::log("Resizing d_temporary_storage"); - common::device_ptr d_temporary_storage(temporary_storage_bytes); + primbench::log("Creating d_keys_input_arrays"); + std::vector> d_keys_input_arrays(num_input_arrays); + for(int i = 0; i < num_input_arrays; ++i) + { + d_keys_input_arrays[i].store(keys_input_arrays[i]); + } - state.set_items(items); - state.add_reads(items); - state.add_reads(items); + primbench::log("Creating d_values_input_arrays"); + std::vector> d_values_input_arrays(num_input_arrays); + for(int i = 0; i < num_input_arrays; ++i) + { + d_values_input_arrays[i].store(values_input_arrays[i]); + } - state.run( - [&] - { - HIP_CHECK(rocprim::segmented_radix_sort_pairs(d_temporary_storage.get(), - temporary_storage_bytes, - d_keys_input.get(), - d_keys_output.get(), - d_values_input.get(), - d_values_output.get(), - items, - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - 0, - sizeof(key_type) * 8, - stream, - false)); - }); - } + // Shared, reused across all runs (sized to the largest) + primbench::log("Creating d_keys_output"); + common::device_ptr d_keys_output(max_items); + primbench::log("Creating d_values_output"); + common::device_ptr d_values_output(max_items); + + // Single call to segmented_radix_sort_pairs for one combo, with its own storage size. + const auto dispatch_input = [&](void* d_temp_storage, + size_t& temp_storage_size_bytes, + key_type* d_keys_input, + value_type* d_values_input, + size_t items, + size_t segments_count, + offset_type* d_offsets) + { + HIP_CHECK(rocprim::segmented_radix_sort_pairs(d_temp_storage, + temp_storage_size_bytes, + d_keys_input, + d_keys_output.get(), + d_values_input, + d_values_output.get(), + items, + segments_count, + d_offsets, + d_offsets + 1, + 0, + sizeof(key_type) * 8, + stream, + false)); + }; + + // Size each combo independently and track the max required storage. + primbench::log("Calculating d_temporary_storage size"); + std::vector temp_storage_bytes_per_run(num_input_arrays); + size_t max_temporary_storage_bytes = 0; - void run(primbench::state& state) override - { - if(m_segment_counts.size() == 1) + for(int i = 0; i < num_input_arrays; ++i) { - run_benchmark(std::forward(state), - m_segment_counts[0], - m_segment_lengths[0]); - return; + size_t temp_storage_size_bytes = 0; + dispatch_input(nullptr, + temp_storage_size_bytes, + d_keys_input_arrays[i].get(), + d_values_input_arrays[i].get(), + items_per_run[i], + segments_counts[i], + d_offsets_arrays[i].get()); + + temp_storage_bytes_per_run[i] = temp_storage_size_bytes; + max_temporary_storage_bytes + = std::max(max_temporary_storage_bytes, temp_storage_size_bytes); } - constexpr size_t min_size = 300000; - constexpr size_t max_size = 33554432; + primbench::log("Resizing d_temporary_storage"); + common::device_ptr d_temporary_storage(max_temporary_storage_bytes); - // TODO: Replace with KernelTuner-based autotuning that generates one benchmark per segment_count+length combo. - for(const auto segment_count : m_segment_counts) - { - for(const auto segment_length : m_segment_lengths) + const size_t total_items + = std::accumulate(items_per_run.begin(), items_per_run.end(), size_t{0}); + state.set_items(total_items); + state.add_reads(total_items); + state.add_reads(total_items); + + state.run( + [&] { - const auto number_of_elements = segment_count * segment_length; - if(number_of_elements < min_size || number_of_elements > max_size) + for(int i = 0; i < num_input_arrays; ++i) { - continue; + dispatch_input(d_temporary_storage.get(), + temp_storage_bytes_per_run[i], + d_keys_input_arrays[i].get(), + d_values_input_arrays[i].get(), + items_per_run[i], + segments_counts[i], + d_offsets_arrays[i].get()); } - - run_benchmark(std::forward(state), segment_count, segment_length); - } - } + }); } private: diff --git a/projects/rocprim/benchmark/benchmark_device_segmented_reduce.hpp b/projects/rocprim/benchmark/benchmark_device_segmented_reduce.hpp index 75f81aa67b70..fe62b6251db4 100644 --- a/projects/rocprim/benchmark/benchmark_device_segmented_reduce.hpp +++ b/projects/rocprim/benchmark/benchmark_device_segmented_reduce.hpp @@ -100,7 +100,7 @@ struct device_segmented_reduce_benchmark : public primbench::benchmark_interface return j; } - void run_benchmark(primbench::state&& state, size_t desired_segment) + void run(primbench::state& state) override { const auto& stream = state.stream; const auto& bytes = state.size; @@ -109,82 +109,109 @@ struct device_segmented_reduce_benchmark : public primbench::benchmark_interface using offset_type = int; using value_type = T; - size_t items = bytes / sizeof(T); + const size_t items = bytes / sizeof(T); + const int num_input_arrays = m_desired_segments.size(); - // Generate data + // Shared input values, generated once engine_type gen(seed); - // The minimal average length should at least be 1 to prevent infinite loop. - const double avg_segment_length - = std::max(1.0, static_cast(items) / desired_segment); - std::uniform_real_distribution segment_length_dis(0, avg_segment_length * 2); + std::vector values_input(items); + std::iota(values_input.begin(), values_input.end(), 0); + common::device_ptr d_values_input(values_input); + + // Build offsets per desired segment count + std::vector> offsets_arrays(num_input_arrays); + std::vector segments_counts(num_input_arrays); + unsigned int max_segments_count = 0; - std::vector offsets; - unsigned int segments_count = 0; - size_t offset = 0; - while(offset < items) + for(int i = 0; i < num_input_arrays; ++i) { - const size_t segment_length = std::round(segment_length_dis(gen)); - offsets.push_back(offset); - segments_count++; - offset += segment_length; - } - offsets.push_back(items); + const size_t desired_segment = m_desired_segments[i]; - std::vector values_input(items); - std::iota(values_input.begin(), values_input.end(), 0); + const double avg_segment_length + = std::max(1.0, static_cast(items) / desired_segment); + std::uniform_real_distribution segment_length_dis(0, avg_segment_length * 2); - common::device_ptr d_offsets(offsets); + std::vector& offsets = offsets_arrays[i]; + unsigned int segments_count = 0; + size_t offset = 0; + while(offset < items) + { + const size_t segment_length = std::round(segment_length_dis(gen)); + offsets.push_back(offset); + segments_count++; + offset += segment_length; + } + offsets.push_back(items); + + segments_counts[i] = segments_count; + max_segments_count = std::max(max_segments_count, segments_count); + } - common::device_ptr d_values_input(values_input); + std::vector> d_offsets_arrays(num_input_arrays); + for(int i = 0; i < num_input_arrays; ++i) + { + d_offsets_arrays[i].store(offsets_arrays[i]); + } - common::device_ptr d_aggregates_output(segments_count); + // Shared, reused across all runs (sized to the largest) + common::device_ptr d_aggregates_output(max_segments_count); BinaryFunction reduce_op; value_type init(0); - size_t temporary_storage_bytes = 0; - - HIP_CHECK(rocprim::segmented_reduce(nullptr, - temporary_storage_bytes, - d_values_input.get(), - d_aggregates_output.get(), - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - reduce_op, - init, - stream)); + const auto dispatch_input = [&](void* d_temp_storage, + size_t& temp_storage_size_bytes, + offset_type* d_offsets, + unsigned int segments_count) + { + HIP_CHECK(rocprim::segmented_reduce(d_temp_storage, + temp_storage_size_bytes, + d_values_input.get(), + d_aggregates_output.get(), + segments_count, + d_offsets, + d_offsets + 1, + reduce_op, + init, + stream)); + }; + + // Query temp storage size per run, track the max + std::vector temp_storage_bytes_per_run(num_input_arrays); + size_t max_temporary_storage_bytes = 0; + + for(int i = 0; i < num_input_arrays; ++i) + { + size_t temp_storage_size_bytes = 0; + dispatch_input(nullptr, + temp_storage_size_bytes, + d_offsets_arrays[i].get(), + segments_counts[i]); + + temp_storage_bytes_per_run[i] = temp_storage_size_bytes; + max_temporary_storage_bytes + = std::max(max_temporary_storage_bytes, temp_storage_size_bytes); + } - common::device_ptr d_temporary_storage(temporary_storage_bytes); + common::device_ptr d_temp_storage(max_temporary_storage_bytes); state.set_items(items); - state.add_reads(items); + state.add_reads(items * num_input_arrays); state.run( [&] { - HIP_CHECK(rocprim::segmented_reduce(d_temporary_storage.get(), - temporary_storage_bytes, - d_values_input.get(), - d_aggregates_output.get(), - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - reduce_op, - init, - stream)); + for(int i = 0; i < num_input_arrays; ++i) + { + dispatch_input(d_temp_storage.get(), + temp_storage_bytes_per_run[i], + d_offsets_arrays[i].get(), + segments_counts[i]); + } }); } - void run(primbench::state& state) override - { - for(const auto desired_segment : m_desired_segments) - { - run_benchmark(std::forward(state), desired_segment); - } - } - private: std::vector m_desired_segments; }; diff --git a/projects/rocprim/benchmark/benchmark_device_segmented_scan.hpp b/projects/rocprim/benchmark/benchmark_device_segmented_scan.hpp index e2a5a8609608..c12767dbb3e1 100644 --- a/projects/rocprim/benchmark/benchmark_device_segmented_scan.hpp +++ b/projects/rocprim/benchmark/benchmark_device_segmented_scan.hpp @@ -138,7 +138,7 @@ struct device_segmented_scan_benchmark : public primbench::benchmark_interface } } - void run_benchmark(primbench::state&& state, size_t desired_segment) + void run(primbench::state& state) override { const auto& stream = state.stream; const auto& bytes = state.size; @@ -147,80 +147,108 @@ struct device_segmented_scan_benchmark : public primbench::benchmark_interface using offset_type = int; using value_type = T; - size_t items = bytes / sizeof(T); + const size_t items = bytes / sizeof(T); + const int num_input_arrays = m_desired_segments.size(); - // Generate data + // Shared input values, generated once engine_type gen(seed); - // The minimal average length should at least be 1 to prevent infinite loop. - const double avg_segment_length - = std::max(1.0, static_cast(items) / desired_segment); - std::uniform_real_distribution segment_length_dis(0, avg_segment_length * 2); + std::vector values_input(items); + std::iota(values_input.begin(), values_input.end(), 0); + common::device_ptr d_values_input(values_input); + + // Build offsets per desired segment count + std::vector> offsets_arrays(num_input_arrays); + std::vector segments_counts(num_input_arrays); + unsigned int max_segments_count = 0; - // Create random sizes for segments. - std::vector offsets; - unsigned int segments_count = 0; - size_t offset = 0; - while(offset < items) + for(int i = 0; i < num_input_arrays; ++i) { - const size_t segment_length = std::round(segment_length_dis(gen)); - offsets.push_back(offset); - segments_count++; - offset += segment_length; - } - offsets.push_back(items); + const size_t desired_segment = m_desired_segments[i]; - std::vector values_input(items); - std::iota(values_input.begin(), values_input.end(), 0); + // The minimal average length should at least be 1 to prevent infinite loop. + const double avg_segment_length + = std::max(1.0, static_cast(items) / desired_segment); + std::uniform_real_distribution segment_length_dis(0, avg_segment_length * 2); - common::device_ptr d_offsets(offsets); + std::vector& offsets = offsets_arrays[i]; + unsigned int segments_count = 0; + size_t offset = 0; + while(offset < items) + { + const size_t segment_length = std::round(segment_length_dis(gen)); + offsets.push_back(offset); + segments_count++; + offset += segment_length; + } + offsets.push_back(items); + + segments_counts[i] = segments_count; + max_segments_count = std::max(max_segments_count, segments_count); + } - common::device_ptr d_values_input(values_input); + std::vector> d_offsets_arrays(num_input_arrays); + for(int i = 0; i < num_input_arrays; ++i) + { + d_offsets_arrays[i].store(offsets_arrays[i]); + } + // Shared, reused across all runs (sized to the largest) common::device_ptr d_values_output(items); value_type init(5); - size_t temporary_storage_bytes = 0; - - HIP_CHECK(run_device_segmented_scan(nullptr, - temporary_storage_bytes, - d_values_input.get(), - d_values_output.get(), - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - init, - stream)); + const auto dispatch_input = [&](void* d_temp_storage, + size_t& temp_storage_size_bytes, + offset_type* d_offsets, + unsigned int segments_count) + { + HIP_CHECK(run_device_segmented_scan(d_temp_storage, + temp_storage_size_bytes, + d_values_input.get(), + d_values_output.get(), + segments_count, + d_offsets, + d_offsets + 1, + init, + stream)); + }; + + // Query temp storage size per run, track the max + std::vector temp_storage_bytes_per_run(num_input_arrays); + size_t max_temporary_storage_bytes = 0; + + for(int i = 0; i < num_input_arrays; ++i) + { + size_t temp_storage_size_bytes = 0; + dispatch_input(nullptr, + temp_storage_size_bytes, + d_offsets_arrays[i].get(), + segments_counts[i]); + + temp_storage_bytes_per_run[i] = temp_storage_size_bytes; + max_temporary_storage_bytes + = std::max(max_temporary_storage_bytes, temp_storage_size_bytes); + } - common::device_ptr d_temporary_storage(temporary_storage_bytes); + common::device_ptr d_temp_storage(max_temporary_storage_bytes); state.set_items(items); - state.add_reads(items); + state.add_reads(items * num_input_arrays); state.run( [&] { - HIP_CHECK(run_device_segmented_scan(d_temporary_storage.get(), - temporary_storage_bytes, - d_values_input.get(), - d_values_output.get(), - segments_count, - d_offsets.get(), - d_offsets.get() + 1, - init, - stream)); + for(int i = 0; i < num_input_arrays; ++i) + { + dispatch_input(d_temp_storage.get(), + temp_storage_bytes_per_run[i], + d_offsets_arrays[i].get(), + segments_counts[i]); + } }); } - void run(primbench::state& state) override - { - for(const auto desired_segment : m_desired_segments) - { - run_benchmark(std::forward(state), desired_segment); - } - } - private: std::vector m_desired_segments; };