From b1dbf64c91f711f712818a0a7ec7c7c813a14fe3 Mon Sep 17 00:00:00 2001 From: "Po-Yen, Chen" Date: Wed, 13 Mar 2024 03:55:07 -0400 Subject: [PATCH] Some minor changes --- include/ck_tile/core/config.hpp | 8 ----- include/ck_tile/core/container/array.hpp | 35 +++++-------------- include/ck_tile/core/numeric/vector_type.hpp | 3 +- .../ck_tile/core/tensor/tensor_adaptor.hpp | 4 +-- .../core/tensor/tensor_adaptor_coordinate.hpp | 2 +- include/ck_tile/core/utility/random.hpp | 8 +---- include/ck_tile/ops/epilogue.hpp | 1 - 7 files changed, 14 insertions(+), 47 deletions(-) diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index 0ef51a5552..965281d98b 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -35,14 +35,6 @@ #define CK_TILE_FLOAT_TO_FP8_DEFAULT CK_TILE_FLOAT_TO_FP8_STANDARD #endif -#ifndef STATIC_ASSERT -#ifndef NDEBUG -#define STATIC_ASSERT(...) static_assert(__VA_ARGS__) -#else -#define STATIC_ASSERT(...) -#endif -#endif // #ifndef STATIC_ASSERT - // in the old rocm period, we have to use tuple array implementation to implement this // so turn on the _USE_TUPLE if meet compiler error, otherwise _USE_ARRAY by default. #define CK_TILE_STATICALLY_INDEXED_ARRAY_USE_ARRAY 0 diff --git a/include/ck_tile/core/container/array.hpp b/include/ck_tile/core/container/array.hpp index a75af1f1ec..c272b01f54 100644 --- a/include/ck_tile/core/container/array.hpp +++ b/include/ck_tile/core/container/array.hpp @@ -3,6 +3,8 @@ #pragma once +#include + #include "ck_tile/core/config.hpp" #include "ck_tile/core/numeric/integer.hpp" #include "ck_tile/core/numeric/integral_constant.hpp" @@ -96,13 +98,13 @@ struct array CK_TILE_HOST_DEVICE constexpr value_type& operator[](index_t i) { return get(i); } CK_TILE_HOST_DEVICE constexpr value_type& operator()(index_t i) { return get(i); } // TODO: compatible #if 0 - template - CK_TILE_HOST_DEVICE constexpr auto operator=(const ArrayType& a) + template + CK_TILE_HOST_DEVICE constexpr auto operator=(const ArrayLike& arr) { - static_assert(ArrayType::size() == size(), "wrong! size not the same"); + static_assert(ArrayLike::size() == size(), "wrong! size not the same"); for(index_t i = 0; i < size(); ++i) { - data[i] = a[i]; + data[i] = arr[i]; } return *this; } @@ -205,31 +207,10 @@ CK_TILE_HOST_DEVICE constexpr details::return_type make_array(Ts&&... // compatible with old ck's initializer, make an array and fill it withe the last element from // initializer_list -#include template CK_TILE_HOST_DEVICE constexpr auto make_array_with(std::initializer_list ilist) { - constexpr index_t list_size = std::initializer_list{}.size(); - - static_assert(list_size <= Size, "out of bound"); - - index_t i = 0; - T vlast = T{}; - array arr; - - for(const T& val : ilist) - { - arr.data[i] = val; - vlast = val; - ++i; - } - - for(; i < Size; ++i) - { - arr.data[i] = vlast; - } - - return arr; + return array(ilist); } template @@ -258,7 +239,7 @@ CK_TILE_HOST_DEVICE constexpr bool operator!=(const array& a, const arr template CK_TILE_HOST_DEVICE constexpr auto to_array(const X& x) { - STATIC_ASSERT(N <= X::size(), ""); + static_assert(N <= X::size(), ""); array arr; diff --git a/include/ck_tile/core/numeric/vector_type.hpp b/include/ck_tile/core/numeric/vector_type.hpp index 0449a085ab..4d02937992 100644 --- a/include/ck_tile/core/numeric/vector_type.hpp +++ b/include/ck_tile/core/numeric/vector_type.hpp @@ -24,7 +24,8 @@ struct ext_vector { static constexpr index_t N = N_; using value_type = T_; - using type = value_type __attribute__((ext_vector_type(N))); // this is danguous + static_assert(!std::is_class_v); + using type = value_type __attribute__((ext_vector_type(N))); // this is danguous }; } // namespace impl diff --git a/include/ck_tile/core/tensor/tensor_adaptor.hpp b/include/ck_tile/core/tensor/tensor_adaptor.hpp index 8a5c7218c6..d8dc4e5b53 100644 --- a/include/ck_tile/core/tensor/tensor_adaptor.hpp +++ b/include/ck_tile/core/tensor/tensor_adaptor.hpp @@ -725,7 +725,7 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. constexpr auto num_low_dim = encoded_transforms[i].template at<2>(); \ constexpr auto num_up_dim = encoded_transforms[i].template at<4>(); \ \ - STATIC_ASSERT(name == coord_transform_enum::pass_through || \ + static_assert(name == coord_transform_enum::pass_through || \ name == coord_transform_enum::pad || \ name == coord_transform_enum::embed || \ name == coord_transform_enum::merge || \ @@ -849,7 +849,7 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. constexpr auto num_low_dim = encoded_transforms[i].template at<2>(); \ constexpr auto num_up_dim = encoded_transforms[i].template at<4>(); \ \ - STATIC_ASSERT(name == coord_transform_enum::pass_through || \ + static_assert(name == coord_transform_enum::pass_through || \ name == coord_transform_enum::pad || \ name == coord_transform_enum::embed || \ name == coord_transform_enum::merge || \ diff --git a/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp b/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp index c4528fbc4b..0d398d4237 100644 --- a/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp +++ b/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp @@ -104,7 +104,7 @@ CK_TILE_HOST_DEVICE constexpr void move_tensor_adaptor_coordinate(const Adaptor& // constexpr index_t ndim_bottom = Adaptor::get_num_of_bottom_dimension(); constexpr index_t ntransform = Adaptor::get_num_of_transform(); - // STATIC_ASSERT(TopIndex::size() == ndim_top && BottomIndex::size() == ndim_bottom, ""); + // static_assert(TopIndex::size() == ndim_top && BottomIndex::size() == ndim_bottom, ""); // judge whether calculation of lower diff is needed for each transform // use index_t for boolean type diff --git a/include/ck_tile/core/utility/random.hpp b/include/ck_tile/core/utility/random.hpp index e7dc34e6d5..f7fbfad4dd 100644 --- a/include/ck_tile/core/utility/random.hpp +++ b/include/ck_tile/core/utility/random.hpp @@ -15,13 +15,7 @@ namespace ck_tile { template struct prand_generator_t { - CK_TILE_HOST_DEVICE uint32_t operator()(int id, T val, uint32_t seed = seed_) - { - std::ignore = id; - std::ignore = val; - std::ignore = seed; - return 0; - } + CK_TILE_HOST_DEVICE uint32_t operator()(int, T, uint32_t = seed_) { return 0; } }; // version for fp32 diff --git a/include/ck_tile/ops/epilogue.hpp b/include/ck_tile/ops/epilogue.hpp index 388f52c898..8e0d3e449b 100644 --- a/include/ck_tile/ops/epilogue.hpp +++ b/include/ck_tile/ops/epilogue.hpp @@ -4,4 +4,3 @@ #pragma once #include "ck_tile/ops/epilogue/default_2d_epilogue.hpp" -#include "ck_tile/ops/common/tensor_layout.hpp"