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
4a5b2257
Commit
4a5b2257
authored
Mar 08, 2023
by
Rosty Geyyer
Browse files
Format
parent
3e002b60
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
8 additions
and
6 deletions
+8
-6
include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle.hpp
..._batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle.hpp
+8
-6
No files found.
include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle.hpp
View file @
4a5b2257
...
@@ -789,13 +789,13 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
...
@@ -789,13 +789,13 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
}
}
// check if DsLayout is supported
// check if DsLayout is supported
template
<
typename
RefLayout
,
typename
DsLayout
,
const
index_t
NumDTensor
>
template
<
typename
RefLayout
,
typename
DsLayout
,
const
index_t
NumDTensor
>
static
bool
CheckDLayout
()
static
bool
CheckDLayout
()
{
{
static
bool
valid
=
true
;
static
bool
valid
=
true
;
// iterate over DLayout tuple
// iterate over DLayout tuple
static_for
<
0
,
NumDTensor
,
1
>
{}([
&
](
auto
i
)
{
static_for
<
0
,
NumDTensor
,
1
>
{}([
&
](
auto
i
)
{
using
DLayout
=
remove_cvref_t
<
tuple_element_t
<
i
.
value
,
DsLayout
>>
;
using
DLayout
=
remove_cvref_t
<
tuple_element_t
<
i
.
value
,
DsLayout
>>
;
// if RefLayout and DLayout are same, keep valid true, otherwise false
// if RefLayout and DLayout are same, keep valid true, otherwise false
valid
=
valid
&&
is_same_v
<
RefLayout
,
DLayout
>
;
valid
=
valid
&&
is_same_v
<
RefLayout
,
DLayout
>
;
});
});
...
@@ -816,12 +816,14 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
...
@@ -816,12 +816,14 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
// B1 - Row or Col
// B1 - Row or Col
// D1s - Rows
// D1s - Rows
// E1 - Row
// E1 - Row
if
(
!
(
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
A0Layout
>
&&
if
(
!
(
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
A0Layout
>
&&
is_same_v
<
tensor_layout
::
gemm
::
ColumnMajor
,
B0Layout
>
&&
is_same_v
<
tensor_layout
::
gemm
::
ColumnMajor
,
B0Layout
>
&&
CheckDLayout
<
tensor_layout
::
gemm
::
RowMajor
,
D0sLayout
,
NumD0Tensor
>
()
&&
CheckDLayout
<
tensor_layout
::
gemm
::
RowMajor
,
D0sLayout
,
NumD0Tensor
>
()
&&
(
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
B1Layout
>
||
(
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
B1Layout
>
||
is_same_v
<
tensor_layout
::
gemm
::
ColumnMajor
,
B1Layout
>
)
&&
is_same_v
<
tensor_layout
::
gemm
::
ColumnMajor
,
CheckDLayout
<
tensor_layout
::
gemm
::
RowMajor
,
D1sLayout
,
NumD1Tensor
>
()
&&
B1Layout
>
)
&&
CheckDLayout
<
tensor_layout
::
gemm
::
RowMajor
,
D1sLayout
,
NumD1Tensor
>
()
&&
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
E1Layout
>
))
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
E1Layout
>
))
{
{
return
false
;
return
false
;
...
...
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