Files
composable_kernel/test/ck_tile/utility/test_sequence.cpp
Christopher Millette 56e1d5da08 [rocm-libraries] ROCm/rocm-libraries#5028 (commit 5131491)
[CK_TILE] Optimize ck_tile::sequence to reduce template
 instantiation depth [2A] (#5028)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Summary

### Rationale
`ck_tile::sequence` is the most fundamental metaprogramming type in
ck_tile — it underpins
tensor dimensions, strides, loop bounds, and index calculations. Six of
its metafunctions
use recursive template instantiation, producing O(N) to O(N²)
intermediate types that
the compiler must process. When these are used inside deeply nested GEMM
pipelines with
large dimension counts, the cumulative instantiation overhead becomes a
significant
contributor to frontend compile time.

Measurements on `test_gemm_pipeline_compv6` show 84,288
`InstantiateFunction` calls in
the frontend alone. Reducing template instantiation depth in these core
utilities has a
multiplicative effect because they are called from hundreds of sites.

### What changed

| Metafunction | Before | After |
|---|---|---|
| `sequence::modify` | O(N) recursive split/merge | O(1) pack expansion
|
| `sequence_gen` | O(log N) recursive binary split | O(1) via
`__make_integer_seq` |
| `uniform_sequence_gen` | Delegates to `sequence_gen` | O(1) via
`__make_integer_seq` |
| `sequence_reverse_inclusive_scan` | O(N) recursive | O(1) constexpr
for-loop + pack expansion |
| `sequence_inclusive_scan` | Computed via reverse + flip | O(1)
constexpr for-loop (unified impl) |
| `sequence_exclusive_scan` | O(N) recursive merge chain | O(1)
constexpr for-loop + pack expansion |
| `sequence_map_inverse` | O(N²) recursive modify calls | O(1) constexpr
for-loop + pack expansion |

Supporting changes:
- Portable `__type_pack_element` fallback with `__has_builtin` guard
(hipRTC-safe, no `<tuple>` dependency)
- Renamed reserved `__integer_sequence` to `integer_sequence_wrapper`
- Adopted `static_array` from develop (PR #4355) for constexpr
computation
- Unified forward and reverse inclusive scan into a single
`sequence_inclusive_scan_impl` with `bool Reverse` template parameter
- Added `sequence_inclusive_scan` struct (new public API for forward
scan direction)
- Replaced recursive `sequence_exclusive_scan` (3 template
specializations) with `sequence_exclusive_scan_impl` using the same
constexpr for-loop pattern as inclusive scan
- Rewired `exclusive_scan_sequence` and `prefix_sum_sequence` to use new
impl
- Added `CK_TILE_HOST_DEVICE` to `exclusive_scan_sequence` and
`prefix_sum_sequence` to match sibling scan function annotations

### Technical debt and housekeeping
- Unified all `namespace impl` to `namespace detail` across sequence.hpp
for consistency
- Removed dead comment block (orphaned `integer_sequence` alternative)
- Added defensive `static_assert(sizeof...(Is) > 0)` in
`sequence_map_inverse::build_inverse`
- Converted all multi-line Doxygen blocks from `///` to `/** */` per
style guide
- Corrected `constexpr static` to `static constexpr` keyword ordering in
`static_array`
- Added blank line between `#pragma once` and first `#include` in
`static_array.hpp`
- Trimmed redundant 4-line comment on `sequence_gen_helper` to a
one-liner
- Moved `sequence_gen` Doxygen comment below `namespace detail` block so
it directly precedes the struct it documents
- Added Doxygen `@brief`/`@tparam`/`@pre` documentation for
`sequence_gen` and `sequence_map_inverse` public APIs
- Added `@brief` documentation to `static_array` explaining relationship
to `ck_tile::array`
- Added scope comment at `namespace detail` openings

**Note:** `private:`/`public:` access modifier indentation is enforced
at 4 spaces by
`.clang-format`. The style guide calls for left-alignment, but the
formatter overrides
this. Requires a `.clang-format` config change to resolve — not
addressable in code.

### `static_array` hardening (from develop's PR #4355)
- Added zero-length array guard (`T elems[N > 0 ? N : 1]`)
- Added `CK_TILE_HOST_DEVICE` annotations to `operator[]` and `size()`
- Added `#include "ck_tile/core/config.hpp"` (IWYU for
`CK_TILE_HOST_DEVICE`)

### Value
Combined with the `static_ford` changes, measured impact on
`test_gemm_pipeline_compv6`:
- **Frontend: -28.9%** (InstantiateFunction: 84,288 → 69,439)
- **Backend: -13.1%** (CodeGen Functions: 3,170 → 2,203)
- **Wall-clock: -16.3%** (611.6s → 512.2s)

### Files changed (4)
- `sequence.hpp`: Metafunction optimizations, namespace unification,
documentation, style fixes
- `static_array.hpp`: Zero-length guard, `CK_TILE_HOST_DEVICE`,
documentation, style fixes
- `test_sequence.cpp`: 50 unit tests with runtime `EXPECT_EQ` assertions
(new file)
- `CMakeLists.txt`: Register new test target

## Test plan
- [x] 50 runtime unit tests covering all optimized and pre-existing
sequence APIs
- [x] Edge cases: empty sequences, single-element, larger sizes (N=8),
negative values, non-trivial init values
- [x] Both functor signatures tested (`operator()(index_t)` and
`operator()(number<I>)`)
- [x] Both scan reducers (`plus`, `multiplies`) with forward, reverse,
inclusive, and exclusive directions
- [x] Exclusive scan: sum, product, single, empty, non-zero init
- [x] Prefix sum: N+1 output verification, single, empty
- [x] Permutation round-trip verification for `sequence_map_inverse`
- [x] Full sequence public API coverage: modify, gen, uniform_gen, scans
(inclusive, exclusive, prefix sum), map_inverse, make_index_sequence,
size/sum/product, push/pop, reverse, extract, merge, arithmetic
operators, equality, transform
- [x] Portable `__type_pack_element` fallback tested implicitly (same
`at_index_t` interface)

🤖 Generated with [Claude Code](https://claude.com/claude-code)

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
2026-03-11 20:26:11 +00:00

649 lines
19 KiB
C++

// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <gtest/gtest.h>
#include "ck_tile/core/container/sequence.hpp"
#include "ck_tile/core/utility/functional.hpp"
#include "ck_tile/core/numeric/math.hpp"
using namespace ck_tile;
// ============================================================================
// sequence::modify tests
// ============================================================================
TEST(CkTileSequence, ModifyFirstElement)
{
constexpr auto result = sequence<1, 2, 3, 4>{}.modify(number<0>{}, number<99>{});
EXPECT_EQ(result.at(0), 99);
EXPECT_EQ(result.at(1), 2);
EXPECT_EQ(result.at(2), 3);
EXPECT_EQ(result.at(3), 4);
}
TEST(CkTileSequence, ModifyLastElement)
{
constexpr auto result = sequence<1, 2, 3, 4>{}.modify(number<3>{}, number<99>{});
EXPECT_EQ(result.at(0), 1);
EXPECT_EQ(result.at(3), 99);
}
TEST(CkTileSequence, ModifyMiddleElement)
{
constexpr auto result = sequence<5, 5, 5>{}.modify(number<1>{}, number<0>{});
EXPECT_EQ(result.at(0), 5);
EXPECT_EQ(result.at(1), 0);
EXPECT_EQ(result.at(2), 5);
}
TEST(CkTileSequence, ModifySingleElement)
{
constexpr auto result = sequence<42>{}.modify(number<0>{}, number<99>{});
EXPECT_EQ(result.at(0), 99);
}
TEST(CkTileSequence, ModifyNegativeValue)
{
constexpr auto result = sequence<1, 2, 3>{}.modify(number<1>{}, number<-1>{});
EXPECT_EQ(result.at(0), 1);
EXPECT_EQ(result.at(1), -1);
EXPECT_EQ(result.at(2), 3);
}
// ============================================================================
// sequence_gen tests
// ============================================================================
TEST(CkTileSequence, SequenceGenZero)
{
using Result = typename sequence_gen<0, identity>::type;
EXPECT_EQ(Result::size(), 0);
}
TEST(CkTileSequence, SequenceGenZeroNonIdentityFunctor)
{
// N=0 specialization should produce empty sequence regardless of functor.
// Use sequence_gen<1, F> to exercise the functor (suppresses -Wunused-member-function),
// then verify that N=0 still produces an empty sequence with the same functor type.
struct F
{
constexpr index_t operator()(index_t) const { return 999; }
};
using ResultOne = typename sequence_gen<1, F>::type;
using ResultZero = typename sequence_gen<0, F>::type;
EXPECT_EQ(ResultOne{}.at(0), 999);
EXPECT_EQ(ResultZero::size(), 0);
}
TEST(CkTileSequence, SequenceGenIdentity)
{
struct F
{
constexpr index_t operator()(index_t i) const { return i; }
};
using Result = typename sequence_gen<5, F>::type;
EXPECT_EQ(Result::size(), 5);
for(index_t i = 0; i < 5; ++i)
{
EXPECT_EQ(Result{}.at(i), i);
}
}
TEST(CkTileSequence, SequenceGenDouble)
{
struct F
{
constexpr index_t operator()(index_t i) const { return i * 2; }
};
using Result = typename sequence_gen<4, F>::type;
EXPECT_EQ(Result{}.at(0), 0);
EXPECT_EQ(Result{}.at(1), 2);
EXPECT_EQ(Result{}.at(2), 4);
EXPECT_EQ(Result{}.at(3), 6);
}
TEST(CkTileSequence, SequenceGenSingle)
{
struct F
{
constexpr index_t operator()(index_t) const { return 42; }
};
using Result = typename sequence_gen<1, F>::type;
EXPECT_EQ(Result::size(), 1);
EXPECT_EQ(Result{}.at(0), 42);
}
TEST(CkTileSequence, SequenceGenLarger)
{
struct F
{
constexpr index_t operator()(index_t i) const { return i * i; }
};
using Result = typename sequence_gen<8, F>::type;
EXPECT_EQ(Result{}.at(7), 49);
}
// Defined at namespace scope because template members are not allowed in local classes.
namespace {
struct NumberParamFunctor
{
template <index_t I>
constexpr index_t operator()(number<I>) const
{
return I + 10;
}
};
} // anonymous namespace
TEST(CkTileSequence, SequenceGenWithNumberParam)
{
// Verify functor taking number<I> directly (the documented API contract)
using Result = typename sequence_gen<4, NumberParamFunctor>::type;
EXPECT_EQ(Result{}.at(0), 10);
EXPECT_EQ(Result{}.at(3), 13);
}
// ============================================================================
// uniform_sequence_gen tests
// ============================================================================
TEST(CkTileSequence, UniformSequenceGenZero)
{
using Result = typename uniform_sequence_gen<0, 7>::type;
EXPECT_EQ(Result::size(), 0);
}
TEST(CkTileSequence, UniformSequenceGenSingle)
{
using Result = typename uniform_sequence_gen<1, 99>::type;
EXPECT_EQ(Result{}.at(0), 99);
}
TEST(CkTileSequence, UniformSequenceGenMultiple)
{
using Result = typename uniform_sequence_gen<4, 0>::type;
for(index_t i = 0; i < 4; ++i)
{
EXPECT_EQ(Result{}.at(i), 0);
}
}
TEST(CkTileSequence, UniformSequenceGenLarger)
{
using Result = typename uniform_sequence_gen<8, 3>::type;
for(index_t i = 0; i < 8; ++i)
{
EXPECT_EQ(Result{}.at(i), 3);
}
}
// ============================================================================
// sequence_reverse_inclusive_scan tests — runtime value verification
// ============================================================================
TEST(CkTileSequence, ReverseInclusiveScanProduct)
{
using Result = typename sequence_reverse_inclusive_scan<sequence<1, 2, 3, 4>,
multiplies<index_t>,
1>::type;
// result[3]=4*1=4, result[2]=3*4=12, result[1]=2*12=24, result[0]=1*24=24
EXPECT_EQ(Result{}.at(0), 24);
EXPECT_EQ(Result{}.at(1), 24);
EXPECT_EQ(Result{}.at(2), 12);
EXPECT_EQ(Result{}.at(3), 4);
}
TEST(CkTileSequence, ReverseInclusiveScanSum)
{
using Result =
typename sequence_reverse_inclusive_scan<sequence<1, 2, 3, 4>, plus<index_t>, 0>::type;
// result[3]=4, result[2]=7, result[1]=9, result[0]=10
EXPECT_EQ(Result{}.at(0), 10);
EXPECT_EQ(Result{}.at(1), 9);
EXPECT_EQ(Result{}.at(2), 7);
EXPECT_EQ(Result{}.at(3), 4);
}
TEST(CkTileSequence, ReverseInclusiveScanSingleElement)
{
using Result = typename sequence_reverse_inclusive_scan<sequence<5>, plus<index_t>, 0>::type;
EXPECT_EQ(Result{}.at(0), 5);
}
TEST(CkTileSequence, ReverseInclusiveScanEmpty)
{
using Result = typename sequence_reverse_inclusive_scan<sequence<>, plus<index_t>, 0>::type;
EXPECT_EQ(Result::size(), 0);
}
// ============================================================================
// sequence_inclusive_scan (forward) tests — runtime value verification
// ============================================================================
TEST(CkTileSequence, ForwardInclusiveScanSum)
{
using Result = typename sequence_inclusive_scan<sequence<1, 2, 3, 4>, plus<index_t>, 0>::type;
// result[0]=1, result[1]=3, result[2]=6, result[3]=10
EXPECT_EQ(Result{}.at(0), 1);
EXPECT_EQ(Result{}.at(1), 3);
EXPECT_EQ(Result{}.at(2), 6);
EXPECT_EQ(Result{}.at(3), 10);
}
TEST(CkTileSequence, ForwardInclusiveScanProduct)
{
using Result =
typename sequence_inclusive_scan<sequence<1, 2, 3, 4>, multiplies<index_t>, 1>::type;
// result[0]=1, result[1]=2, result[2]=6, result[3]=24
EXPECT_EQ(Result{}.at(0), 1);
EXPECT_EQ(Result{}.at(1), 2);
EXPECT_EQ(Result{}.at(2), 6);
EXPECT_EQ(Result{}.at(3), 24);
}
TEST(CkTileSequence, ForwardInclusiveScanNonTrivialInit)
{
using Result = typename sequence_inclusive_scan<sequence<1, 2, 3>, plus<index_t>, 10>::type;
// init=10: result[0]=1+10=11, result[1]=2+11=13, result[2]=3+13=16
EXPECT_EQ(Result{}.at(0), 11);
EXPECT_EQ(Result{}.at(1), 13);
EXPECT_EQ(Result{}.at(2), 16);
}
TEST(CkTileSequence, ReverseInclusiveScanNonTrivialInit)
{
using Result =
typename sequence_reverse_inclusive_scan<sequence<1, 2, 3>, plus<index_t>, 10>::type;
// init=10: result[2]=3+10=13, result[1]=2+13=15, result[0]=1+15=16
EXPECT_EQ(Result{}.at(0), 16);
EXPECT_EQ(Result{}.at(1), 15);
EXPECT_EQ(Result{}.at(2), 13);
}
TEST(CkTileSequence, ForwardInclusiveScanSingleElement)
{
using Result = typename sequence_inclusive_scan<sequence<5>, plus<index_t>, 0>::type;
EXPECT_EQ(Result{}.at(0), 5);
}
TEST(CkTileSequence, ForwardInclusiveScanEmpty)
{
using Result = typename sequence_inclusive_scan<sequence<>, plus<index_t>, 0>::type;
EXPECT_EQ(Result::size(), 0);
}
// ============================================================================
// sequence_map_inverse tests — runtime round-trip verification
// ============================================================================
TEST(CkTileSequence, MapInverseIdentity)
{
using Result = typename sequence_map_inverse<sequence<0, 1, 2, 3>>::type;
for(index_t i = 0; i < 4; ++i)
{
EXPECT_EQ(Result{}.at(i), i);
}
}
TEST(CkTileSequence, MapInverseSwap)
{
using Result = typename sequence_map_inverse<sequence<1, 0>>::type;
EXPECT_EQ(Result{}.at(0), 1);
EXPECT_EQ(Result{}.at(1), 0);
}
TEST(CkTileSequence, MapInversePermutation)
{
using Input = sequence<2, 0, 1>;
using Result = typename sequence_map_inverse<Input>::type;
EXPECT_EQ(Result{}.at(0), 1);
EXPECT_EQ(Result{}.at(1), 2);
EXPECT_EQ(Result{}.at(2), 0);
// Verify round-trip: input[result[i]] == i for all i
for(index_t i = 0; i < 3; ++i)
{
EXPECT_EQ(Input{}.at(Result{}.at(i)), i);
}
}
TEST(CkTileSequence, MapInverseEmpty)
{
using Result = typename sequence_map_inverse<sequence<>>::type;
EXPECT_EQ(Result::size(), 0);
}
TEST(CkTileSequence, MapInverseSingle)
{
using Result = typename sequence_map_inverse<sequence<0>>::type;
EXPECT_EQ(Result{}.at(0), 0);
}
TEST(CkTileSequence, MapInverseRotation)
{
using Input = sequence<1, 2, 0>;
using Result = typename sequence_map_inverse<Input>::type;
for(index_t i = 0; i < 3; ++i)
{
EXPECT_EQ(Input{}.at(Result{}.at(i)), i);
}
}
TEST(CkTileSequence, MapInverse4D)
{
using Input = sequence<2, 0, 3, 1>;
using Result = typename sequence_map_inverse<Input>::type;
EXPECT_EQ(Result{}.at(0), 1);
EXPECT_EQ(Result{}.at(1), 3);
EXPECT_EQ(Result{}.at(2), 0);
EXPECT_EQ(Result{}.at(3), 2);
// Verify round-trip: input[result[i]] == i for all i
for(index_t i = 0; i < 4; ++i)
{
EXPECT_EQ(Input{}.at(Result{}.at(i)), i);
}
}
// ============================================================================
// make_index_sequence tests
// ============================================================================
TEST(CkTileSequence, MakeIndexSequenceZero)
{
using Result = make_index_sequence<0>;
EXPECT_EQ(Result::size(), 0);
}
TEST(CkTileSequence, MakeIndexSequenceOne)
{
using Result = make_index_sequence<1>;
EXPECT_EQ(Result::size(), 1);
EXPECT_EQ(Result{}.at(0), 0);
}
TEST(CkTileSequence, MakeIndexSequenceSmall)
{
using Result = make_index_sequence<5>;
EXPECT_EQ(Result::size(), 5);
for(index_t i = 0; i < 5; ++i)
{
EXPECT_EQ(Result{}.at(i), i);
}
}
// ============================================================================
// sequence basic accessors tests
// ============================================================================
TEST(CkTileSequence, SizeAndIsStatic)
{
EXPECT_EQ((sequence<1, 2, 3>::size()), 3);
EXPECT_EQ((sequence<>::size()), 0);
EXPECT_TRUE((sequence<1, 2, 3>::is_static()));
}
TEST(CkTileSequence, FrontAndBack)
{
constexpr auto s = sequence<10, 20, 30>{};
EXPECT_EQ(s.at(0), 10);
EXPECT_EQ(s.at(2), 30);
}
TEST(CkTileSequence, SumAndProduct)
{
EXPECT_EQ((sequence<1, 2, 3, 4>::sum()), 10);
EXPECT_EQ((sequence<1, 2, 3, 4>::product()), 24);
EXPECT_EQ((sequence<>::sum()), 0);
EXPECT_EQ((sequence<>::product()), 1);
EXPECT_EQ((sequence<5>::sum()), 5);
EXPECT_EQ((sequence<5>::product()), 5);
}
// ============================================================================
// sequence push/pop tests
// ============================================================================
TEST(CkTileSequence, PushFrontSequence)
{
constexpr auto result = sequence<3, 4>{}.push_front(sequence<1, 2>{});
EXPECT_EQ(result.at(0), 1);
EXPECT_EQ(result.at(1), 2);
EXPECT_EQ(result.at(2), 3);
EXPECT_EQ(result.at(3), 4);
}
TEST(CkTileSequence, PushBackSequence)
{
constexpr auto result = sequence<1, 2>{}.push_back(sequence<3, 4>{});
EXPECT_EQ(result.at(0), 1);
EXPECT_EQ(result.at(1), 2);
EXPECT_EQ(result.at(2), 3);
EXPECT_EQ(result.at(3), 4);
}
TEST(CkTileSequence, PopFront)
{
constexpr auto result = sequence_pop_front(sequence<1, 2, 3>{});
EXPECT_EQ(decltype(result)::size(), 2);
EXPECT_EQ(result.at(0), 2);
EXPECT_EQ(result.at(1), 3);
}
TEST(CkTileSequence, PopBack)
{
constexpr auto result = sequence_pop_back(sequence<1, 2, 3>{});
EXPECT_EQ(decltype(result)::size(), 2);
EXPECT_EQ(result.at(0), 1);
EXPECT_EQ(result.at(1), 2);
}
// ============================================================================
// sequence reverse tests
// ============================================================================
TEST(CkTileSequence, Reverse)
{
constexpr auto result = sequence<1, 2, 3, 4>{}.reverse();
EXPECT_EQ(result.at(0), 4);
EXPECT_EQ(result.at(1), 3);
EXPECT_EQ(result.at(2), 2);
EXPECT_EQ(result.at(3), 1);
}
TEST(CkTileSequence, ReverseSingle)
{
constexpr auto result = sequence<42>{}.reverse();
EXPECT_EQ(result.at(0), 42);
}
// ============================================================================
// sequence extract tests
// ============================================================================
TEST(CkTileSequence, Extract)
{
constexpr auto result = sequence<10, 20, 30, 40>{}.extract(sequence<2, 0, 3>{});
EXPECT_EQ(result.at(0), 30);
EXPECT_EQ(result.at(1), 10);
EXPECT_EQ(result.at(2), 40);
}
// ============================================================================
// sequence_merge tests
// ============================================================================
TEST(CkTileSequence, MergeTwoSequences)
{
constexpr auto result = merge_sequences(sequence<1, 2>{}, sequence<3, 4>{});
EXPECT_EQ(decltype(result)::size(), 4);
EXPECT_EQ(result.at(0), 1);
EXPECT_EQ(result.at(3), 4);
}
TEST(CkTileSequence, MergeWithEmpty)
{
constexpr auto result = merge_sequences(sequence<1, 2>{}, sequence<>{});
EXPECT_EQ(decltype(result)::size(), 2);
EXPECT_EQ(result.at(0), 1);
}
// ============================================================================
// sequence arithmetic operator tests
// ============================================================================
TEST(CkTileSequence, OperatorAdd)
{
constexpr auto result = sequence<1, 2, 3>{} + sequence<10, 20, 30>{};
EXPECT_EQ(result.at(0), 11);
EXPECT_EQ(result.at(1), 22);
EXPECT_EQ(result.at(2), 33);
}
TEST(CkTileSequence, OperatorSubtract)
{
constexpr auto result = sequence<10, 20, 30>{} - sequence<1, 2, 3>{};
EXPECT_EQ(result.at(0), 9);
EXPECT_EQ(result.at(1), 18);
EXPECT_EQ(result.at(2), 27);
}
TEST(CkTileSequence, OperatorMultiply)
{
constexpr auto result = sequence<2, 3, 4>{} * sequence<5, 6, 7>{};
EXPECT_EQ(result.at(0), 10);
EXPECT_EQ(result.at(1), 18);
EXPECT_EQ(result.at(2), 28);
}
TEST(CkTileSequence, OperatorAddScalar)
{
constexpr auto result = sequence<1, 2, 3>{} + number<10>{};
EXPECT_EQ(result.at(0), 11);
EXPECT_EQ(result.at(1), 12);
EXPECT_EQ(result.at(2), 13);
}
TEST(CkTileSequence, OperatorMultiplyScalar)
{
constexpr auto result = sequence<1, 2, 3>{} * number<10>{};
EXPECT_EQ(result.at(0), 10);
EXPECT_EQ(result.at(1), 20);
EXPECT_EQ(result.at(2), 30);
}
TEST(CkTileSequence, ScalarOperatorAdd)
{
constexpr auto result = number<100>{} + sequence<1, 2, 3>{};
EXPECT_EQ(result.at(0), 101);
EXPECT_EQ(result.at(1), 102);
EXPECT_EQ(result.at(2), 103);
}
// ============================================================================
// sequence equality tests
// ============================================================================
TEST(CkTileSequence, EqualityTrue) { EXPECT_TRUE((sequence<1, 2, 3>{} == sequence<1, 2, 3>{})); }
TEST(CkTileSequence, EqualityFalse) { EXPECT_FALSE((sequence<1, 2, 3>{} == sequence<1, 2, 4>{})); }
TEST(CkTileSequence, InequalityTrue) { EXPECT_TRUE((sequence<1, 2, 3>{} != sequence<1, 2, 4>{})); }
TEST(CkTileSequence, EqualityEmpty) { EXPECT_TRUE((sequence<>{} == sequence<>{})); }
// ============================================================================
// sequence transform tests
// ============================================================================
TEST(CkTileSequence, Transform)
{
struct Double
{
constexpr index_t operator()(index_t x) const { return x * 2; }
};
constexpr auto result = sequence<1, 2, 3>{}.transform(Double{});
EXPECT_EQ(result.at(0), 2);
EXPECT_EQ(result.at(1), 4);
EXPECT_EQ(result.at(2), 6);
}
// ============================================================================
// exclusive_scan_sequence tests
// ============================================================================
TEST(CkTileSequence, ExclusiveScanSum)
{
// <2, 3, 4> with Init=0, Add -> <0, 2, 5>
constexpr auto result =
exclusive_scan_sequence(sequence<2, 3, 4>{}, plus<index_t>{}, number<0>{});
EXPECT_EQ(result.at(0), 0);
EXPECT_EQ(result.at(1), 2);
EXPECT_EQ(result.at(2), 5);
}
TEST(CkTileSequence, ExclusiveScanProduct)
{
// <2, 3, 4> with Init=1, Mul -> <1, 2, 6>
constexpr auto result =
exclusive_scan_sequence(sequence<2, 3, 4>{}, multiplies<index_t>{}, number<1>{});
EXPECT_EQ(result.at(0), 1);
EXPECT_EQ(result.at(1), 2);
EXPECT_EQ(result.at(2), 6);
}
TEST(CkTileSequence, ExclusiveScanSingle)
{
constexpr auto result = exclusive_scan_sequence(sequence<5>{}, plus<index_t>{}, number<0>{});
EXPECT_EQ(decltype(result)::size(), 1);
EXPECT_EQ(result.at(0), 0);
}
TEST(CkTileSequence, ExclusiveScanEmpty)
{
constexpr auto result = exclusive_scan_sequence(sequence<>{}, plus<index_t>{}, number<0>{});
EXPECT_EQ(decltype(result)::size(), 0);
}
TEST(CkTileSequence, ExclusiveScanNonZeroInit)
{
// <1, 2, 3> with Init=10, Add -> <10, 11, 13>
constexpr auto result =
exclusive_scan_sequence(sequence<1, 2, 3>{}, plus<index_t>{}, number<10>{});
EXPECT_EQ(result.at(0), 10);
EXPECT_EQ(result.at(1), 11);
EXPECT_EQ(result.at(2), 13);
}
// ============================================================================
// prefix_sum_sequence tests
// ============================================================================
TEST(CkTileSequence, PrefixSumSequence)
{
// <2, 3, 4> -> <0, 2, 5, 9> (N+1 elements)
constexpr auto result = prefix_sum_sequence(sequence<2, 3, 4>{});
EXPECT_EQ(decltype(result)::size(), 4);
EXPECT_EQ(result.at(0), 0);
EXPECT_EQ(result.at(1), 2);
EXPECT_EQ(result.at(2), 5);
EXPECT_EQ(result.at(3), 9);
}
TEST(CkTileSequence, PrefixSumSingle)
{
// <5> -> <0, 5>
constexpr auto result = prefix_sum_sequence(sequence<5>{});
EXPECT_EQ(decltype(result)::size(), 2);
EXPECT_EQ(result.at(0), 0);
EXPECT_EQ(result.at(1), 5);
}
TEST(CkTileSequence, PrefixSumEmpty)
{
// <> -> <0>
constexpr auto result = prefix_sum_sequence(sequence<>{});
EXPECT_EQ(decltype(result)::size(), 1);
EXPECT_EQ(result.at(0), 0);
}