// MIT License // // Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal // in the Software without restriction, including without limitation the rights // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell // copies of the Software, and to permit persons to whom the Software is // furnished to do so, subject to the following conditions: // // The above copyright notice and this permission notice shall be included in all // copies or substantial portions of the Software. // // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. #pragma once #include "primbench.hpp" #include "benchmark_utils.hpp" #include "../common/predicate_iterator.hpp" #include "../common/utils_custom_type.hpp" #include "../common/utils_device_ptr.hpp" #include #include #include #include #include #include #include #include #include #include template struct less_than { __device__ bool operator()(T value) const { return value < T{C}; } }; template struct transform_op { __device__ auto operator()(T v) const { return Predicate{}(v) ? Transform{}(v) : v; } }; template struct transform_it { using value_type = T; void operator()(T* d_input, T* d_output, const size_t items, const hipStream_t stream) { auto t_it = rocprim::make_transform_iterator(d_input, transform_op{}); HIP_CHECK(rocprim::transform(t_it, d_output, items, rocprim::identity{}, stream)); } }; template struct read_predicate_it { using value_type = T; void operator()(T* d_input, T* d_output, const size_t items, const hipStream_t stream) { auto t_it = rocprim::make_transform_iterator(d_input, Transform{}); auto r_it = rocprim::make_predicate_iterator(t_it, d_input, Predicate{}); HIP_CHECK(rocprim::transform(r_it, d_output, items, rocprim::identity{}, stream)); } }; template struct write_predicate_it { using value_type = T; void operator()(T* d_input, T* d_output, const size_t items, const hipStream_t stream) { auto t_it = rocprim::make_transform_iterator(d_input, Transform{}); auto w_it = rocprim::make_predicate_iterator(d_output, d_input, Predicate{}); HIP_CHECK(rocprim::transform(t_it, w_it, items, rocprim::identity{}, stream)); } }; template struct subalgo_name; template struct subalgo_name> { static std::string value() { return "transform"; } }; template struct subalgo_name> { static std::string value() { return "read_predicate"; } }; template struct subalgo_name> { static std::string value() { return "write_predicate"; } }; template std::string get_subalgo_name() { return subalgo_name::value(); } template struct predicate_iterator_benchmark : public primbench::benchmark_interface { primbench::json meta() const override { return primbench::json{} .add("lvl", "device") .add("algo", "predicate_iterator") .add("subalgo", get_subalgo_name()) .add("percent", C) .add("key_type", primbench::name()) .add("cfg", "default"); } void run(primbench::state& state) override { const auto& stream = state.stream; const auto& bytes = state.size; const auto& seed = state.seed; size_t items = bytes / sizeof(T); const auto random_range = limit_random_range(0, 99); std::vector input = get_random_data(items, random_range.first, random_range.second, seed); common::device_ptr d_input(input); common::device_ptr d_output(items); state.set_items(items); state.add_reads(items); state.run([&] { IteratorBenchmark{}(d_input.get(), d_output.get(), items, stream); }); } private: using T = typename IteratorBenchmark::value_type; };