ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-12 17:26:00 +00:00
Code Issues Packages Projects Releases Wiki Activity

Some minor changes

Browse Source
This commit is contained in:
Po-Yen, Chen
2024-03-13 03:55:07 -04:00
parent 8d1631adc9
commit b1dbf64c91
7 changed files with 14 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
include/ck_tile/core/config.hpp
View File

@@ -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

35
include/ck_tile/core/container/array.hpp
View File

@@ -3,6 +3,8 @@
#pragma once
#include <initializer_list>
#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 <typename ArrayType>
CK_TILE_HOST_DEVICE constexpr auto operator=(const ArrayType& a)
template <typename ArrayLike>
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<D, Ts...> make_array(Ts&&...
// compatible with old ck's initializer, make an array and fill it withe the last element from
// initializer_list
#include <initializer_list>
template <typename T, index_t Size>
CK_TILE_HOST_DEVICE constexpr auto make_array_with(std::initializer_list<T> ilist)
{
constexpr index_t list_size = std::initializer_list<T>{}.size();
static_assert(list_size <= Size, "out of bound");
index_t i = 0;
T vlast = T{};
array<T, Size> 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<T, Size>(ilist);
}
template <typename T, index_t Size>
@@ -258,7 +239,7 @@ CK_TILE_HOST_DEVICE constexpr bool operator!=(const array<T, Size>& a, const arr
template <typename T, index_t N, typename X>
CK_TILE_HOST_DEVICE constexpr auto to_array(const X& x)
{
STATIC_ASSERT(N <= X::size(), "");
static_assert(N <= X::size(), "");
array<T, N> arr;

3
include/ck_tile/core/numeric/vector_type.hpp
View File

@@ -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<value_type>);
using type = value_type __attribute__((ext_vector_type(N))); // this is danguous
};
} // namespace impl

4
include/ck_tile/core/tensor/tensor_adaptor.hpp
View File

@@ -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 || \

2
include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp
View File

@@ -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

8
include/ck_tile/core/utility/random.hpp
View File

@@ -15,13 +15,7 @@ namespace ck_tile {
template <typename T, uint32_t seed_>
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

1
include/ck_tile/ops/epilogue.hpp
View File

@@ -4,4 +4,3 @@
#pragma once
#include "ck_tile/ops/epilogue/default_2d_epilogue.hpp"
#include "ck_tile/ops/common/tensor_layout.hpp"