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..e0fb28322a6 --- /dev/null +++ b/test/util/unit_container_helper.cpp @@ -0,0 +1,145 @@ +// 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<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); +} 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..4e796f5081c --- /dev/null +++ b/test/util/unit_sequence_helper.cpp @@ -0,0 +1,153 @@ +// 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<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, 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, 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, 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..5f4d9e08be0 --- /dev/null +++ b/test/util/unit_tensor_descriptor_helper.cpp @@ -0,0 +1,232 @@ +// 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 make_naive_tensor_descriptor (public API) +// Formula for element_space_size: 1 + sum((length[i] - 1) * stride[i]) + +TEST(MakeNaiveTensorDescriptor, ElementSpaceSize2D) +{ + // 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(MakeNaiveTensorDescriptor, ElementSpaceSize3D) +{ + // 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(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 desc = make_naive_tensor_descriptor(lengths, strides); + 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) + // 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(), 649); +} + +TEST(MakeNaiveTensorDescriptor, WithPaddingStrideSlice) +{ + // 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)*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(), 156); +} + +TEST(MakeNaiveTensorDescriptor, ColumnMajor) +{ + // 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(), 21); +} + +// Test with runtime values (index_t instead of Number<>) +TEST(MakeNaiveTensorDescriptorRuntime, Simple2D) +{ + // 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(), 36); +} + +TEST(MakeNaiveTensorDescriptorRuntime, WithPadding) +{ + // 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(), 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); +} + +// 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); +}