diff --git a/include/ck_tile/core.hpp b/include/ck_tile/core.hpp index 24f84dc813..e4b0ac02a7 100644 --- a/include/ck_tile/core.hpp +++ b/include/ck_tile/core.hpp @@ -39,9 +39,9 @@ #include "ck_tile/core/container/multi_index.hpp" #include "ck_tile/core/container/sequence.hpp" #include "ck_tile/core/container/span.hpp" +#include "ck_tile/core/container/static_array.hpp" #include "ck_tile/core/container/statically_indexed_array.hpp" #include "ck_tile/core/container/thread_buffer.hpp" -#include "ck_tile/core/container/trivial_array.hpp" #include "ck_tile/core/container/tuple.hpp" #include "ck_tile/core/numeric/bfloat16.hpp" #include "ck_tile/core/numeric/e8m0.hpp" diff --git a/include/ck_tile/core/container/sequence.hpp b/include/ck_tile/core/container/sequence.hpp index ee4805234c..14c60da293 100644 --- a/include/ck_tile/core/container/sequence.hpp +++ b/include/ck_tile/core/container/sequence.hpp @@ -4,7 +4,7 @@ #pragma once #include "ck_tile/core/config.hpp" -#include "ck_tile/core/container/trivial_array.hpp" +#include "ck_tile/core/container/static_array.hpp" #include "ck_tile/core/numeric/integer.hpp" #include "ck_tile/core/numeric/integral_constant.hpp" #include "ck_tile/core/numeric/math.hpp" @@ -353,8 +353,8 @@ struct sequence_inclusive_scan_impl, Reduce, Init, Reverse> else { constexpr auto arr = []() { - trivial_array values = {Is...}; - trivial_array result = {0}; + static_array values = {Is...}; + static_array result = {0}; if constexpr(Reverse) { // Reverse scan: right to left diff --git a/include/ck_tile/core/container/static_array.hpp b/include/ck_tile/core/container/static_array.hpp new file mode 100644 index 0000000000..abc6bc8615 --- /dev/null +++ b/include/ck_tile/core/container/static_array.hpp @@ -0,0 +1,30 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once +#include "ck_tile/core/numeric/integer.hpp" + +namespace ck_tile { +// Fixed-size array with aggregate initialization +// +// This is a minimal array type designed for: +// - Constexpr/compile-time computation +// - GPU kernel code (trivially copyable) +// - Template metaprogramming +// +// Unlike ck_tile::array, this has no custom constructors, +// making it a literal type suitable for constexpr contexts. +// Use aggregate initialization: static_array arr{1, 2, 3}; +template +struct static_array +{ + // Public aggregate initialization makes this a literal type + T elems[N]; + + // Basic constexpr accessors + constexpr const T& operator[](index_t i) const { return elems[i]; } + constexpr T& operator[](index_t i) { return elems[i]; } + + constexpr static index_t size() { return N; } +}; +} // namespace ck_tile diff --git a/include/ck_tile/core/container/trivial_array.hpp b/include/ck_tile/core/container/trivial_array.hpp deleted file mode 100644 index 0223b20145..0000000000 --- a/include/ck_tile/core/container/trivial_array.hpp +++ /dev/null @@ -1,21 +0,0 @@ -// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT - -#pragma once -#include "ck_tile/core/numeric/integer.hpp" - -namespace ck_tile { -// trivial_array has no custom constructor, allowing it to be used as a constexpr variable type -template -struct trivial_array -{ - // Public aggregate initialization makes this a literal type - T data[N]; - - // Basic constexpr accessors - constexpr const T& operator[](index_t i) const { return data[i]; } - constexpr T& operator[](index_t i) { return data[i]; } - - constexpr static index_t size() { return N; } -}; -} // namespace ck_tile