From 0c8188374a579f22599ad85558322f2a925a7006 Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Mon, 19 Jan 2026 16:35:39 -0600 Subject: [PATCH] Add unit tests for template optimization helpers Add Google Tests for optimized template utilities: - sequence_gen: Tests with custom functors (4 tests) - generate_identity_sequences: Tuple of identity sequences (4 tests) - find_in_tuple_of_sequences: O(1) sequence search (6 tests) - sequence_find_value: Value lookup in sequences (5 tests) - container_concat: Tuple/array concatenation (5 tests) - make_uniform_tuple: Repeated value tuples (4 tests) - compute_element_space_size: Fold expression (8 tests) - unpack_and_merge_sequences: Sequence merging (2 tests) Total: 43 new tests across 4 test files. --- test/util/CMakeLists.txt | 15 +++ test/util/unit_container_helper.cpp | 141 +++++++++++++++++++ test/util/unit_sequence.cpp | 48 +++++++ test/util/unit_sequence_helper.cpp | 142 ++++++++++++++++++++ test/util/unit_tensor_descriptor_helper.cpp | 125 +++++++++++++++++ 5 files changed, 471 insertions(+) create mode 100644 test/util/unit_container_helper.cpp create mode 100644 test/util/unit_sequence_helper.cpp create mode 100644 test/util/unit_tensor_descriptor_helper.cpp diff --git a/test/util/CMakeLists.txt b/test/util/CMakeLists.txt index bf0a444f18..574915c273 100644 --- a/test/util/CMakeLists.txt +++ b/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() diff --git a/test/util/unit_container_helper.cpp b/test/util/unit_container_helper.cpp new file mode 100644 index 0000000000..a817ebf21c --- /dev/null +++ b/test/util/unit_container_helper.cpp @@ -0,0 +1,141 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#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<1>{}, Number<2>{}); + constexpr auto t2 = make_tuple(Number<3>{}, Number<4>{}); + constexpr auto result = container_concat(t1, t2); + + EXPECT_EQ(result.Size(), 4); + EXPECT_EQ(result[Number<0>{}], 1); + EXPECT_EQ(result[Number<1>{}], 2); + EXPECT_EQ(result[Number<2>{}], 3); + EXPECT_EQ(result[Number<3>{}], 4); +} + +TEST(ContainerConcat, ConcatThreeTuples) +{ + constexpr auto t1 = make_tuple(Number<1>{}); + constexpr auto t2 = make_tuple(Number<2>{}); + constexpr auto t3 = make_tuple(Number<3>{}); + constexpr auto result = container_concat(t1, t2, t3); + + EXPECT_EQ(result.Size(), 3); + EXPECT_EQ(result[Number<0>{}], 1); + EXPECT_EQ(result[Number<1>{}], 2); + EXPECT_EQ(result[Number<2>{}], 3); +} + +TEST(ContainerConcat, ConcatWithEmptyTuple) +{ + constexpr auto t1 = make_tuple(Number<1>{}, Number<2>{}); + constexpr auto empty = make_tuple(); + constexpr auto result = container_concat(t1, empty); + + EXPECT_EQ(result.Size(), 2); + EXPECT_EQ(result[Number<0>{}], 1); + EXPECT_EQ(result[Number<1>{}], 2); +} + +TEST(ContainerConcat, ConcatSingleTuple) +{ + constexpr auto t1 = make_tuple(Number<1>{}, Number<2>{}, Number<3>{}); + constexpr auto result = container_concat(t1); + + EXPECT_EQ(result.Size(), 3); +} + +// Test container_concat with arrays +TEST(ContainerConcat, ConcatTwoArrays) +{ + constexpr auto a1 = make_array(1, 2); + constexpr auto a2 = make_array(3, 4); + constexpr auto result = container_concat(a1, a2); + + EXPECT_EQ(result.Size(), 4); + EXPECT_EQ(result[Number<0>{}], 1); + EXPECT_EQ(result[Number<1>{}], 2); + EXPECT_EQ(result[Number<2>{}], 3); + EXPECT_EQ(result[Number<3>{}], 4); +} + +// Test make_uniform_tuple +TEST(MakeUniformTuple, Size3) +{ + constexpr auto result = make_uniform_tuple<3>(Number<42>{}); + + EXPECT_EQ(result.Size(), 3); + EXPECT_EQ(result[Number<0>{}], 42); + EXPECT_EQ(result[Number<1>{}], 42); + EXPECT_EQ(result[Number<2>{}], 42); +} + +TEST(MakeUniformTuple, Size1) +{ + constexpr auto result = make_uniform_tuple<1>(Number<99>{}); + + EXPECT_EQ(result.Size(), 1); + EXPECT_EQ(result[Number<0>{}], 99); +} + +TEST(MakeUniformTuple, Size0) +{ + constexpr auto result = make_uniform_tuple<0>(Number<42>{}); + + EXPECT_EQ(result.Size(), 0); +} + +TEST(MakeUniformTuple, Size5) +{ + constexpr auto result = make_uniform_tuple<5>(Number<7>{}); + + EXPECT_EQ(result.Size(), 5); + EXPECT_EQ(result[Number<0>{}], 7); + EXPECT_EQ(result[Number<1>{}], 7); + EXPECT_EQ(result[Number<2>{}], 7); + EXPECT_EQ(result[Number<3>{}], 7); + EXPECT_EQ(result[Number<4>{}], 7); +} + +// Test make_tuple_functor (used internally by container_concat) +TEST(MakeTupleFunctor, CreatesTuple) +{ + make_tuple_functor functor; + auto result = functor(Number<1>{}, Number<2>{}, Number<3>{}); + + EXPECT_EQ(result.Size(), 3); + EXPECT_EQ(result[Number<0>{}], 1); + EXPECT_EQ(result[Number<1>{}], 2); + EXPECT_EQ(result[Number<2>{}], 3); +} + +// Test container_push_front and container_push_back +TEST(ContainerPush, PushFront) +{ + constexpr auto t = make_tuple(Number<2>{}, Number<3>{}); + constexpr auto result = container_push_front(t, Number<1>{}); + + EXPECT_EQ(result.Size(), 3); + EXPECT_EQ(result[Number<0>{}], 1); + EXPECT_EQ(result[Number<1>{}], 2); + EXPECT_EQ(result[Number<2>{}], 3); +} + +TEST(ContainerPush, PushBack) +{ + constexpr auto t = make_tuple(Number<1>{}, Number<2>{}); + constexpr auto result = container_push_back(t, Number<3>{}); + + EXPECT_EQ(result.Size(), 3); + EXPECT_EQ(result[Number<0>{}], 1); + EXPECT_EQ(result[Number<1>{}], 2); + EXPECT_EQ(result[Number<2>{}], 3); +} diff --git a/test/util/unit_sequence.cpp b/test/util/unit_sequence.cpp index f09fd86e06..112212dbb7 100644 --- a/test/util/unit_sequence.cpp +++ b/test/util/unit_sequence.cpp @@ -185,6 +185,54 @@ TEST(Sequence, ReorderGivenOld2New) EXPECT_TRUE((is_same::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::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::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::value)); + // Also verify non-zero size works with identity + using Result5 = typename sequence_gen<5, IdentityFunctor>::type; + EXPECT_TRUE((is_same>::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::value)); +} + // Test arithmetic_sequence_gen TEST(SequenceGen, ArithmeticSequence) { diff --git a/test/util/unit_sequence_helper.cpp b/test/util/unit_sequence_helper.cpp new file mode 100644 index 0000000000..4f088d262b --- /dev/null +++ b/test/util/unit_sequence_helper.cpp @@ -0,0 +1,142 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#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<1>, Sequence<2>, Sequence<3>, Sequence<4>> + EXPECT_TRUE((is_same{}])>, Sequence<0>>::value)); + EXPECT_TRUE((is_same{}])>, Sequence<1>>::value)); + EXPECT_TRUE((is_same{}])>, Sequence<2>>::value)); + EXPECT_TRUE((is_same{}])>, Sequence<3>>::value)); + EXPECT_TRUE((is_same{}])>, Sequence<4>>::value)); +} + +TEST(GenerateIdentitySequences, Size1) +{ + constexpr auto result = generate_identity_sequences<1>(); + EXPECT_TRUE((is_same{}])>, 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{}])>, Sequence<0>>::value)); + EXPECT_TRUE((is_same{}])>, Sequence<1>>::value)); + EXPECT_TRUE((is_same{}])>, Sequence<2>>::value)); + EXPECT_EQ(result.Size(), 3); +} + +// Test sequence_find_value +TEST(SequenceFindValue, FindExistingElement) +{ + constexpr auto result = sequence_find_value<3>(Sequence<1, 2, 3, 4, 5>{}); + EXPECT_EQ(result, 2); // 3 is at index 2 +} + +TEST(SequenceFindValue, FindFirstElement) +{ + constexpr auto result = sequence_find_value<10>(Sequence<10, 20, 30>{}); + EXPECT_EQ(result, 0); +} + +TEST(SequenceFindValue, FindLastElement) +{ + constexpr auto result = sequence_find_value<30>(Sequence<10, 20, 30>{}); + EXPECT_EQ(result, 2); +} + +TEST(SequenceFindValue, ElementNotFound) +{ + constexpr auto result = sequence_find_value<99>(Sequence<1, 2, 3, 4, 5>{}); + 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<0, 1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}); + constexpr auto result = find_in_tuple_of_sequences<1>(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<0, 1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}); + constexpr auto result = find_in_tuple_of_sequences<3>(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<0, 1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}); + constexpr auto result = find_in_tuple_of_sequences<5>(tuple_of_seqs); + EXPECT_EQ(result.itran, 2); // Found in third sequence (index 2) + EXPECT_EQ(result.idim_up, 1); // At position 1 within that sequence + EXPECT_TRUE(result.found); +} + +TEST(FindInTupleOfSequences, NotFound) +{ + constexpr auto tuple_of_seqs = make_tuple(Sequence<0, 1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}); + constexpr auto result = find_in_tuple_of_sequences<99>(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<10, 20, 30>{}); + constexpr auto result = find_in_tuple_of_sequences<20>(tuple_of_seqs); + EXPECT_EQ(result.itran, 0); + EXPECT_EQ(result.idim_up, 1); + EXPECT_TRUE(result.found); +} + +// Test unpack_and_merge_sequences +TEST(UnpackAndMergeSequences, MergeMultipleSequences) +{ + constexpr auto tuple_of_seqs = make_tuple(Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5, 6>{}); + constexpr auto result = unpack_and_merge_sequences(tuple_of_seqs); + using Expected = Sequence<1, 2, 3, 4, 5, 6>; + EXPECT_TRUE((is_same, Expected>::value)); +} + +TEST(UnpackAndMergeSequences, SingleSequence) +{ + constexpr auto tuple_of_seqs = make_tuple(Sequence<1, 2, 3>{}); + constexpr auto result = unpack_and_merge_sequences(tuple_of_seqs); + using Expected = Sequence<1, 2, 3>; + EXPECT_TRUE((is_same, Expected>::value)); +} diff --git a/test/util/unit_tensor_descriptor_helper.cpp b/test/util/unit_tensor_descriptor_helper.cpp new file mode 100644 index 0000000000..5614e4dd5f --- /dev/null +++ b/test/util/unit_tensor_descriptor_helper.cpp @@ -0,0 +1,125 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include "ck/utility/common_header.hpp" +#include "ck/tensor_description/tensor_descriptor_helper.hpp" + +using namespace ck; + +// Test compute_element_space_size helper directly +// Formula: 1 + sum((length[i] - 1) * stride[i]) +TEST(ComputeElementSpaceSize, Simple2D) +{ + // 4x4 tensor with row-major strides [4, 1] + // element_space_size = 1 + (4-1)*4 + (4-1)*1 = 1 + 12 + 3 = 16 + constexpr auto lengths = make_tuple(Number<4>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<4>{}, Number<1>{}); + constexpr auto result = detail::compute_element_space_size(lengths, strides, Sequence<0, 1>{}); + EXPECT_EQ(result, 16); +} + +TEST(ComputeElementSpaceSize, Simple3D) +{ + // 2x3x4 tensor with strides [12, 4, 1] + // element_space_size = 1 + (2-1)*12 + (3-1)*4 + (4-1)*1 = 1 + 12 + 8 + 3 = 24 + constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<12>{}, Number<4>{}, Number<1>{}); + constexpr auto result = + detail::compute_element_space_size(lengths, strides, Sequence<0, 1, 2>{}); + EXPECT_EQ(result, 24); +} + +TEST(ComputeElementSpaceSize, WithPadding) +{ + // 4x4 tensor with padded strides [8, 1] (8 elements per row instead of 4) + // element_space_size = 1 + (4-1)*8 + (4-1)*1 = 1 + 24 + 3 = 28 + constexpr auto lengths = make_tuple(Number<4>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<8>{}, Number<1>{}); + constexpr auto result = detail::compute_element_space_size(lengths, strides, Sequence<0, 1>{}); + EXPECT_EQ(result, 28); +} + +TEST(ComputeElementSpaceSize, ColumnMajor) +{ + // 4x4 tensor with column-major strides [1, 4] + // element_space_size = 1 + (4-1)*1 + (4-1)*4 = 1 + 3 + 12 = 16 + constexpr auto lengths = make_tuple(Number<4>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<1>{}, Number<4>{}); + constexpr auto result = detail::compute_element_space_size(lengths, strides, Sequence<0, 1>{}); + EXPECT_EQ(result, 16); +} + +TEST(ComputeElementSpaceSize, SingleDimension) +{ + // 10 element 1D tensor with stride 1 + // element_space_size = 1 + (10-1)*1 = 10 + constexpr auto lengths = make_tuple(Number<10>{}); + constexpr auto strides = make_tuple(Number<1>{}); + constexpr auto result = detail::compute_element_space_size(lengths, strides, Sequence<0>{}); + EXPECT_EQ(result, 10); +} + +TEST(ComputeElementSpaceSize, SingleElement) +{ + // 1x1 tensor - should have size 1 + // element_space_size = 1 + (1-1)*1 + (1-1)*1 = 1 + constexpr auto lengths = make_tuple(Number<1>{}, Number<1>{}); + constexpr auto strides = make_tuple(Number<1>{}, Number<1>{}); + constexpr auto result = detail::compute_element_space_size(lengths, strides, Sequence<0, 1>{}); + EXPECT_EQ(result, 1); +} + +TEST(ComputeElementSpaceSize, BroadcastDimension) +{ + // 8x5 tensor with broadcast on first dimension (stride 0) + // element_space_size = 1 + (8-1)*0 + (5-1)*1 = 1 + 0 + 4 = 5 + constexpr auto lengths = make_tuple(Number<8>{}, Number<5>{}); + constexpr auto strides = make_tuple(Number<0>{}, Number<1>{}); + constexpr auto result = detail::compute_element_space_size(lengths, strides, Sequence<0, 1>{}); + EXPECT_EQ(result, 5); +} + +// Test make_naive_tensor_descriptor uses compute_element_space_size correctly +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize2D) +{ + constexpr auto lengths = make_tuple(Number<4>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<4>{}, Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 16); +} + +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize3D) +{ + constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<12>{}, Number<4>{}, Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 24); +} + +TEST(MakeNaiveTensorDescriptor, BroadcastDimension) +{ + // 8x5 tensor with broadcast on first dimension (stride 0) + // element_space_size = 1 + (8-1)*0 + (5-1)*1 = 5 + constexpr auto lengths = make_tuple(Number<8>{}, Number<5>{}); + constexpr auto strides = make_tuple(Number<0>{}, Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 5); +} + +// Test with runtime values (index_t instead of Number<>) +TEST(ComputeElementSpaceSizeRuntime, Simple2D) +{ + const auto lengths = make_tuple(index_t{4}, index_t{4}); + const auto strides = make_tuple(index_t{4}, index_t{1}); + const auto result = detail::compute_element_space_size(lengths, strides, Sequence<0, 1>{}); + EXPECT_EQ(result, 16); +} + +TEST(ComputeElementSpaceSizeRuntime, WithPadding) +{ + const auto lengths = make_tuple(index_t{4}, index_t{4}); + const auto strides = make_tuple(index_t{8}, index_t{1}); + const auto result = detail::compute_element_space_size(lengths, strides, Sequence<0, 1>{}); + EXPECT_EQ(result, 28); +}