ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-30 11:47:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

Address review feedback on tensor descriptor helper tests

Browse Source 
- 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
This commit is contained in:
Max Podkorytov
2026-01-19 23:31:13 -06:00
parent 0c8188374a
commit b60d14ba89
1 changed files with 57 additions and 90 deletions
Download Patch File Download Diff File Expand all files Collapse all files

147
test/util/unit_tensor_descriptor_helper.cpp
View File

@@ -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 make_naive_tensor_descriptor (public API)
// Formula for element_space_size: 1 + sum((length[i] - 1) * stride[i])
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>{});
// 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(), 16);
EXPECT_EQ(desc.GetElementSpaceSize(), 20);
}
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>{});
// 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(), 24);
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 = 5
// 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 with runtime values (index_t instead of Number<>)
TEST(ComputeElementSpaceSizeRuntime, Simple2D)
TEST(MakeNaiveTensorDescriptor, WithPaddingArbitrary)
{
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);
// 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(ComputeElementSpaceSizeRuntime, WithPadding)
TEST(MakeNaiveTensorDescriptor, WithPaddingStrideSlice)
{
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);
// 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<1>{}, Number<5>{}, Number<35>{});
constexpr auto desc = make_naive_tensor_descriptor(lengths, strides);
EXPECT_EQ(desc.GetElementSpaceSize(), 117);
}
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>{});
constexpr auto desc = make_naive_tensor_descriptor(lengths, strides);
EXPECT_EQ(desc.GetElementSpaceSize(), 20);
}
// 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});
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);
}
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});
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);
}