Add unit tests for template optimization helpers

test/util/CMakeLists.txt

@@ -5,3 +5,18 @@ add_gtest_executable(unit_sequence unit_sequence.cpp)
 if(result EQUAL 0)
   target_link_libraries(unit_sequence PRIVATE utility)
 endif()
+
+add_gtest_executable(unit_sequence_helper unit_sequence_helper.cpp)
+if(result EQUAL 0)
+  target_link_libraries(unit_sequence_helper PRIVATE utility)
+endif()
+
+add_gtest_executable(unit_container_helper unit_container_helper.cpp)
+if(result EQUAL 0)
+  target_link_libraries(unit_container_helper PRIVATE utility)
+endif()
+
+add_gtest_executable(unit_tensor_descriptor_helper unit_tensor_descriptor_helper.cpp)
+if(result EQUAL 0)
+  target_link_libraries(unit_tensor_descriptor_helper PRIVATE utility)
+endif()

test/util/unit_container_helper.cpp

@@ -0,0 +1,145 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include <gtest/gtest.h>
+#include "ck/utility/container_helper.hpp"
+#include "ck/utility/tuple_helper.hpp"
+
+using namespace ck;
+
+// Test container_concat with tuples
+TEST(ContainerConcat, ConcatTwoTuples)
+{
+    constexpr auto t1     = make_tuple(Number<7>{}, Number<11>{});
+    constexpr auto t2     = make_tuple(Number<13>{}, Number<17>{});
+    constexpr auto result = container_concat(t1, t2);
+
+    EXPECT_EQ(result.Size(), 4);
+    EXPECT_EQ(result[Number<0>{}], 7);
+    EXPECT_EQ(result[Number<1>{}], 11);
+    EXPECT_EQ(result[Number<2>{}], 13);
+    EXPECT_EQ(result[Number<3>{}], 17);
+}
+
+TEST(ContainerConcat, ConcatThreeTuples)
+{
+    constexpr auto t1     = make_tuple(Number<19>{});
+    constexpr auto t2     = make_tuple(Number<23>{}, Number<29>{});
+    constexpr auto t3     = make_tuple(Number<31>{});
+    constexpr auto result = container_concat(t1, t2, t3);
+
+    EXPECT_EQ(result.Size(), 4);
+    EXPECT_EQ(result[Number<0>{}], 19);
+    EXPECT_EQ(result[Number<1>{}], 23);
+    EXPECT_EQ(result[Number<2>{}], 29);
+    EXPECT_EQ(result[Number<3>{}], 31);
+}
+
+TEST(ContainerConcat, ConcatWithEmptyTuple)
+{
+    constexpr auto t1     = make_tuple(Number<37>{}, Number<41>{});
+    constexpr auto empty  = make_tuple();
+    constexpr auto result = container_concat(t1, empty);
+
+    EXPECT_EQ(result.Size(), 2);
+    EXPECT_EQ(result[Number<0>{}], 37);
+    EXPECT_EQ(result[Number<1>{}], 41);
+}
+
+TEST(ContainerConcat, ConcatSingleTuple)
+{
+    constexpr auto t1     = make_tuple(Number<43>{}, Number<47>{}, Number<53>{});
+    constexpr auto result = container_concat(t1);
+
+    EXPECT_EQ(result.Size(), 3);
+    EXPECT_EQ(result[Number<0>{}], 43);
+    EXPECT_EQ(result[Number<1>{}], 47);
+    EXPECT_EQ(result[Number<2>{}], 53);
+}
+
+// Test container_concat with arrays
+TEST(ContainerConcat, ConcatTwoArrays)
+{
+    constexpr auto a1     = make_array(59, 61);
+    constexpr auto a2     = make_array(67, 71);
+    constexpr auto result = container_concat(a1, a2);
+
+    EXPECT_EQ(result.Size(), 4);
+    EXPECT_EQ(result[Number<0>{}], 59);
+    EXPECT_EQ(result[Number<1>{}], 61);
+    EXPECT_EQ(result[Number<2>{}], 67);
+    EXPECT_EQ(result[Number<3>{}], 71);
+}
+
+// Test make_uniform_tuple
+TEST(MakeUniformTuple, Size3)
+{
+    constexpr auto result = make_uniform_tuple<3>(Number<73>{});
+
+    EXPECT_EQ(result.Size(), 3);
+    EXPECT_EQ(result[Number<0>{}], 73);
+    EXPECT_EQ(result[Number<1>{}], 73);
+    EXPECT_EQ(result[Number<2>{}], 73);
+}
+
+TEST(MakeUniformTuple, Size1)
+{
+    constexpr auto result = make_uniform_tuple<1>(Number<79>{});
+
+    EXPECT_EQ(result.Size(), 1);
+    EXPECT_EQ(result[Number<0>{}], 79);
+}
+
+TEST(MakeUniformTuple, Size0)
+{
+    constexpr auto result = make_uniform_tuple<0>(Number<83>{});
+
+    EXPECT_EQ(result.Size(), 0);
+}
+
+TEST(MakeUniformTuple, Size5)
+{
+    constexpr auto result = make_uniform_tuple<5>(Number<89>{});
+
+    EXPECT_EQ(result.Size(), 5);
+    EXPECT_EQ(result[Number<0>{}], 89);
+    EXPECT_EQ(result[Number<1>{}], 89);
+    EXPECT_EQ(result[Number<2>{}], 89);
+    EXPECT_EQ(result[Number<3>{}], 89);
+    EXPECT_EQ(result[Number<4>{}], 89);
+}
+
+// Test make_tuple_functor (used internally by container_concat)
+TEST(MakeTupleFunctor, CreatesTuple)
+{
+    make_tuple_functor functor;
+    auto result = functor(Number<97>{}, Number<101>{}, Number<103>{});
+
+    EXPECT_EQ(result.Size(), 3);
+    EXPECT_EQ(result[Number<0>{}], 97);
+    EXPECT_EQ(result[Number<1>{}], 101);
+    EXPECT_EQ(result[Number<2>{}], 103);
+}
+
+// Test container_push_front and container_push_back
+TEST(ContainerPush, PushFront)
+{
+    constexpr auto t      = make_tuple(Number<109>{}, Number<113>{});
+    constexpr auto result = container_push_front(t, Number<107>{});
+
+    EXPECT_EQ(result.Size(), 3);
+    EXPECT_EQ(result[Number<0>{}], 107);
+    EXPECT_EQ(result[Number<1>{}], 109);
+    EXPECT_EQ(result[Number<2>{}], 113);
+}
+
+TEST(ContainerPush, PushBack)
+{
+    constexpr auto t      = make_tuple(Number<127>{}, Number<131>{});
+    constexpr auto result = container_push_back(t, Number<137>{});
+
+    EXPECT_EQ(result.Size(), 3);
+    EXPECT_EQ(result[Number<0>{}], 127);
+    EXPECT_EQ(result[Number<1>{}], 131);
+    EXPECT_EQ(result[Number<2>{}], 137);
+}

test/util/unit_sequence.cpp

@@ -185,6 +185,54 @@ TEST(Sequence, ReorderGivenOld2New)
     EXPECT_TRUE((is_same<Result, Expected>::value));
 }
 
+// Test sequence_gen with custom functor
+TEST(SequenceGen, SequenceGenWithDoubleFunctor)
+{
+    struct DoubleFunctor
+    {
+        __host__ __device__ constexpr index_t operator()(index_t i) const { return i * 2; }
+    };
+    using Result   = typename sequence_gen<5, DoubleFunctor>::type;
+    using Expected = Sequence<0, 2, 4, 6, 8>;
+    EXPECT_TRUE((is_same<Result, Expected>::value));
+}
+
+TEST(SequenceGen, SequenceGenWithSquareFunctor)
+{
+    struct SquareFunctor
+    {
+        __host__ __device__ constexpr index_t operator()(index_t i) const { return i * i; }
+    };
+    using Result   = typename sequence_gen<5, SquareFunctor>::type;
+    using Expected = Sequence<0, 1, 4, 9, 16>;
+    EXPECT_TRUE((is_same<Result, Expected>::value));
+}
+
+TEST(SequenceGen, SequenceGenZeroSize)
+{
+    struct IdentityFunctor
+    {
+        __host__ __device__ constexpr index_t operator()(index_t i) const { return i; }
+    };
+    using Result   = typename sequence_gen<0, IdentityFunctor>::type;
+    using Expected = Sequence<>;
+    EXPECT_TRUE((is_same<Result, Expected>::value));
+    // Also verify non-zero size works with identity
+    using Result5 = typename sequence_gen<5, IdentityFunctor>::type;
+    EXPECT_TRUE((is_same<Result5, Sequence<0, 1, 2, 3, 4>>::value));
+}
+
+TEST(SequenceGen, SequenceGenSingleElement)
+{
+    struct ConstantFunctor
+    {
+        __host__ __device__ constexpr index_t operator()(index_t) const { return 42; }
+    };
+    using Result   = typename sequence_gen<1, ConstantFunctor>::type;
+    using Expected = Sequence<42>;
+    EXPECT_TRUE((is_same<Result, Expected>::value));
+}
+
 // Test arithmetic_sequence_gen
 TEST(SequenceGen, ArithmeticSequence)
 {

test/util/unit_sequence_helper.cpp

@@ -0,0 +1,153 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include <gtest/gtest.h>
+#include "ck/utility/sequence.hpp"
+#include "ck/utility/sequence_helper.hpp"
+#include "ck/utility/tuple_helper.hpp"
+
+using namespace ck;
+
+// Test generate_identity_sequences
+TEST(GenerateIdentitySequences, Size5)
+{
+    constexpr auto result = generate_identity_sequences<5>();
+    // Should produce Tuple<Sequence<0>, Sequence<1>, Sequence<2>, Sequence<3>, Sequence<4>>
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<0>{}])>, Sequence<0>>::value));
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<1>{}])>, Sequence<1>>::value));
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<2>{}])>, Sequence<2>>::value));
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<3>{}])>, Sequence<3>>::value));
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<4>{}])>, Sequence<4>>::value));
+}
+
+TEST(GenerateIdentitySequences, Size1)
+{
+    constexpr auto result = generate_identity_sequences<1>();
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<0>{}])>, Sequence<0>>::value));
+    EXPECT_EQ(result.Size(), 1);
+}
+
+TEST(GenerateIdentitySequences, Size0)
+{
+    constexpr auto result = generate_identity_sequences<0>();
+    EXPECT_EQ(result.Size(), 0);
+}
+
+TEST(GenerateIdentitySequences, WithNumber)
+{
+    constexpr auto result = generate_identity_sequences(Number<3>{});
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<0>{}])>, Sequence<0>>::value));
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<1>{}])>, Sequence<1>>::value));
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result[Number<2>{}])>, Sequence<2>>::value));
+    EXPECT_EQ(result.Size(), 3);
+}
+
+// Test sequence_find_value
+TEST(SequenceFindValue, FindExistingElement)
+{
+    constexpr auto result = sequence_find_value<17>(Sequence<5, 11, 17, 23, 29>{});
+    EXPECT_EQ(result, 2); // 17 is at index 2
+}
+
+TEST(SequenceFindValue, FindFirstElement)
+{
+    constexpr auto result = sequence_find_value<7>(Sequence<7, 13, 19, 31>{});
+    EXPECT_EQ(result, 0);
+}
+
+TEST(SequenceFindValue, FindLastElement)
+{
+    constexpr auto result = sequence_find_value<41>(Sequence<3, 11, 23, 41>{});
+    EXPECT_EQ(result, 3);
+}
+
+TEST(SequenceFindValue, ElementNotFound)
+{
+    constexpr auto result = sequence_find_value<50>(Sequence<2, 8, 14, 26>{});
+    EXPECT_EQ(result, -1);
+}
+
+TEST(SequenceFindValue, EmptySequence)
+{
+    constexpr auto result = sequence_find_value<1>(Sequence<>{});
+    EXPECT_EQ(result, -1);
+}
+
+// Test find_in_tuple_of_sequences
+TEST(FindInTupleOfSequences, FindInFirstSequence)
+{
+    constexpr auto tuple_of_seqs =
+        make_tuple(Sequence<5, 11>{}, Sequence<17, 23>{}, Sequence<29, 37>{});
+    constexpr auto result = find_in_tuple_of_sequences<11>(tuple_of_seqs);
+    EXPECT_EQ(result.itran, 0);   // Found in first sequence (index 0)
+    EXPECT_EQ(result.idim_up, 1); // At position 1 within that sequence
+    EXPECT_TRUE(result.found);
+}
+
+TEST(FindInTupleOfSequences, FindInMiddleSequence)
+{
+    constexpr auto tuple_of_seqs =
+        make_tuple(Sequence<2, 4, 6>{}, Sequence<8, 10>{}, Sequence<12>{});
+    constexpr auto result = find_in_tuple_of_sequences<10>(tuple_of_seqs);
+    EXPECT_EQ(result.itran, 1);   // Found in second sequence (index 1)
+    EXPECT_EQ(result.idim_up, 1); // At position 1 within that sequence
+    EXPECT_TRUE(result.found);
+}
+
+TEST(FindInTupleOfSequences, FindInLastSequence)
+{
+    constexpr auto tuple_of_seqs = make_tuple(Sequence<3>{}, Sequence<7>{}, Sequence<13, 19, 31>{});
+    constexpr auto result        = find_in_tuple_of_sequences<31>(tuple_of_seqs);
+    EXPECT_EQ(result.itran, 2);   // Found in third sequence (index 2)
+    EXPECT_EQ(result.idim_up, 2); // At position 2 within that sequence
+    EXPECT_TRUE(result.found);
+}
+
+TEST(FindInTupleOfSequences, NotFound)
+{
+    constexpr auto tuple_of_seqs = make_tuple(Sequence<1, 3>{}, Sequence<5, 7, 9>{});
+    constexpr auto result        = find_in_tuple_of_sequences<100>(tuple_of_seqs);
+    EXPECT_FALSE(result.found);
+}
+
+TEST(FindInTupleOfSequences, EmptyTuple)
+{
+    constexpr auto tuple_of_seqs = make_tuple();
+    constexpr auto result        = find_in_tuple_of_sequences<1>(tuple_of_seqs);
+    EXPECT_FALSE(result.found);
+}
+
+TEST(FindInTupleOfSequences, SingleSequence)
+{
+    constexpr auto tuple_of_seqs = make_tuple(Sequence<41, 43, 47, 53>{});
+    constexpr auto result        = find_in_tuple_of_sequences<47>(tuple_of_seqs);
+    EXPECT_EQ(result.itran, 0);
+    EXPECT_EQ(result.idim_up, 2);
+    EXPECT_TRUE(result.found);
+}
+
+// Test unpack_and_merge_sequences
+TEST(UnpackAndMergeSequences, MergeMultipleSequences)
+{
+    constexpr auto tuple_of_seqs =
+        make_tuple(Sequence<7, 11>{}, Sequence<13>{}, Sequence<17, 19, 23>{});
+    constexpr auto result = unpack_and_merge_sequences(tuple_of_seqs);
+    using Expected        = Sequence<7, 11, 13, 17, 19, 23>;
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result)>, Expected>::value));
+}
+
+TEST(UnpackAndMergeSequences, SingleSequence)
+{
+    constexpr auto tuple_of_seqs = make_tuple(Sequence<29, 31, 37, 41>{});
+    constexpr auto result        = unpack_and_merge_sequences(tuple_of_seqs);
+    using Expected               = Sequence<29, 31, 37, 41>;
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result)>, Expected>::value));
+}
+
+TEST(UnpackAndMergeSequences, TwoSequences)
+{
+    constexpr auto tuple_of_seqs = make_tuple(Sequence<2, 3, 5>{}, Sequence<7, 11>{});
+    constexpr auto result        = unpack_and_merge_sequences(tuple_of_seqs);
+    using Expected               = Sequence<2, 3, 5, 7, 11>;
+    EXPECT_TRUE((is_same<remove_cvref_t<decltype(result)>, Expected>::value));
+}