Skip to content

GitLab

Commit 6cd94d98 authored by Chao Liu's avatar Chao Liu
Browse files

refactoring array

parent 55d461a0
Show whitespace changes
Inline Side-by-side
Showing with 31 additions and 31 deletions
composable_kernel/include/tensor_description/dynamic_multi_index_transform.hpp
View file @ 6cd94d98
...@@ -211,8 +211,8 @@ struct DynamicEmbed ...@@ -211,8 +211,8 @@ struct DynamicEmbed
} }
__host__ __device__ explicit constexpr DynamicEmbed() __host__ __device__ explicit constexpr DynamicEmbed()
: up_lengths_{make_zero_array<index_t, NDimUp>()}, : up_lengths_{make_zero_multi_index<NDimUp>()},
coefficients_{make_zero_array<index_t, NDimUp>()} coefficients_{make_zero_multi_index<NDimUp>()}
{ {
} }
...@@ -288,8 +288,8 @@ struct DynamicMerge ...@@ -288,8 +288,8 @@ struct DynamicMerge
} }
__host__ __device__ explicit constexpr DynamicMerge() __host__ __device__ explicit constexpr DynamicMerge()
: low_lengths_{make_zero_array<index_t, NDimLow>()}, : low_lengths_{make_zero_multi_index<NDimLow>()},
low_lengths_scan_{make_zero_array<index_t, NDimLow>()}, low_lengths_scan_{make_zero_multi_index<NDimLow>()},
up_lengths_{0} up_lengths_{0}
{ {
} }
...@@ -429,8 +429,8 @@ struct DynamicUnMerge ...@@ -429,8 +429,8 @@ struct DynamicUnMerge
} }
__host__ __device__ explicit constexpr DynamicUnMerge() __host__ __device__ explicit constexpr DynamicUnMerge()
: up_lengths_{make_zero_array<index_t, NDimUp>()}, : up_lengths_{make_zero_multi_index<NDimUp>()},
up_lengths_scan_{make_zero_array<index_t, NDimUp>()} up_lengths_scan_{make_zero_multi_index<NDimUp>()}
{ {
} }
......
composable_kernel/include/tensor_description/multi_index.hpp
View file @ 6cd94d98
...@@ -27,9 +27,9 @@ template <index_t N> ...@@ -27,9 +27,9 @@ template <index_t N>
using MultiIndex = StaticallyIndexedArray<index_t, N>; using MultiIndex = StaticallyIndexedArray<index_t, N>;
template <typename... Xs> template <typename... Xs>
__host__ __device__ constexpr auto make_multi_index(Xs... xs) __host__ __device__ constexpr auto make_multi_index(const Xs&... xs)
{ {
return make_statically_indexed_array<index_t>(xs...); return make_statically_indexed_array<const index_t>(std::forward<const Xs>(xs)...);
} }
#endif #endif
......
composable_kernel/include/utility/statically_indexed_array.hpp
View file @ 6cd94d98
...@@ -29,7 +29,7 @@ struct StaticallyIndexedArray<T, 1> : public Tuple<T> ...@@ -29,7 +29,7 @@ struct StaticallyIndexedArray<T, 1> : public Tuple<T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -42,7 +42,7 @@ struct StaticallyIndexedArray<T, 2> : public Tuple<T, T> ...@@ -42,7 +42,7 @@ struct StaticallyIndexedArray<T, 2> : public Tuple<T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -55,7 +55,7 @@ struct StaticallyIndexedArray<T, 3> : public Tuple<T, T, T> ...@@ -55,7 +55,7 @@ struct StaticallyIndexedArray<T, 3> : public Tuple<T, T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -68,7 +68,7 @@ struct StaticallyIndexedArray<T, 4> : public Tuple<T, T, T, T> ...@@ -68,7 +68,7 @@ struct StaticallyIndexedArray<T, 4> : public Tuple<T, T, T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -81,7 +81,7 @@ struct StaticallyIndexedArray<T, 5> : public Tuple<T, T, T, T, T> ...@@ -81,7 +81,7 @@ struct StaticallyIndexedArray<T, 5> : public Tuple<T, T, T, T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -94,7 +94,7 @@ struct StaticallyIndexedArray<T, 6> : public Tuple<T, T, T, T, T, T> ...@@ -94,7 +94,7 @@ struct StaticallyIndexedArray<T, 6> : public Tuple<T, T, T, T, T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -107,7 +107,7 @@ struct StaticallyIndexedArray<T, 7> : public Tuple<T, T, T, T, T, T, T> ...@@ -107,7 +107,7 @@ struct StaticallyIndexedArray<T, 7> : public Tuple<T, T, T, T, T, T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -120,7 +120,7 @@ struct StaticallyIndexedArray<T, 8> : public Tuple<T, T, T, T, T, T, T, T> ...@@ -120,7 +120,7 @@ struct StaticallyIndexedArray<T, 8> : public Tuple<T, T, T, T, T, T, T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -133,7 +133,7 @@ struct StaticallyIndexedArray<T, 9> : public Tuple<T, T, T, T, T, T, T, T, T> ...@@ -133,7 +133,7 @@ struct StaticallyIndexedArray<T, 9> : public Tuple<T, T, T, T, T, T, T, T, T>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -146,7 +146,7 @@ struct StaticallyIndexedArray<T, 10> : public Tuple<T, T, T, T, T, T, T, T, T, T ...@@ -146,7 +146,7 @@ struct StaticallyIndexedArray<T, 10> : public Tuple<T, T, T, T, T, T, T, T, T, T
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -159,7 +159,7 @@ struct StaticallyIndexedArray<T, 11> : public Tuple<T, T, T, T, T, T, T, T, T, T ...@@ -159,7 +159,7 @@ struct StaticallyIndexedArray<T, 11> : public Tuple<T, T, T, T, T, T, T, T, T, T
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -172,7 +172,7 @@ struct StaticallyIndexedArray<T, 12> : public Tuple<T, T, T, T, T, T, T, T, T, T ...@@ -172,7 +172,7 @@ struct StaticallyIndexedArray<T, 12> : public Tuple<T, T, T, T, T, T, T, T, T, T
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -185,7 +185,7 @@ struct StaticallyIndexedArray<T, 13> : public Tuple<T, T, T, T, T, T, T, T, T, T ...@@ -185,7 +185,7 @@ struct StaticallyIndexedArray<T, 13> : public Tuple<T, T, T, T, T, T, T, T, T, T
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -198,7 +198,7 @@ struct StaticallyIndexedArray<T, 14> : public Tuple<T, T, T, T, T, T, T, T, T, T ...@@ -198,7 +198,7 @@ struct StaticallyIndexedArray<T, 14> : public Tuple<T, T, T, T, T, T, T, T, T, T
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -211,7 +211,7 @@ struct StaticallyIndexedArray<T, 15> : public Tuple<T, T, T, T, T, T, T, T, T, T ...@@ -211,7 +211,7 @@ struct StaticallyIndexedArray<T, 15> : public Tuple<T, T, T, T, T, T, T, T, T, T
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -224,7 +224,7 @@ struct StaticallyIndexedArray<T, 16> : public Tuple<T, T, T, T, T, T, T, T, T, T ...@@ -224,7 +224,7 @@ struct StaticallyIndexedArray<T, 16> : public Tuple<T, T, T, T, T, T, T, T, T, T
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -238,7 +238,7 @@ struct StaticallyIndexedArray<T, 17> ...@@ -238,7 +238,7 @@ struct StaticallyIndexedArray<T, 17>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -252,7 +252,7 @@ struct StaticallyIndexedArray<T, 18> ...@@ -252,7 +252,7 @@ struct StaticallyIndexedArray<T, 18>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -266,7 +266,7 @@ struct StaticallyIndexedArray<T, 19> ...@@ -266,7 +266,7 @@ struct StaticallyIndexedArray<T, 19>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -280,7 +280,7 @@ struct StaticallyIndexedArray<T, 20> ...@@ -280,7 +280,7 @@ struct StaticallyIndexedArray<T, 20>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -294,7 +294,7 @@ struct StaticallyIndexedArray<T, 21> ...@@ -294,7 +294,7 @@ struct StaticallyIndexedArray<T, 21>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
...@@ -308,7 +308,7 @@ struct StaticallyIndexedArray<T, 22> ...@@ -308,7 +308,7 @@ struct StaticallyIndexedArray<T, 22>
template <typename... Ys> template <typename... Ys>
__host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys) __host__ __device__ explicit constexpr StaticallyIndexedArray(Ys&&... ys)
: base(static_cast<T&&>(ys)...) : base(std::forward<Ys>(ys)...)
{ {
} }
}; };
......
driver/src/conv_driver.cpp
View file @ 6cd94d98
...@@ -549,7 +549,7 @@ int main(int argc, char* argv[]) ...@@ -549,7 +549,7 @@ int main(int argc, char* argv[])
#endif #endif
} }
#if 1 #if 0
device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(in_nchw_desc, device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(in_nchw_desc,
in_nchw, in_nchw,
wei_kcyx_desc, wei_kcyx_desc,
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment