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; }; diff --git a/shared/primbench/primbench.hpp b/shared/primbench/primbench.hpp index 323c8cc677a4..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(); @@ -2716,8 +2721,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;