TheRock/rocm-libraries/projects/rocthrust/benchmark/bench/equal/basic.cu

// SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved.
// SPDX-FileCopyrightText: Modifications Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

// Benchmark utils
#include "../../bench_utils/bench_utils.hpp"

// rocThrust
#include <thrust/device_vector.h>
#include <thrust/equal.h>
#include <thrust/execution_policy.h>

// Google Benchmark
#include <benchmark/benchmark.h>

// STL
#include <algorithm>
#include <cstddef>
#include <string>
#include <vector>

struct basic
{
  template <typename T, typename Policy>
  float64_t run(thrust::device_vector<T>& a, thrust::device_vector<T>& b, Policy policy)
  {
    bench_utils::do_not_optimize(thrust::equal(policy, a.begin(), a.end(), b.begin()));

    bench_utils::gpu_timer d_timer;

    d_timer.start(0);
    bench_utils::do_not_optimize(thrust::equal(policy, a.begin(), a.end(), b.begin()));
    d_timer.stop(0);

    return d_timer.get_duration();
  }
};

template <class Benchmark, class T>
void run_benchmark(benchmark::State& state,
                   const std::size_t elements,
                   const std::string, /*seed_type*/
                   const float64_t common_prefix_ratio)
{
  // Benchmark object
  Benchmark benchmark{};

  // GPU times
  std::vector<double> gpu_times;

  thrust::device_vector<T> a(elements, T{1});
  thrust::device_vector<T> b(elements, T{1});

  const auto same_elements = std::min(static_cast<std::size_t>(elements * common_prefix_ratio), elements);

  bench_utils::caching_allocator_t alloc;
  thrust::detail::device_t policy{};

  thrust::fill(policy(alloc), b.begin() + same_elements, b.end(), T{2});

  for (auto _ : state)
  {
    float64_t duration = benchmark.template run<T>(a, b, policy(alloc));
    state.SetIterationTime(duration);
    gpu_times.push_back(duration);
  }

  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
  // it will actually be the global memory bandwidth gotten.
  // using `same_elements` instead of `elements` corresponds to the
  // actual elements read in an early exit
  state.SetBytesProcessed(state.iterations() * 2 * std::max(same_elements, std::size_t(1)) * sizeof(T));
  state.SetItemsProcessed(state.iterations() * std::max(same_elements, std::size_t(1)));

  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
  state.counters["gpu_noise"] = gpu_cv;
}

#define CREATE_BENCHMARK(T, Elements, CommonPrefixRatio)                                                    \
  benchmark::RegisterBenchmark(                                                                             \
    bench_utils::bench_naming::format_name(                                                                 \
      "{algo:equal,subalgo:" + name + ",input_type:" #T + ",elements:" + bench_utils::format_pow2(Elements) \
      + ", common_prefix_ratio:" #CommonPrefixRatio)                                                        \
      .c_str(),                                                                                             \
    run_benchmark<Benchmark, T>,                                                                            \
    Elements,                                                                                               \
    seed_type,                                                                                              \
    CommonPrefixRatio)

#define BENCHMARK_ELEMENTS(type, elements)             \
  bs.push_back(CREATE_BENCHMARK(type, elements, 1.0)); \
  bs.push_back(CREATE_BENCHMARK(type, elements, 0.5)); \
  bs.push_back(CREATE_BENCHMARK(type, elements, 0.0));

#define BENCHMARK_TYPE(type)                                               \
  for (size_t size : bench_utils::sizes)                                   \
  {                                                                        \
    if (sizeof(type) * size <= bench_utils::system.devProp.totalGlobalMem) \
      BENCHMARK_ELEMENTS(type, size)                                       \
  }

template <class Benchmark>
void add_benchmarks(
  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
{
  std::vector<benchmark::internal::Benchmark*> bs;
  BENCHMARK_TYPE(int8_t)
  BENCHMARK_TYPE(int16_t)
  BENCHMARK_TYPE(int32_t)
  BENCHMARK_TYPE(uint32_t)
  BENCHMARK_TYPE(int64_t)
  BENCHMARK_TYPE(uint64_t)

  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
}

int main(int argc, char* argv[])
{
  cli::Parser parser(argc, argv);
  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
  parser.run_and_exit_if_error();

  // Parse argv
  benchmark::Initialize(&argc, argv);
  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
  const std::string seed_type = parser.get<std::string>("seed");

  // Benchmark info
  bench_utils::add_common_benchmark_info();
  benchmark::AddCustomContext("seed", seed_type);

  // Add benchmark
  std::vector<benchmark::internal::Benchmark*> benchmarks;
  add_benchmarks<basic>("basic", benchmarks, seed_type);

  // Use manual timing
  for (auto& b : benchmarks)
  {
    b->UseManualTime();
    b->Unit(benchmark::kMicrosecond);
    b->MinTime(0.4); // in seconds
  }

  // Run benchmarks
  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());

  // Finish
  benchmark::Shutdown();
  return 0;
}