From 7b4de775704d8de67c6785a852127390b93aa5a8 Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Wed, 4 May 2022 03:00:16 +0000 Subject: [PATCH] use remove_cvref_t --- include/ck/config.hpp | 2 +- .../gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp | 4 ++-- .../gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp | 4 ++-- .../gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp | 4 ++-- .../gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp | 6 +++--- .../gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp | 8 ++++---- .../gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp | 4 ++-- .../gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp | 4 ++-- .../gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp | 2 +- include/ck/utility/tuple.hpp | 6 +++--- 10 files changed, 22 insertions(+), 22 deletions(-) diff --git a/include/ck/config.hpp b/include/ck/config.hpp index 22e2ac943e..6b0710a795 100644 --- a/include/ck/config.hpp +++ b/include/ck/config.hpp @@ -26,7 +26,7 @@ #endif #endif -// buffer resourse +// buffer resource #ifndef __HIP_DEVICE_COMPILE__ // for host code #define CK_BUFFER_RESOURCE_3RD_DWORD -1 #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx908__) || \ diff --git a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp index 2358cdc98b..8306fa93ff 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp @@ -56,8 +56,8 @@ struct BlockwiseTensorSliceTransfer_v4r1 dst_element_op) { - static_assert(nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && + static_assert(nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && nDim == ThreadClusterLengths::Size() && nDim == ThreadClusterArrangeOrder::Size() && nDim == SrcDimAccessOrder::Size() && nDim == DstDimAccessOrder::Size(), diff --git a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp index acd99132cc..93fe5da723 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp @@ -45,8 +45,8 @@ struct BlockwiseTensorSliceTransfer_v5r1 src_desc, make_zero_multi_index(), dst_desc, make_zero_multi_index()) { - static_assert(nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && + static_assert(nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && nDim == BlockSliceLengths::Size() && nDim == ThreadSliceLengths::Size() && nDim == ThreadClusterLengths::Size() && nDim == ThreadClusterArrangeOrder::Size() && diff --git a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp index 957c8f522c..3bf86d6785 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp @@ -48,8 +48,8 @@ struct BlockwiseTensorSliceTransfer_v6r1 element_op) { - static_assert(nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && + static_assert(nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && nDim == ThreadClusterLengths::Size() && nDim == ThreadClusterArrangeOrder::Size() && nDim == DimAccessOrder::Size(), diff --git a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp index 2e06214b8c..575a015802 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp @@ -55,9 +55,9 @@ struct BlockwiseTensorSliceTransfer_v6r2 element_op) { - static_assert(nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && + static_assert(nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && nDim == ThreadClusterLengths::Size() && nDim == ThreadClusterArrangeOrder::Size() && nDim == DimAccessOrder::Size(), diff --git a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp index 085981736b..4a1d82000a 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp @@ -62,10 +62,10 @@ struct BlockwiseTensorSliceTransfer_v6r3 element_op) { - static_assert(nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && + static_assert(nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && nDim == ThreadClusterLengths::Size() && nDim == ThreadClusterArrangeOrder::Size() && nDim == DimAccessOrder::Size(), diff --git a/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp b/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp index 8a6748eaf2..cbabbaf47d 100644 --- a/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp +++ b/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp @@ -56,8 +56,8 @@ struct ThreadGroupTensorSliceTransfer_v4r1 dst_element_op) { - static_assert(nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && + static_assert(nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && nDim == ThreadClusterLengths::Size() && nDim == ThreadClusterArrangeOrder::Size() && nDim == SrcDimAccessOrder::Size() && nDim == DstDimAccessOrder::Size(), diff --git a/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp b/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp index cd6881dabe..0baebc5e3b 100644 --- a/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp +++ b/include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp @@ -46,8 +46,8 @@ struct ThreadGroupTensorSliceTransfer_v6r1 element_op) { - static_assert(nDim == remove_reference_t>::GetNumOfDimension() && - nDim == remove_reference_t>::GetNumOfDimension() && + static_assert(nDim == remove_cvref_t::GetNumOfDimension() && + nDim == remove_cvref_t::GetNumOfDimension() && nDim == ThreadClusterLengths::Size() && nDim == ThreadClusterArrangeOrder::Size() && nDim == DimAccessOrder::Size(), diff --git a/include/ck/tensor_operation/gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp b/include/ck/tensor_operation/gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp index 70b1b0fe22..4010965312 100644 --- a/include/ck/tensor_operation/gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp +++ b/include/ck/tensor_operation/gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp @@ -289,7 +289,7 @@ struct DeviceGemmXdl_C_Shuffle_Bias_2d arg.N01_)) { throw std::runtime_error( - "wrong! GridwiseGemm_km_kn_m0m1n0n1_xdlops_v3r2 has invalid setting"); + "wrong! GridwiseGemm_k0mk1_k0nk1_mn_xdlops_v3r2 has invalid setting"); } const index_t grid_size = GridwiseGemm::CalculateGridSize(arg.c_grid_desc_m_n_); diff --git a/include/ck/utility/tuple.hpp b/include/ck/utility/tuple.hpp index 96cab4b99e..9fa77d1932 100644 --- a/include/ck/utility/tuple.hpp +++ b/include/ck/utility/tuple.hpp @@ -22,7 +22,7 @@ struct TupleElement __host__ __device__ constexpr TupleElement() = default; template >, TupleElement>::value, + typename enable_if, TupleElement>::value, bool>::type = false> __host__ __device__ constexpr TupleElement(T&& v) : mData(std::forward(v)) { @@ -60,7 +60,7 @@ struct TupleImpl, Xs...> : TupleElement, Xs> template >, TupleImpl>::value, + !is_same, TupleImpl>::value, bool>::type = false> __host__ __device__ constexpr TupleImpl(Y&& y) : TupleElement, Xs>(std::forward(y))... @@ -102,7 +102,7 @@ struct Tuple : detail::TupleImpl>, Tuple>::value, + !is_same, Tuple>::value, bool>::type = false> __host__ __device__ constexpr Tuple(Y&& y) : base(std::forward(y)) {