// SPDX-License-Identifier: MIT // Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #pragma once #include "ck/utility/sequence.hpp" #include "ck/utility/array.hpp" #include "ck/utility/tuple.hpp" #include "ck/utility/macro_func_array_to_sequence.hpp" namespace ck { template __host__ __device__ constexpr auto make_sequence(Number...) { return Sequence{}; } // F() returns index_t // F use default constructor, so F cannot be lambda function template __host__ __device__ constexpr auto generate_sequence(F, Number) { return typename sequence_gen::type{}; } // F() returns Number<> // F could be lambda function template __host__ __device__ constexpr auto generate_sequence_v2(F&& f, Number) { return unpack([&f](auto&&... xs) { return make_sequence(f(xs)...); }, typename arithmetic_sequence_gen<0, N, 1>::type{}); } template __host__ __device__ constexpr auto to_sequence(Tuple...>) { return Sequence{}; } namespace detail { template struct sorted_sequence_histogram; template struct sorted_sequence_histogram, Sequence> { template constexpr auto operator()(Histogram& h) { if constexpr(x < r) { h.template At() += 1; sorted_sequence_histogram, Sequence>{}(h); } else { h.template At() = 1; sorted_sequence_histogram, Sequence>{}(h); } } }; template struct sorted_sequence_histogram, Sequence> { template constexpr auto operator()(Histogram& h) { if constexpr(x < r) { h.template At() += 1; } } }; } // namespace detail // SeqSortedSamples: <0, 2, 3, 5, 7>, SeqRange: <0, 3, 6, 9> -> SeqHistogram : <2, 2, 1> template constexpr auto histogram_sorted_sequence(SeqSortedSamples, Sequence) { constexpr auto bins = sizeof...(rs); // or categories constexpr auto histogram = [&]() { Array h{0}; // make sure this can clear all element to zero detail::sorted_sequence_histogram<0, SeqSortedSamples, Sequence>{}(h); return h; }(); return TO_SEQUENCE(histogram, bins); } } // namespace ck