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 1/8] 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 bf0a444f18b..574915c2731 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 00000000000..a817ebf21c5 --- /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 f09fd86e063..112212dbb72 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 00000000000..4f088d262b1 --- /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 00000000000..5614e4dd5ff --- /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); +} From b60d14ba898f3fbb92247761c398d75c87d2a3bc Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Mon, 19 Jan 2026 23:31:13 -0600 Subject: [PATCH 2/8] Address review feedback on tensor descriptor helper tests - Remove tests of implementation details (detail::compute_element_space_size) - Use public API (make_naive_tensor_descriptor) for all tests - Avoid square/cube shapes that could hide row/column major bugs - Use prime numbers for padding tests to catch index calculation errors - Add two padding test cases: arbitrary offsets and stride slice --- test/util/unit_tensor_descriptor_helper.cpp | 131 ++++++++------------ 1 file changed, 49 insertions(+), 82 deletions(-) diff --git a/test/util/unit_tensor_descriptor_helper.cpp b/test/util/unit_tensor_descriptor_helper.cpp index 5614e4dd5ff..d0d67b38e31 100644 --- a/test/util/unit_tensor_descriptor_helper.cpp +++ b/test/util/unit_tensor_descriptor_helper.cpp @@ -7,119 +7,86 @@ 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 make_naive_tensor_descriptor (public API) +// Formula for element_space_size: 1 + sum((length[i] - 1) * stride[i]) -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) +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize2D) { - // 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); + // 5x4 tensor with row-major strides [4, 1] + // element_space_size = 1 + (5-1)*4 + (4-1)*1 = 1 + 16 + 3 = 20 + constexpr auto lengths = make_tuple(Number<5>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<4>{}, Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 20); } -TEST(ComputeElementSpaceSize, SingleElement) +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize3D) { - // 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); + // 3x8x11 tensor with strides [88, 11, 1] + // element_space_size = 1 + (3-1)*88 + (8-1)*11 + (11-1)*1 = 1 + 176 + 77 + 10 = 264 + constexpr auto lengths = make_tuple(Number<3>{}, Number<8>{}, Number<11>{}); + constexpr auto strides = make_tuple(Number<88>{}, Number<11>{}, Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 264); } -TEST(ComputeElementSpaceSize, BroadcastDimension) +TEST(MakeNaiveTensorDescriptor, 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); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 5); } -// Test make_naive_tensor_descriptor uses compute_element_space_size correctly -TEST(MakeNaiveTensorDescriptor, ElementSpaceSize2D) +TEST(MakeNaiveTensorDescriptor, WithPaddingArbitrary) { - constexpr auto lengths = make_tuple(Number<4>{}, Number<4>{}); - constexpr auto strides = make_tuple(Number<4>{}, Number<1>{}); + // 11x7x3 tensor with arbitrary strides [2, 97, 23] (prime numbers, no common factors) + // This tests padding where offsets don't correspond to any packed array + // element_space_size = 1 + (11-1)*2 + (7-1)*97 + (3-1)*23 = 1 + 20 + 582 + 46 = 649 + constexpr auto lengths = make_tuple(Number<11>{}, Number<7>{}, Number<3>{}); + constexpr auto strides = make_tuple(Number<2>{}, Number<97>{}, Number<23>{}); constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); - EXPECT_EQ(desc.GetElementSpaceSize(), 16); + EXPECT_EQ(desc.GetElementSpaceSize(), 649); } -TEST(MakeNaiveTensorDescriptor, ElementSpaceSize3D) +TEST(MakeNaiveTensorDescriptor, WithPaddingStrideSlice) { + // 2x3x4 tensor with strides [1, 5, 35] - like a slice from a 5x7xN column-major tensor + // This tests padding where there's space for extra elements + // element_space_size = 1 + (2-1)*1 + (3-1)*5 + (4-1)*35 = 1 + 1 + 10 + 105 = 117 constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<4>{}); - constexpr auto strides = make_tuple(Number<12>{}, Number<4>{}, Number<1>{}); + constexpr auto strides = make_tuple(Number<1>{}, Number<5>{}, Number<35>{}); constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); - EXPECT_EQ(desc.GetElementSpaceSize(), 24); + EXPECT_EQ(desc.GetElementSpaceSize(), 117); } -TEST(MakeNaiveTensorDescriptor, BroadcastDimension) +TEST(MakeNaiveTensorDescriptor, ColumnMajor) { - // 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>{}); + // 5x4 tensor with column-major strides [1, 5] + // element_space_size = 1 + (5-1)*1 + (4-1)*5 = 1 + 4 + 15 = 20 + constexpr auto lengths = make_tuple(Number<5>{}, Number<4>{}); + constexpr auto strides = make_tuple(Number<1>{}, Number<5>{}); constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); - EXPECT_EQ(desc.GetElementSpaceSize(), 5); + EXPECT_EQ(desc.GetElementSpaceSize(), 20); } // Test with runtime values (index_t instead of Number<>) -TEST(ComputeElementSpaceSizeRuntime, Simple2D) +TEST(MakeNaiveTensorDescriptorRuntime, Simple2D) { - const auto lengths = make_tuple(index_t{4}, index_t{4}); + // 5x4 tensor with row-major strides + const auto lengths = make_tuple(index_t{5}, 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); + const auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 20); } -TEST(ComputeElementSpaceSizeRuntime, WithPadding) +TEST(MakeNaiveTensorDescriptorRuntime, 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); + // 11x7x3 tensor with arbitrary strides (using prime numbers) + const auto lengths = make_tuple(index_t{11}, index_t{7}, index_t{3}); + const auto strides = make_tuple(index_t{2}, index_t{97}, index_t{23}); + const auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 649); } From 97873bc0d53e0faffe92dccbfd47ab3d6914d412 Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Mon, 19 Jan 2026 23:41:39 -0600 Subject: [PATCH 3/8] Expand tensor descriptor test coverage - Add tests for make_naive_tensor_descriptor_packed (1D, 2D, 3D) - Add tests for make_naive_tensor_descriptor_aligned (2D, 3D) - Add 1D tensor tests with explicit strides - Ensure all shapes use distinct, coprime dimensions --- test/util/unit_tensor_descriptor_helper.cpp | 102 +++++++++++++++++--- 1 file changed, 86 insertions(+), 16 deletions(-) diff --git a/test/util/unit_tensor_descriptor_helper.cpp b/test/util/unit_tensor_descriptor_helper.cpp index d0d67b38e31..b4ce884b786 100644 --- a/test/util/unit_tensor_descriptor_helper.cpp +++ b/test/util/unit_tensor_descriptor_helper.cpp @@ -53,40 +53,110 @@ TEST(MakeNaiveTensorDescriptor, WithPaddingArbitrary) TEST(MakeNaiveTensorDescriptor, WithPaddingStrideSlice) { - // 2x3x4 tensor with strides [1, 5, 35] - like a slice from a 5x7xN column-major tensor + // 2x3x5 tensor with strides [1, 7, 35] - like a slice from a 7x7xN column-major tensor // This tests padding where there's space for extra elements - // element_space_size = 1 + (2-1)*1 + (3-1)*5 + (4-1)*35 = 1 + 1 + 10 + 105 = 117 - constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<4>{}); - constexpr auto strides = make_tuple(Number<1>{}, Number<5>{}, Number<35>{}); + // element_space_size = 1 + (2-1)*1 + (3-1)*7 + (5-1)*35 = 1 + 1 + 14 + 140 = 156 + constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<5>{}); + constexpr auto strides = make_tuple(Number<1>{}, Number<7>{}, Number<35>{}); constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); - EXPECT_EQ(desc.GetElementSpaceSize(), 117); + EXPECT_EQ(desc.GetElementSpaceSize(), 156); } TEST(MakeNaiveTensorDescriptor, ColumnMajor) { - // 5x4 tensor with column-major strides [1, 5] - // element_space_size = 1 + (5-1)*1 + (4-1)*5 = 1 + 4 + 15 = 20 - constexpr auto lengths = make_tuple(Number<5>{}, Number<4>{}); - constexpr auto strides = make_tuple(Number<1>{}, Number<5>{}); + // 7x3 tensor with column-major strides [1, 7] + // element_space_size = 1 + (7-1)*1 + (3-1)*7 = 1 + 6 + 14 = 21 + constexpr auto lengths = make_tuple(Number<7>{}, Number<3>{}); + constexpr auto strides = make_tuple(Number<1>{}, Number<7>{}); constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); - EXPECT_EQ(desc.GetElementSpaceSize(), 20); + EXPECT_EQ(desc.GetElementSpaceSize(), 21); } // Test with runtime values (index_t instead of Number<>) TEST(MakeNaiveTensorDescriptorRuntime, Simple2D) { - // 5x4 tensor with row-major strides - const auto lengths = make_tuple(index_t{5}, index_t{4}); + // 9x4 tensor with row-major strides + // element_space_size = 1 + (9-1)*4 + (4-1)*1 = 1 + 32 + 3 = 36 + const auto lengths = make_tuple(index_t{9}, index_t{4}); const auto strides = make_tuple(index_t{4}, index_t{1}); const auto desc = make_naive_tensor_descriptor(lengths, strides); - EXPECT_EQ(desc.GetElementSpaceSize(), 20); + EXPECT_EQ(desc.GetElementSpaceSize(), 36); } TEST(MakeNaiveTensorDescriptorRuntime, WithPadding) { - // 11x7x3 tensor with arbitrary strides (using prime numbers) - const auto lengths = make_tuple(index_t{11}, index_t{7}, index_t{3}); + // 13x5x2 tensor with arbitrary strides (using prime numbers) + // element_space_size = 1 + (13-1)*2 + (5-1)*97 + (2-1)*23 = 1 + 24 + 388 + 23 = 436 + const auto lengths = make_tuple(index_t{13}, index_t{5}, index_t{2}); const auto strides = make_tuple(index_t{2}, index_t{97}, index_t{23}); const auto desc = make_naive_tensor_descriptor(lengths, strides); - EXPECT_EQ(desc.GetElementSpaceSize(), 649); + EXPECT_EQ(desc.GetElementSpaceSize(), 436); +} + +// Test 1D tensors with explicit strides +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize1D) +{ + // 13-element 1D tensor with stride 1 + constexpr auto lengths = make_tuple(Number<13>{}); + constexpr auto strides = make_tuple(Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 13); +} + +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize1DStrided) +{ + // 7-element 1D tensor with stride 3 (every 3rd element) + // element_space_size = 1 + (7-1)*3 = 19 + constexpr auto lengths = make_tuple(Number<7>{}); + constexpr auto strides = make_tuple(Number<3>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 19); +} + +// Test make_naive_tensor_descriptor_packed (contiguous layout) +// element_space_size = product of all lengths +TEST(MakeNaiveTensorDescriptorPacked, Simple1D) +{ + // 17-element packed tensor - element_space_size = 17 + constexpr auto lengths = make_tuple(Number<17>{}); + constexpr auto desc = make_naive_tensor_descriptor_packed(lengths); + EXPECT_EQ(desc.GetElementSpaceSize(), 17); +} + +TEST(MakeNaiveTensorDescriptorPacked, Simple2D) +{ + // 6x5 packed tensor - element_space_size = 6*5 = 30 + constexpr auto lengths = make_tuple(Number<6>{}, Number<5>{}); + constexpr auto desc = make_naive_tensor_descriptor_packed(lengths); + EXPECT_EQ(desc.GetElementSpaceSize(), 30); +} + +TEST(MakeNaiveTensorDescriptorPacked, Simple3D) +{ + // 4x5x9 packed tensor - element_space_size = 4*5*9 = 180 + constexpr auto lengths = make_tuple(Number<4>{}, Number<5>{}, Number<9>{}); + constexpr auto desc = make_naive_tensor_descriptor_packed(lengths); + EXPECT_EQ(desc.GetElementSpaceSize(), 180); +} + +// Test make_naive_tensor_descriptor_aligned (stride alignment for memory access) +// Aligns the second-to-last stride to be a multiple of 'align' +TEST(MakeNaiveTensorDescriptorAligned, Align4) +{ + // 5x3 tensor aligned to 4 elements + // strides[1] = 1, strides[0] = integer_least_multiple(3, 4) = 4 + // element_space_size = 1 + (5-1)*4 + (3-1)*1 = 1 + 16 + 2 = 19 + constexpr auto lengths = make_tuple(Number<5>{}, Number<3>{}); + constexpr auto desc = make_naive_tensor_descriptor_aligned(lengths, Number<4>{}); + EXPECT_EQ(desc.GetElementSpaceSize(), 19); +} + +TEST(MakeNaiveTensorDescriptorAligned, Align8) +{ + // 3x5x7 tensor aligned to 8 elements + // strides[2] = 1, strides[1] = integer_least_multiple(7, 8) = 8, strides[0] = 5*8 = 40 + // element_space_size = 1 + (3-1)*40 + (5-1)*8 + (7-1)*1 = 1 + 80 + 32 + 6 = 119 + constexpr auto lengths = make_tuple(Number<3>{}, Number<5>{}, Number<7>{}); + constexpr auto desc = make_naive_tensor_descriptor_aligned(lengths, Number<8>{}); + EXPECT_EQ(desc.GetElementSpaceSize(), 119); } From 1be9e524ceb1abde4338b18427abc3fa24461afa Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Tue, 20 Jan 2026 10:36:10 -0600 Subject: [PATCH 4/8] Use unique input values in sequence helper tests - Vary sequence values across all tests to avoid repetition - Use prime numbers for distinct, coprime test data - Add TwoSequences test for unpack_and_merge_sequences --- test/util/unit_sequence_helper.cpp | 52 +++++++++++++++++------------- 1 file changed, 30 insertions(+), 22 deletions(-) diff --git a/test/util/unit_sequence_helper.cpp b/test/util/unit_sequence_helper.cpp index 4f088d262b1..2f5894de003 100644 --- a/test/util/unit_sequence_helper.cpp +++ b/test/util/unit_sequence_helper.cpp @@ -45,25 +45,25 @@ TEST(GenerateIdentitySequences, WithNumber) // 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 + 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<10>(Sequence<10, 20, 30>{}); + 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<30>(Sequence<10, 20, 30>{}); - EXPECT_EQ(result, 2); + 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<99>(Sequence<1, 2, 3, 4, 5>{}); + constexpr auto result = sequence_find_value<50>(Sequence<2, 8, 14, 26>{}); EXPECT_EQ(result, -1); } @@ -76,8 +76,8 @@ TEST(SequenceFindValue, EmptySequence) // 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); + 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); @@ -85,8 +85,8 @@ TEST(FindInTupleOfSequences, FindInFirstSequence) 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); + 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); @@ -94,17 +94,17 @@ TEST(FindInTupleOfSequences, FindInMiddleSequence) 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); + 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, 1); // At position 1 within that sequence + 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<0, 1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}); - constexpr auto result = find_in_tuple_of_sequences<99>(tuple_of_seqs); + 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); } @@ -117,26 +117,34 @@ TEST(FindInTupleOfSequences, EmptyTuple) 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); + 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, 1); + 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<1, 2>{}, Sequence<3, 4>{}, Sequence<5, 6>{}); + 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<1, 2, 3, 4, 5, 6>; + using Expected = Sequence<7, 11, 13, 17, 19, 23>; EXPECT_TRUE((is_same, Expected>::value)); } TEST(UnpackAndMergeSequences, SingleSequence) { - constexpr auto tuple_of_seqs = make_tuple(Sequence<1, 2, 3>{}); + 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, 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<1, 2, 3>; + using Expected = Sequence<2, 3, 5, 7, 11>; EXPECT_TRUE((is_same, Expected>::value)); } From 1a5d956be7f124df46982a19d1fd536cff1c11c5 Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Tue, 20 Jan 2026 10:37:36 -0600 Subject: [PATCH 5/8] Use unique prime values in container helper tests - Replace repeated 1,2,3,4 values with unique primes - Each test now uses distinct input values --- test/util/unit_container_helper.cpp | 104 +++++++++++++++------------- 1 file changed, 54 insertions(+), 50 deletions(-) diff --git a/test/util/unit_container_helper.cpp b/test/util/unit_container_helper.cpp index a817ebf21c5..e0fb28322a6 100644 --- a/test/util/unit_container_helper.cpp +++ b/test/util/unit_container_helper.cpp @@ -10,132 +10,136 @@ 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 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>{}], 1); - EXPECT_EQ(result[Number<1>{}], 2); - EXPECT_EQ(result[Number<2>{}], 3); - EXPECT_EQ(result[Number<3>{}], 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<1>{}); - constexpr auto t2 = make_tuple(Number<2>{}); - constexpr auto t3 = make_tuple(Number<3>{}); + 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(), 3); - EXPECT_EQ(result[Number<0>{}], 1); - EXPECT_EQ(result[Number<1>{}], 2); - EXPECT_EQ(result[Number<2>{}], 3); + 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<1>{}, Number<2>{}); + 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>{}], 1); - EXPECT_EQ(result[Number<1>{}], 2); + EXPECT_EQ(result[Number<0>{}], 37); + EXPECT_EQ(result[Number<1>{}], 41); } TEST(ContainerConcat, ConcatSingleTuple) { - constexpr auto t1 = make_tuple(Number<1>{}, Number<2>{}, Number<3>{}); + 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(1, 2); - constexpr auto a2 = make_array(3, 4); + 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>{}], 1); - EXPECT_EQ(result[Number<1>{}], 2); - EXPECT_EQ(result[Number<2>{}], 3); - EXPECT_EQ(result[Number<3>{}], 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<42>{}); + constexpr auto result = make_uniform_tuple<3>(Number<73>{}); EXPECT_EQ(result.Size(), 3); - EXPECT_EQ(result[Number<0>{}], 42); - EXPECT_EQ(result[Number<1>{}], 42); - EXPECT_EQ(result[Number<2>{}], 42); + 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<99>{}); + constexpr auto result = make_uniform_tuple<1>(Number<79>{}); EXPECT_EQ(result.Size(), 1); - EXPECT_EQ(result[Number<0>{}], 99); + EXPECT_EQ(result[Number<0>{}], 79); } TEST(MakeUniformTuple, Size0) { - constexpr auto result = make_uniform_tuple<0>(Number<42>{}); + 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<7>{}); + constexpr auto result = make_uniform_tuple<5>(Number<89>{}); 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); + 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<1>{}, Number<2>{}, Number<3>{}); + auto result = functor(Number<97>{}, Number<101>{}, Number<103>{}); EXPECT_EQ(result.Size(), 3); - EXPECT_EQ(result[Number<0>{}], 1); - EXPECT_EQ(result[Number<1>{}], 2); - EXPECT_EQ(result[Number<2>{}], 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<2>{}, Number<3>{}); - constexpr auto result = container_push_front(t, Number<1>{}); + 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>{}], 1); - EXPECT_EQ(result[Number<1>{}], 2); - EXPECT_EQ(result[Number<2>{}], 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<1>{}, Number<2>{}); - constexpr auto result = container_push_back(t, Number<3>{}); + 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>{}], 1); - EXPECT_EQ(result[Number<1>{}], 2); - EXPECT_EQ(result[Number<2>{}], 3); + EXPECT_EQ(result[Number<0>{}], 127); + EXPECT_EQ(result[Number<1>{}], 131); + EXPECT_EQ(result[Number<2>{}], 137); } From e1f2d45ef44388dcbd9cb9734a0ceb9f57df9d49 Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Tue, 20 Jan 2026 10:43:49 -0600 Subject: [PATCH 6/8] Add 8D tensor descriptor tests Test high-dimensional tensors to verify no integer overflow in element space size calculations. Includes: - 8D packed test with prime dimensions (180180 elements) - 8D permuted layout test with non-monotonous strides (memory order differs from logical order) --- test/util/unit_tensor_descriptor_helper.cpp | 26 +++++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/test/util/unit_tensor_descriptor_helper.cpp b/test/util/unit_tensor_descriptor_helper.cpp index b4ce884b786..427a5faa117 100644 --- a/test/util/unit_tensor_descriptor_helper.cpp +++ b/test/util/unit_tensor_descriptor_helper.cpp @@ -160,3 +160,29 @@ TEST(MakeNaiveTensorDescriptorAligned, Align8) constexpr auto desc = make_naive_tensor_descriptor_aligned(lengths, Number<8>{}); EXPECT_EQ(desc.GetElementSpaceSize(), 119); } + +// Test high-dimensional tensors (8D) to verify no integer overflow +TEST(MakeNaiveTensorDescriptorPacked, Simple8D) +{ + // 8D packed tensor with small prime dimensions: 2x3x5x7x11x13x2x3 + // element_space_size = 2*3*5*7*11*13*2*3 = 180180 + constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<5>{}, Number<7>{}, + Number<11>{}, Number<13>{}, Number<2>{}, Number<3>{}); + constexpr auto desc = make_naive_tensor_descriptor_packed(lengths); + EXPECT_EQ(desc.GetElementSpaceSize(), 180180); +} + +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize8D) +{ + // 8D tensor with permuted layout (non-monotonous strides): 2x3x5x7x11x13x17x19 + // Memory order: [dim7, dim3, dim1, dim5, dim2, dim0, dim6, dim4] + // This gives strides: [25935, 133, 5187, 19, 881790, 399, 51870, 1] + // Note: strides go up/down/up/down - not monotonously increasing or decreasing + // element_space_size = 2*3*5*7*11*13*17*19 = 9699690 + constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<5>{}, Number<7>{}, + Number<11>{}, Number<13>{}, Number<17>{}, Number<19>{}); + constexpr auto strides = make_tuple(Number<25935>{}, Number<133>{}, Number<5187>{}, Number<19>{}, + Number<881790>{}, Number<399>{}, Number<51870>{}, Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 9699690); +} From af1f05efada73fda5b4550eff0f8f23de5bad278 Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Tue, 20 Jan 2026 10:52:54 -0600 Subject: [PATCH 7/8] Add 7D broadcast test with non-adjacent dimensions Tests element space size calculation when two non-adjacent middle dimensions (dims 2 and 4) have stride 0 (broadcast). --- test/util/unit_tensor_descriptor_helper.cpp | 56 ++++++++++++++++++--- 1 file changed, 50 insertions(+), 6 deletions(-) diff --git a/test/util/unit_tensor_descriptor_helper.cpp b/test/util/unit_tensor_descriptor_helper.cpp index 427a5faa117..5f4d9e08be0 100644 --- a/test/util/unit_tensor_descriptor_helper.cpp +++ b/test/util/unit_tensor_descriptor_helper.cpp @@ -40,6 +40,32 @@ TEST(MakeNaiveTensorDescriptor, BroadcastDimension) EXPECT_EQ(desc.GetElementSpaceSize(), 5); } +TEST(MakeNaiveTensorDescriptor, BroadcastNonAdjacentDims7D) +{ + // 7D tensor: 2x3x5x7x11x13x17 with broadcast on dims 2 and 4 (non-adjacent) + // Underlying data is 2x3x7x13x17 = 9282 elements, broadcast over dims 2 and 4 + // strides: [4641, 1547, 0, 221, 0, 17, 1] + // element_space_size = 1 + (2-1)*4641 + (3-1)*1547 + (5-1)*0 + (7-1)*221 + // + (11-1)*0 + (13-1)*17 + (17-1)*1 + // = 1 + 4641 + 3094 + 0 + 1326 + 0 + 204 + 16 = 9282 + constexpr auto lengths = make_tuple(Number<2>{}, + Number<3>{}, + Number<5>{}, + Number<7>{}, + Number<11>{}, + Number<13>{}, + Number<17>{}); + constexpr auto strides = make_tuple(Number<4641>{}, + Number<1547>{}, + Number<0>{}, + Number<221>{}, + Number<0>{}, + Number<17>{}, + Number<1>{}); + constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); + EXPECT_EQ(desc.GetElementSpaceSize(), 9282); +} + TEST(MakeNaiveTensorDescriptor, WithPaddingArbitrary) { // 11x7x3 tensor with arbitrary strides [2, 97, 23] (prime numbers, no common factors) @@ -166,8 +192,14 @@ TEST(MakeNaiveTensorDescriptorPacked, Simple8D) { // 8D packed tensor with small prime dimensions: 2x3x5x7x11x13x2x3 // element_space_size = 2*3*5*7*11*13*2*3 = 180180 - constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<5>{}, Number<7>{}, - Number<11>{}, Number<13>{}, Number<2>{}, Number<3>{}); + constexpr auto lengths = make_tuple(Number<2>{}, + Number<3>{}, + Number<5>{}, + Number<7>{}, + Number<11>{}, + Number<13>{}, + Number<2>{}, + Number<3>{}); constexpr auto desc = make_naive_tensor_descriptor_packed(lengths); EXPECT_EQ(desc.GetElementSpaceSize(), 180180); } @@ -179,10 +211,22 @@ TEST(MakeNaiveTensorDescriptor, ElementSpaceSize8D) // This gives strides: [25935, 133, 5187, 19, 881790, 399, 51870, 1] // Note: strides go up/down/up/down - not monotonously increasing or decreasing // element_space_size = 2*3*5*7*11*13*17*19 = 9699690 - constexpr auto lengths = make_tuple(Number<2>{}, Number<3>{}, Number<5>{}, Number<7>{}, - Number<11>{}, Number<13>{}, Number<17>{}, Number<19>{}); - constexpr auto strides = make_tuple(Number<25935>{}, Number<133>{}, Number<5187>{}, Number<19>{}, - Number<881790>{}, Number<399>{}, Number<51870>{}, Number<1>{}); + constexpr auto lengths = make_tuple(Number<2>{}, + Number<3>{}, + Number<5>{}, + Number<7>{}, + Number<11>{}, + Number<13>{}, + Number<17>{}, + Number<19>{}); + constexpr auto strides = make_tuple(Number<25935>{}, + Number<133>{}, + Number<5187>{}, + Number<19>{}, + Number<881790>{}, + Number<399>{}, + Number<51870>{}, + Number<1>{}); constexpr auto desc = make_naive_tensor_descriptor(lengths, strides); EXPECT_EQ(desc.GetElementSpaceSize(), 9699690); } From aa2a5778e6a68cd9b1d2ffe25a39439b6a1959df Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Tue, 20 Jan 2026 11:20:03 -0600 Subject: [PATCH 8/8] Apply clang-format to sequence helper tests --- test/util/unit_sequence_helper.cpp | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/test/util/unit_sequence_helper.cpp b/test/util/unit_sequence_helper.cpp index 2f5894de003..4e796f5081c 100644 --- a/test/util/unit_sequence_helper.cpp +++ b/test/util/unit_sequence_helper.cpp @@ -76,8 +76,9 @@ TEST(SequenceFindValue, EmptySequence) // 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); + 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); @@ -85,8 +86,9 @@ TEST(FindInTupleOfSequences, FindInFirstSequence) 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); + 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); @@ -127,9 +129,10 @@ TEST(FindInTupleOfSequences, SingleSequence) // 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>; + 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, Expected>::value)); }