Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
7b4de775
Commit
7b4de775
authored
May 04, 2022
by
Chao Liu
Browse files
use remove_cvref_t
parent
e86f3769
Changes
10
Hide whitespace changes
Inline
Side-by-side
Showing
10 changed files
with
22 additions
and
22 deletions
+22
-22
include/ck/config.hpp
include/ck/config.hpp
+1
-1
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp
...ration/gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp
+2
-2
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp
...ration/gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp
+2
-2
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp
...ration/gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp
+2
-2
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp
...ration/gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp
+3
-3
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp
...ration/gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp
+4
-4
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp
...ion/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp
+2
-2
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp
...ion/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp
+2
-2
include/ck/tensor_operation/gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp
...peration/gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp
+1
-1
include/ck/utility/tuple.hpp
include/ck/utility/tuple.hpp
+3
-3
No files found.
include/ck/config.hpp
View file @
7b4de775
...
...
@@ -26,7 +26,7 @@
#endif
#endif
// buffer resour
s
e
// buffer resour
c
e
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_BUFFER_RESOURCE_3RD_DWORD -1
#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx908__) || \
...
...
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v4r1.hpp
View file @
7b4de775
...
...
@@ -56,8 +56,8 @@ struct BlockwiseTensorSliceTransfer_v4r1
dst_element_op
)
{
static_assert
(
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
SrcDesc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
DstDesc
>
>
::
GetNumOfDimension
()
&&
static_assert
(
nDim
==
remove_
cv
ref_t
<
SrcDesc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
DstDesc
>::
GetNumOfDimension
()
&&
nDim
==
ThreadClusterLengths
::
Size
()
&&
nDim
==
ThreadClusterArrangeOrder
::
Size
()
&&
nDim
==
SrcDimAccessOrder
::
Size
()
&&
nDim
==
DstDimAccessOrder
::
Size
(),
...
...
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v5r1.hpp
View file @
7b4de775
...
...
@@ -45,8 +45,8 @@ struct BlockwiseTensorSliceTransfer_v5r1
src_desc
,
make_zero_multi_index
<
nDim
>
(),
dst_desc
,
make_zero_multi_index
<
nDim
>
())
{
static_assert
(
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
SrcDesc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
DstDesc
>
>
::
GetNumOfDimension
()
&&
static_assert
(
nDim
==
remove_
cv
ref_t
<
SrcDesc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
DstDesc
>::
GetNumOfDimension
()
&&
nDim
==
BlockSliceLengths
::
Size
()
&&
nDim
==
ThreadSliceLengths
::
Size
()
&&
nDim
==
ThreadClusterLengths
::
Size
()
&&
nDim
==
ThreadClusterArrangeOrder
::
Size
()
&&
...
...
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r1.hpp
View file @
7b4de775
...
...
@@ -48,8 +48,8 @@ struct BlockwiseTensorSliceTransfer_v6r1
element_op
)
{
static_assert
(
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
SrcDesc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
DstDesc
>
>
::
GetNumOfDimension
()
&&
static_assert
(
nDim
==
remove_
cv
ref_t
<
SrcDesc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
DstDesc
>::
GetNumOfDimension
()
&&
nDim
==
ThreadClusterLengths
::
Size
()
&&
nDim
==
ThreadClusterArrangeOrder
::
Size
()
&&
nDim
==
DimAccessOrder
::
Size
(),
...
...
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r2.hpp
View file @
7b4de775
...
...
@@ -55,9 +55,9 @@ struct BlockwiseTensorSliceTransfer_v6r2
element_op
)
{
static_assert
(
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
Src0Desc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
Src1Desc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
DstDesc
>
>
::
GetNumOfDimension
()
&&
static_assert
(
nDim
==
remove_
cv
ref_t
<
Src0Desc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
Src1Desc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
DstDesc
>::
GetNumOfDimension
()
&&
nDim
==
ThreadClusterLengths
::
Size
()
&&
nDim
==
ThreadClusterArrangeOrder
::
Size
()
&&
nDim
==
DimAccessOrder
::
Size
(),
...
...
include/ck/tensor_operation/gpu/block/blockwise_tensor_slice_transfer_v6r3.hpp
View file @
7b4de775
...
...
@@ -62,10 +62,10 @@ struct BlockwiseTensorSliceTransfer_v6r3
element_op
)
{
static_assert
(
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
Src0Desc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
Src1Desc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
Src2Desc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
DstDesc
>
>
::
GetNumOfDimension
()
&&
static_assert
(
nDim
==
remove_
cv
ref_t
<
Src0Desc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
Src1Desc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
Src2Desc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
DstDesc
>::
GetNumOfDimension
()
&&
nDim
==
ThreadClusterLengths
::
Size
()
&&
nDim
==
ThreadClusterArrangeOrder
::
Size
()
&&
nDim
==
DimAccessOrder
::
Size
(),
...
...
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp
View file @
7b4de775
...
...
@@ -56,8 +56,8 @@ struct ThreadGroupTensorSliceTransfer_v4r1
dst_element_op
)
{
static_assert
(
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
SrcDesc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
DstDesc
>
>
::
GetNumOfDimension
()
&&
static_assert
(
nDim
==
remove_
cv
ref_t
<
SrcDesc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
DstDesc
>::
GetNumOfDimension
()
&&
nDim
==
ThreadClusterLengths
::
Size
()
&&
nDim
==
ThreadClusterArrangeOrder
::
Size
()
&&
nDim
==
SrcDimAccessOrder
::
Size
()
&&
nDim
==
DstDimAccessOrder
::
Size
(),
...
...
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp
View file @
7b4de775
...
...
@@ -46,8 +46,8 @@ struct ThreadGroupTensorSliceTransfer_v6r1
element_op
)
{
static_assert
(
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
SrcDesc
>
>
::
GetNumOfDimension
()
&&
nDim
==
remove_ref
erence_t
<
remove_cv
_t
<
DstDesc
>
>
::
GetNumOfDimension
()
&&
static_assert
(
nDim
==
remove_
cv
ref_t
<
SrcDesc
>::
GetNumOfDimension
()
&&
nDim
==
remove_
cv
ref_t
<
DstDesc
>::
GetNumOfDimension
()
&&
nDim
==
ThreadClusterLengths
::
Size
()
&&
nDim
==
ThreadClusterArrangeOrder
::
Size
()
&&
nDim
==
DimAccessOrder
::
Size
(),
...
...
include/ck/tensor_operation/gpu/device/device_gemm_xdl_c_shuffle_bias_2d.hpp
View file @
7b4de775
...
...
@@ -289,7 +289,7 @@ struct DeviceGemmXdl_C_Shuffle_Bias_2d
arg
.
N01_
))
{
throw
std
::
runtime_error
(
"wrong! GridwiseGemm_k
m_kn_m0m1n0n1
_xdlops_v3r2 has invalid setting"
);
"wrong! GridwiseGemm_k
0mk1_k0nk1_mn
_xdlops_v3r2 has invalid setting"
);
}
const
index_t
grid_size
=
GridwiseGemm
::
CalculateGridSize
(
arg
.
c_grid_desc_m_n_
);
...
...
include/ck/utility/tuple.hpp
View file @
7b4de775
...
...
@@ -22,7 +22,7 @@ struct TupleElement
__host__
__device__
constexpr
TupleElement
()
=
default
;
template
<
typename
T
,
typename
enable_if
<!
is_same
<
remove_ref
erence_t
<
remove_cv
_t
<
T
>
>
,
TupleElement
>::
value
,
typename
enable_if
<!
is_same
<
remove_
cv
ref_t
<
T
>,
TupleElement
>::
value
,
bool
>::
type
=
false
>
__host__
__device__
constexpr
TupleElement
(
T
&&
v
)
:
mData
(
std
::
forward
<
T
>
(
v
))
{
...
...
@@ -60,7 +60,7 @@ struct TupleImpl<Sequence<Is...>, Xs...> : TupleElement<TupleElementKey<Is>, Xs>
template
<
typename
Y
,
typename
enable_if
<
sizeof
...(
Is
)
==
1
&&
sizeof
...(
Xs
)
==
1
&&
!
is_same
<
remove_ref
erence_t
<
remove_cv
_t
<
Y
>
>
,
TupleImpl
>::
value
,
!
is_same
<
remove_
cv
ref_t
<
Y
>,
TupleImpl
>::
value
,
bool
>::
type
=
false
>
__host__
__device__
constexpr
TupleImpl
(
Y
&&
y
)
:
TupleElement
<
TupleElementKey
<
Is
>
,
Xs
>
(
std
::
forward
<
Y
>
(
y
))...
...
...
@@ -102,7 +102,7 @@ struct Tuple : detail::TupleImpl<typename arithmetic_sequence_gen<0, sizeof...(X
template
<
typename
Y
,
typename
enable_if
<
sizeof
...(
Xs
)
==
1
&&
!
is_same
<
remove_ref
erence_t
<
remove_cv
_t
<
Y
>
>
,
Tuple
>::
value
,
!
is_same
<
remove_
cv
ref_t
<
Y
>,
Tuple
>::
value
,
bool
>::
type
=
false
>
__host__
__device__
constexpr
Tuple
(
Y
&&
y
)
:
base
(
std
::
forward
<
Y
>
(
y
))
{
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment