From 36d5eda3d5916c052505032bf5c3fcc3379dbbb0 Mon Sep 17 00:00:00 2001 From: Cong Ma Date: Sun, 18 Jan 2026 23:02:23 -0500 Subject: [PATCH 1/2] [CK TILE] Refactor sequence_reverse_inclusive_scan --- include/ck_tile/core/container/sequence.hpp | 56 +++++++++++++-------- 1 file changed, 36 insertions(+), 20 deletions(-) diff --git a/include/ck_tile/core/container/sequence.hpp b/include/ck_tile/core/container/sequence.hpp index 763d860d68..1eb276e116 100644 --- a/include/ck_tile/core/container/sequence.hpp +++ b/include/ck_tile/core/container/sequence.hpp @@ -13,6 +13,9 @@ namespace ck_tile { +template +struct array; // declare for later use (array->seq utility) + template struct sequence; @@ -332,29 +335,45 @@ struct uniform_sequence_gen }; // reverse inclusive scan (with init) sequence -template -struct sequence_reverse_inclusive_scan; +namespace impl { +template +struct sequence_reverse_inclusive_scan_impl; -template -struct sequence_reverse_inclusive_scan, Reduce, Init> +template +struct sequence_reverse_inclusive_scan_impl, Reduce, Init> { - using old_scan = typename sequence_reverse_inclusive_scan, Reduce, Init>::type; - - static constexpr index_t new_reduce = Reduce{}(I, old_scan{}.front()); - - using type = typename sequence_merge, old_scan>::type; -}; + template + static constexpr auto compute(sequence) + { + constexpr index_t size = sizeof...(Is); + if constexpr(size == 0) + { + return sequence<>{}; + } + else + { + constexpr array arr = []() { + array values = {Is...}; + array result = {0}; + result[size - 1] = Reduce{}(values[size - 1], Init); + for(index_t i = size - 1; i > 0; --i) + { + result[i - 1] = Reduce{}(values[i - 1], result[i]); + } + return result; + }(); + return sequence{}; + } + } -template -struct sequence_reverse_inclusive_scan, Reduce, Init> -{ - using type = sequence; + using type = decltype(compute(make_index_sequence{})); }; +} // namespace impl -template -struct sequence_reverse_inclusive_scan, Reduce, Init> +template +struct sequence_reverse_inclusive_scan { - using type = sequence<>; + using type = typename impl::sequence_reverse_inclusive_scan_impl::type; }; // split sequence @@ -1105,9 +1124,6 @@ struct sorted_sequence_histogram, sequence> }; } // namespace detail -template -struct array; // declare for later use (array->seq utility) - // SeqSortedSamples: <0, 2, 3, 5, 7>, SeqRange: <0, 3, 6, 9> -> SeqHistogram : <2, 2, 1> template CK_TILE_HOST_DEVICE constexpr auto histogram_sorted_sequence(SeqSortedSamples, sequence) From b337c160ed6ddabdce46ea589893b1d6d55702d8 Mon Sep 17 00:00:00 2001 From: Cong Ma Date: Mon, 19 Jan 2026 22:58:41 -0500 Subject: [PATCH 2/2] [CK TILE] Refactor sequence_reverse_inclusive_scan - Add trivial_array for template programming - Add unit tests of sequence --- include/ck_tile/core.hpp | 1 + include/ck_tile/core/container/sequence.hpp | 18 +- .../ck_tile/core/container/trivial_array.hpp | 21 + test/ck_tile/core/CMakeLists.txt | 1 + test/ck_tile/core/container/CMakeLists.txt | 7 + test/ck_tile/core/container/unit_sequence.cpp | 677 ++++++++++++++++++ 6 files changed, 716 insertions(+), 9 deletions(-) create mode 100644 include/ck_tile/core/container/trivial_array.hpp create mode 100644 test/ck_tile/core/container/CMakeLists.txt create mode 100644 test/ck_tile/core/container/unit_sequence.cpp diff --git a/include/ck_tile/core.hpp b/include/ck_tile/core.hpp index 01e1d00b59..0a0c230157 100644 --- a/include/ck_tile/core.hpp +++ b/include/ck_tile/core.hpp @@ -41,6 +41,7 @@ #include "ck_tile/core/container/span.hpp" #include "ck_tile/core/container/statically_indexed_array.hpp" #include "ck_tile/core/container/thread_buffer.hpp" +#include "ck_tile/core/container/trivial_array.hpp" #include "ck_tile/core/container/tuple.hpp" #include "ck_tile/core/numeric/bfloat16.hpp" #include "ck_tile/core/numeric/e8m0.hpp" diff --git a/include/ck_tile/core/container/sequence.hpp b/include/ck_tile/core/container/sequence.hpp index 1eb276e116..22f95763f3 100644 --- a/include/ck_tile/core/container/sequence.hpp +++ b/include/ck_tile/core/container/sequence.hpp @@ -4,6 +4,7 @@ #pragma once #include "ck_tile/core/config.hpp" +#include "ck_tile/core/container/trivial_array.hpp" #include "ck_tile/core/numeric/integer.hpp" #include "ck_tile/core/numeric/integral_constant.hpp" #include "ck_tile/core/numeric/math.hpp" @@ -13,9 +14,6 @@ namespace ck_tile { -template -struct array; // declare for later use (array->seq utility) - template struct sequence; @@ -38,6 +36,7 @@ template CK_TILE_HOST_DEVICE constexpr auto sequence_pop_back(Seq); namespace impl { + // static_assert(__has_builtin(__type_pack_element), "can't find __type_pack_element"); template using at_index_t = __type_pack_element; @@ -336,6 +335,7 @@ struct uniform_sequence_gen // reverse inclusive scan (with init) sequence namespace impl { + template struct sequence_reverse_inclusive_scan_impl; @@ -352,10 +352,10 @@ struct sequence_reverse_inclusive_scan_impl, Reduce, Init> } else { - constexpr array arr = []() { - array values = {Is...}; - array result = {0}; - result[size - 1] = Reduce{}(values[size - 1], Init); + constexpr auto arr = []() { + trivial_array values = {Is...}; + trivial_array result = {0}; + result[size - 1] = Reduce{}(values[size - 1], Init); for(index_t i = size - 1; i > 0; --i) { result[i - 1] = Reduce{}(values[i - 1], result[i]); @@ -1130,7 +1130,7 @@ CK_TILE_HOST_DEVICE constexpr auto histogram_sorted_sequence(SeqSortedSamples, s { constexpr auto bins = sizeof...(rs); // or categories constexpr auto histogram = [&]() { - array h{0}; // make sure this can clear all element to zero + trivial_array h{0}; // make sure this can clear all element to zero detail::sorted_sequence_histogram<0, SeqSortedSamples, sequence>{}(h); return h; }(); @@ -1143,7 +1143,7 @@ CK_TILE_HOST_DEVICE constexpr auto generate_array(F&& f, number) { using T = remove_cvref_t{}))>; - return unpack([&f](auto&&... is) { return array{f(is)...}; }, + return unpack([&f](auto&&... is) { return trivial_array{f(is)...}; }, typename arithmetic_sequence_gen<0, N, 1>::type{}); } diff --git a/include/ck_tile/core/container/trivial_array.hpp b/include/ck_tile/core/container/trivial_array.hpp new file mode 100644 index 0000000000..e1355a400b --- /dev/null +++ b/include/ck_tile/core/container/trivial_array.hpp @@ -0,0 +1,21 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once +#include "ck_tile/core/numeric/integer.hpp" + +namespace ck_tile { +// trivial_array has no custom constructor, allowing it to be used as a constexpr variable type +template +struct trivial_array +{ + // Public aggregate initialization makes this a literal type + T data[N]; + + // Basic constexpr accessors + constexpr const T& operator[](index_t i) const { return data[i]; } + constexpr T& operator[](index_t i) { return data[i]; } + + constexpr index_t size() const { return N; } +}; +} // namespace ck_tile diff --git a/test/ck_tile/core/CMakeLists.txt b/test/ck_tile/core/CMakeLists.txt index a46c2b1d41..90db6809ec 100644 --- a/test/ck_tile/core/CMakeLists.txt +++ b/test/ck_tile/core/CMakeLists.txt @@ -2,3 +2,4 @@ # SPDX-License-Identifier: MIT add_subdirectory(arch) +add_subdirectory(container) diff --git a/test/ck_tile/core/container/CMakeLists.txt b/test/ck_tile/core/container/CMakeLists.txt new file mode 100644 index 0000000000..02aa8f17c2 --- /dev/null +++ b/test/ck_tile/core/container/CMakeLists.txt @@ -0,0 +1,7 @@ +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +add_gtest_executable(ck_tile_unit_sequence unit_sequence.cpp) +if(result EQUAL 0) + target_link_libraries(ck_tile_unit_sequence PRIVATE utility) +endif() diff --git a/test/ck_tile/core/container/unit_sequence.cpp b/test/ck_tile/core/container/unit_sequence.cpp new file mode 100644 index 0000000000..3769d6ecf9 --- /dev/null +++ b/test/ck_tile/core/container/unit_sequence.cpp @@ -0,0 +1,677 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include "ck_tile/core/container/sequence.hpp" + +using namespace ck_tile; + +// Test basic Sequence construction and properties +TEST(Sequence, BasicConstruction) +{ + using Seq = sequence<1, 2, 3, 4, 5>; + EXPECT_EQ(Seq::size(), 5); +} + +TEST(Sequence, EmptySequence) +{ + using Seq = sequence<>; + EXPECT_EQ(Seq::size(), 0); +} + +// Test at() method +TEST(Sequence, AtRuntime) +{ + using Seq = sequence<10, 20, 30, 40>; + EXPECT_EQ(Seq::at(0), 10); + EXPECT_EQ(Seq::at(1), 20); + EXPECT_EQ(Seq::at(2), 30); + EXPECT_EQ(Seq::at(3), 40); +} + +TEST(Sequence, AtCompileTime) +{ + using Seq = sequence<10, 20, 30, 40>; + EXPECT_EQ(Seq::at(number<0>{}), 10); + EXPECT_EQ(Seq::at(number<1>{}), 20); + EXPECT_EQ(Seq::at(number<2>{}), 30); + EXPECT_EQ(Seq::at(number<3>{}), 40); +} + +TEST(Sequence, OperatorBracket) +{ + constexpr auto seq = sequence<5, 10, 15>{}; + EXPECT_EQ(seq[number<0>{}], 5); + EXPECT_EQ(seq[number<1>{}], 10); + EXPECT_EQ(seq[number<2>{}], 15); +} + +// Test Front() and Back() +TEST(Sequence, FrontBack) +{ + using Seq = sequence<100, 200, 300>; + EXPECT_EQ(Seq::front(), 100); + EXPECT_EQ(Seq::back(), 300); +} + +TEST(Sequence, FrontBackSingleElement) +{ + using Seq = sequence<42>; + EXPECT_EQ(Seq::front(), 42); + EXPECT_EQ(Seq::back(), 42); +} + +// Test PushFront and PushBack +TEST(Sequence, PushFront) +{ + using Seq = sequence<2, 3, 4>; + using Result = decltype(Seq::push_front(sequence<1>{})); + using Expected = sequence<1, 2, 3, 4>; + EXPECT_TRUE((std::is_same_v)); +} + +TEST(Sequence, PushFrontnumbers) +{ + using Seq = sequence<3, 4>; + using Result = decltype(Seq::push_front(number<1>{}, number<2>{})); + using Expected = sequence<1, 2, 3, 4>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(Sequence, PushBack) +{ + using Seq = sequence<1, 2, 3>; + using Result = decltype(Seq::push_back(sequence<4, 5>{})); + using Expected = sequence<1, 2, 3, 4, 5>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(Sequence, PushBacknumbers) +{ + using Seq = sequence<1, 2>; + using Result = decltype(Seq::push_back(number<3>{}, number<4>{})); + using Expected = sequence<1, 2, 3, 4>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test PopFront and PopBack +TEST(Sequence, PopFront) +{ + using Seq = sequence<1, 2, 3, 4>; + using Result = decltype(Seq::pop_front()); + using Expected = sequence<2, 3, 4>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(Sequence, PopBack) +{ + using Seq = sequence<1, 2, 3, 4>; + using Result = decltype(Seq::pop_back()); + using Expected = sequence<1, 2, 3>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test Extract +TEST(Sequence, ExtractBynumbers) +{ + using Seq = sequence<10, 20, 30, 40, 50>; + using Result = decltype(Seq::extract(number<0>{}, number<2>{}, number<4>{})); + using Expected = sequence<10, 30, 50>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(Sequence, ExtractBySequence) +{ + using Seq = sequence<10, 20, 30, 40, 50>; + using Result = decltype(Seq::extract(sequence<1, 3>{})); + using Expected = sequence<20, 40>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test Modify +TEST(Sequence, Modify) +{ + using Seq = sequence<1, 2, 3, 4>; + using Result = decltype(Seq::modify(number<2>{}, number<99>{})); + using Expected = sequence<1, 2, 99, 4>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test Transform +TEST(Sequence, Transform) +{ + using Seq = sequence<1, 2, 3, 4>; + auto double_it = [](auto x) { return 2 * x; }; + using Result = decltype(Seq::transform(double_it)); + using Expected = sequence<2, 4, 6, 8>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test Reverse +TEST(Sequence, Reverse) +{ + using Seq = sequence<1, 2, 3, 4, 5>; + using Result = decltype(Seq::reverse()); + using Expected = sequence<5, 4, 3, 2, 1>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(Sequence, ReverseSingleElement) +{ + using Seq = sequence<42>; + using Result = decltype(Seq::reverse()); + using Expected = sequence<42>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test ReorderGivenNew2Old +TEST(Sequence, ReorderGivenNew2Old) +{ + using Seq = sequence<10, 20, 30, 40>; + using Result = decltype(Seq::reorder_new_to_old(sequence<3, 1, 2, 0>{})); + using Expected = sequence<40, 20, 30, 10>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test ReorderGivenOld2New +TEST(Sequence, ReorderGivenOld2New) +{ + using Seq = sequence<10, 20, 30, 40>; + using Result = decltype(Seq::reorder_old_to_new(sequence<3, 1, 2, 0>{})); + using Expected = sequence<40, 20, 30, 10>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test arithmetic_sequence_gen +TEST(SequenceGen, ArithmeticSequence) +{ + using Result = typename arithmetic_sequence_gen<0, 5, 1>::type; + using Expected = sequence<0, 1, 2, 3, 4>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceGen, ArithmeticSequenceWithIncrement) +{ + using Result = typename arithmetic_sequence_gen<0, 10, 2>::type; + using Expected = sequence<0, 2, 4, 6, 8>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceGen, ArithmeticSequenceNegativeIncrement) +{ + using Result = typename arithmetic_sequence_gen<10, 5, -1>::type; + using Expected = sequence<10, 9, 8, 7, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceGen, ArithmeticSequenceEmpty) +{ + using Result = typename arithmetic_sequence_gen<5, 5, 1>::type; + using Expected = sequence<>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test uniform_sequence_gen +TEST(SequenceGen, UniformSequence) +{ + using Result = typename uniform_sequence_gen<5, 42>::type; + using Expected = sequence<42, 42, 42, 42, 42>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceGen, UniformSequenceZeroSize) +{ + using Result = typename uniform_sequence_gen<0, 42>::type; + using Expected = sequence<>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test make_index_sequence +TEST(SequenceGen, MakeIndexSequence) +{ + using Result = make_index_sequence<5>; + using Expected = sequence<0, 1, 2, 3, 4>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceGen, MakeIndexSequenceZero) +{ + using Result = make_index_sequence<0>; + using Expected = sequence<>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test sequence_merge +TEST(SequenceMerge, MergeTwoSequences) +{ + using Seq1 = sequence<1, 2, 3>; + using Seq2 = sequence<4, 5, 6>; + using Result = typename sequence_merge::type; + using Expected = sequence<1, 2, 3, 4, 5, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceMerge, MergeMultipleSequences) +{ + using Seq1 = sequence<1, 2>; + using Seq2 = sequence<3, 4>; + using Seq3 = sequence<5, 6>; + using Result = typename sequence_merge::type; + using Expected = sequence<1, 2, 3, 4, 5, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceMerge, MergeSingleSequence) +{ + using Seq = sequence<1, 2, 3>; + using Result = typename sequence_merge::type; + using Expected = sequence<1, 2, 3>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test sequence_split +TEST(SequenceSplit, SplitInMiddle) +{ + using Seq = sequence<1, 2, 3, 4, 5, 6>; + using Split = sequence_split; + using ExpectedLeft = sequence<1, 2, 3>; + using ExpectedRight = sequence<4, 5, 6>; + EXPECT_TRUE((std::is_same::value)); + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceSplit, SplitAtBeginning) +{ + using Seq = sequence<1, 2, 3, 4>; + using Split = sequence_split; + using ExpectedLeft = sequence<>; + using ExpectedRight = sequence<1, 2, 3, 4>; + EXPECT_TRUE((std::is_same::value)); + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceSplit, SplitAtEnd) +{ + using Seq = sequence<1, 2, 3, 4>; + using Split = sequence_split; + using ExpectedLeft = sequence<1, 2, 3, 4>; + using ExpectedRight = sequence<>; + EXPECT_TRUE((std::is_same::value)); + EXPECT_TRUE((std::is_same::value)); +} + +// Test sequence_sort +TEST(SequenceSort, SortAscending) +{ + using Seq = sequence<5, 2, 8, 1, 9>; + using Result = typename sequence_sort>::type; + using Expected = sequence<1, 2, 5, 8, 9>; + EXPECT_TRUE((std::is_same::value)); +} + +// TEST(SequenceSort, SortDescending) +// { +// // Create a greater-than comparator +// struct greater +// { +// __host__ constexpr bool operator()(index_t x, index_t y) const { return x > y; } +// }; +// using Seq = sequence<5, 2, 8, 1, 9>; +// using Result = typename sequence_sort::type; +// using Expected = sequence<9, 8, 5, 2, 1>; +// static_assert((std::is_same::value)); +// EXPECT_TRUE((std::is_same::value)); +// } + +TEST(SequenceSort, SortAlreadySorted) +{ + using Seq = sequence<1, 2, 3, 4, 5>; + using Result = typename sequence_sort>::type; + using Expected = sequence<1, 2, 3, 4, 5>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceSort, SortWithDuplicates) +{ + using Seq = sequence<3, 1, 4, 1, 5, 9, 2, 6, 5>; + using Result = typename sequence_sort>::type; + using Expected = sequence<1, 1, 2, 3, 4, 5, 5, 6, 9>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceSort, SortEmptySequence) +{ + using Seq = sequence<>; + using Result = typename sequence_sort>::type; + using Expected = sequence<>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceSort, SortSingleElement) +{ + using Seq = sequence<42>; + using Result = typename sequence_sort>::type; + using Expected = sequence<42>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test sequence_unique_sort +TEST(SequenceUniqueSort, UniqueSort) +{ + using Seq = sequence<3, 1, 4, 1, 5, 9, 2, 6, 5>; + using Result = typename sequence_unique_sort, equal>::type; + using Expected = sequence<1, 2, 3, 4, 5, 6, 9>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceUniqueSort, UniqueSortNoDuplicates) +{ + using Seq = sequence<5, 2, 8, 1, 9>; + using Result = typename sequence_unique_sort, equal>::type; + using Expected = sequence<1, 2, 5, 8, 9>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceUniqueSort, UniqueSortAllSame) +{ + using Seq = sequence<5, 5, 5, 5>; + using Result = typename sequence_unique_sort, equal>::type; + using Expected = sequence<5>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test is_valid_sequence_map +TEST(SequenceMap, ValidMap) +{ + using Map = sequence<0, 1, 2, 3>; + EXPECT_TRUE((is_valid_sequence_map::value)); +} + +TEST(SequenceMap, ValidMapPermuted) +{ + using Map = sequence<2, 0, 3, 1>; + EXPECT_TRUE((is_valid_sequence_map::value)); +} + +TEST(SequenceMap, InvalidMapDuplicate) +{ + using Map = sequence<0, 1, 1, 3>; + EXPECT_FALSE((is_valid_sequence_map::value)); +} + +TEST(SequenceMap, InvalidMapMissing) +{ + using Map = sequence<0, 1, 3, 4>; + EXPECT_FALSE((is_valid_sequence_map::value)); +} + +// Test sequence_map_inverse +// Note: sequence_map_inverse inverts a mapping where Map[i] = j means old position i maps to new +// position j The inverse gives us new position i came from old position inverse[i] +TEST(SequenceMapInverse, InverseMap) +{ + // Map = <2, 0, 3, 1> means: old[0]->new[2], old[1]->new[0], old[2]->new[3], old[3]->new[1] + // Inverse should be: new[0]<-old[1], new[1]<-old[3], new[2]<-old[0], new[3]<-old[2] + using Map = sequence<2, 0, 3, 1>; + using Result = typename sequence_map_inverse::type; + // Verify by checking that Map[Result[i]] == i for all i + EXPECT_EQ((Map::at(number{})>{}) == 0), true); + EXPECT_EQ((Map::at(number{})>{}) == 1), true); + EXPECT_EQ((Map::at(number{})>{}) == 2), true); + EXPECT_EQ((Map::at(number{})>{}) == 3), true); +} + +TEST(SequenceMapInverse, InverseIdentityMap) +{ + using Map = sequence<0, 1, 2, 3>; + using Result = typename sequence_map_inverse::type; + // Verify by checking that Map[Result[i]] == i for all i (same as the other test) + EXPECT_EQ((Map::at(number{})>{}) == 0), true); + EXPECT_EQ((Map::at(number{})>{}) == 1), true); + EXPECT_EQ((Map::at(number{})>{}) == 2), true); + EXPECT_EQ((Map::at(number{})>{}) == 3), true); +} + +// Test sequence operators +TEST(SequenceOperators, Equality) +{ + constexpr auto seq1 = sequence<1, 2, 3>{}; + constexpr auto seq2 = sequence<1, 2, 3>{}; + constexpr auto seq3 = sequence<1, 2, 4>{}; + EXPECT_TRUE(seq1 == seq2); + EXPECT_FALSE(seq1 == seq3); +} + +TEST(SequenceOperators, Addition) +{ + using Seq1 = sequence<1, 2, 3>; + using Seq2 = sequence<4, 5, 6>; + using Result = decltype(Seq1{} + Seq2{}); + using Expected = sequence<5, 7, 9>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, Subtraction) +{ + using Seq1 = sequence<10, 20, 30>; + using Seq2 = sequence<1, 2, 3>; + using Result = decltype(Seq1{} - Seq2{}); + using Expected = sequence<9, 18, 27>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, Multiplication) +{ + using Seq1 = sequence<2, 3, 4>; + using Seq2 = sequence<5, 6, 7>; + using Result = decltype(Seq1{} * Seq2{}); + using Expected = sequence<10, 18, 28>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, Division) +{ + using Seq1 = sequence<10, 20, 30>; + using Seq2 = sequence<2, 4, 5>; + using Result = decltype(Seq1{} / Seq2{}); + using Expected = sequence<5, 5, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, Modulo) +{ + using Seq1 = sequence<10, 20, 30>; + using Seq2 = sequence<3, 7, 8>; + using Result = decltype(Seq1{} % Seq2{}); + using Expected = sequence<1, 6, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, AdditionWithnumber) +{ + using Seq = sequence<1, 2, 3>; + using Result = decltype(Seq{} + number<10>{}); + using Expected = sequence<11, 12, 13>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, SubtractionWithnumber) +{ + using Seq = sequence<10, 20, 30>; + using Result = decltype(Seq{} - number<5>{}); + using Expected = sequence<5, 15, 25>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, MultiplicationWithnumber) +{ + using Seq = sequence<2, 3, 4>; + using Result = decltype(Seq{} * number<3>{}); + using Expected = sequence<6, 9, 12>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, DivisionWithnumber) +{ + using Seq = sequence<10, 20, 30>; + using Result = decltype(Seq{} / number<5>{}); + using Expected = sequence<2, 4, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, numberAddition) +{ + using Seq = sequence<1, 2, 3>; + using Result = decltype(number<10>{} + Seq{}); + using Expected = sequence<11, 12, 13>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceOperators, numberMultiplication) +{ + using Seq = sequence<2, 3, 4>; + using Result = decltype(number<3>{} * Seq{}); + using Expected = sequence<6, 9, 12>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test helper functions +TEST(SequenceHelpers, MergeSequences) +{ + using Seq1 = sequence<1, 2>; + using Seq2 = sequence<3, 4>; + using Seq3 = sequence<5, 6>; + using Result = decltype(merge_sequences(Seq1{}, Seq2{}, Seq3{})); + using Expected = sequence<1, 2, 3, 4, 5, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceHelpers, TransformSequencesSingle) +{ + auto double_it = [](auto x) { return 2 * x; }; + using Seq = sequence<1, 2, 3>; + using Result = decltype(transform_sequences(double_it, Seq{})); + using Expected = sequence<2, 4, 6>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceHelpers, TransformSequencesTwo) +{ + auto add = [](auto x, auto y) { return x + y; }; + using Seq1 = sequence<1, 2, 3>; + using Seq2 = sequence<4, 5, 6>; + using Result = decltype(transform_sequences(add, Seq1{}, Seq2{})); + using Expected = sequence<5, 7, 9>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceHelpers, TransformSequencesThree) +{ + auto add3 = [](auto x, auto y, auto z) { return x + y + z; }; + using Seq1 = sequence<1, 2, 3>; + using Seq2 = sequence<4, 5, 6>; + using Seq3 = sequence<7, 8, 9>; + using Result = decltype(transform_sequences(add3, Seq1{}, Seq2{}, Seq3{})); + using Expected = sequence<12, 15, 18>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceHelpers, ReduceOnSequence) +{ + auto add = [](auto x, auto y) { return x + y; }; + constexpr auto seq = sequence<1, 2, 3, 4, 5>{}; + constexpr auto result = reduce_on_sequence(seq, add, number<0>{}); + EXPECT_EQ(result, 15); +} + +TEST(SequenceHelpers, SequenceAnyOf) +{ + auto is_even = [](auto x) { return x % 2 == 0; }; + constexpr auto seq1 = sequence<1, 3, 5, 7>{}; + constexpr auto seq2 = sequence<1, 3, 4, 7>{}; + EXPECT_FALSE(sequence_any_of(seq1, is_even)); + EXPECT_TRUE(sequence_any_of(seq2, is_even)); +} + +TEST(SequenceHelpers, SequenceAllOf) +{ + auto is_positive = [](auto x) { return x > 0; }; + constexpr auto seq1 = sequence<1, 2, 3, 4>{}; + constexpr auto seq2 = sequence<1, -2, 3, 4>{}; + EXPECT_TRUE(sequence_all_of(seq1, is_positive)); + EXPECT_FALSE(sequence_all_of(seq2, is_positive)); +} + +// Test scan operations +TEST(SequenceScan, ReverseInclusiveScan) +{ + using Seq = sequence<1, 2, 3, 4>; + using Result = decltype(reverse_inclusive_scan_sequence(Seq{}, plus{}, number<0>{})); + using Expected = sequence<10, 9, 7, 4>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceScan, ReverseExclusiveScan) +{ + using Seq = sequence<1, 2, 3, 4>; + using Result = decltype(reverse_exclusive_scan_sequence(Seq{}, plus{}, number<0>{})); + using Expected = sequence<9, 7, 4, 0>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceScan, InclusiveScan) +{ + using Seq = sequence<1, 2, 3, 4>; + using Result = decltype(inclusive_scan_sequence(Seq{}, plus{}, number<0>{})); + using Expected = sequence<1, 3, 6, 10>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test pick and modify operations +TEST(SequencePickModify, PickElementsByIds) +{ + using Seq = sequence<10, 20, 30, 40, 50>; + using Ids = sequence<0, 2, 4>; + using Result = decltype(pick_sequence_elements_by_ids(Seq{}, Ids{})); + using Expected = sequence<10, 30, 50>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequencePickModify, PickElementsByMask) +{ + using Seq = sequence<10, 20, 30, 40, 50>; + using Mask = sequence<1, 0, 1, 0, 1>; + using Result = decltype(pick_sequence_elements_by_mask(Seq{}, Mask{})); + using Expected = sequence<10, 30, 50>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequencePickModify, ModifyElementsByIds) +{ + using Seq = sequence<10, 20, 30, 40, 50>; + using Values = sequence<99, 88>; + using Ids = sequence<1, 3>; + using Result = decltype(modify_sequence_elements_by_ids(Seq{}, Values{}, Ids{})); + using Expected = sequence<10, 99, 30, 88, 50>; + EXPECT_TRUE((std::is_same::value)); +} + +// Test sequence_reduce +TEST(SequenceReduce, ReduceTwoSequences) +{ + using Seq1 = sequence<1, 2, 3>; + using Seq2 = sequence<4, 5, 6>; + using Result = typename sequence_reduce, Seq1, Seq2>::type; + using Expected = sequence<5, 7, 9>; + EXPECT_TRUE((std::is_same::value)); +} + +TEST(SequenceReduce, ReduceMultipleSequences) +{ + using Seq1 = sequence<1, 2>; + using Seq2 = sequence<3, 4>; + using Seq3 = sequence<5, 6>; + using Result = typename sequence_reduce, Seq1, Seq2, Seq3>::type; + using Expected = sequence<9, 12>; + EXPECT_TRUE((std::is_same::value)); +}