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
025cfd86
"docs/source/en/api/pipelines/ddim.md" did not exist on "1a79969d23a18481619db73eec9f4445d6eaed86"
Commit
025cfd86
authored
Mar 08, 2023
by
Rosty Geyyer
Browse files
Add layout check to IsSupportedArgument
parent
1e59eb3b
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
32 additions
and
0 deletions
+32
-0
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
+32
-0
No files found.
include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle.hpp
View file @
025cfd86
...
@@ -788,6 +788,20 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
...
@@ -788,6 +788,20 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
return
true
;
return
true
;
}
}
// check if DsLayout is supported
template
<
typename
RefLayout
,
typename
DsLayout
,
const
index_t
NumDTensor
>
static
bool
CheckDLayout
()
{
static
bool
valid
=
true
;
// iterate over DLayout tuple
static_for
<
0
,
NumDTensor
,
1
>
{}([
&
](
auto
i
)
{
using
DLayout
=
remove_cvref_t
<
tuple_element_t
<
i
.
value
,
DsLayout
>>
;
// if RefLayout and DLayout are same, keep valid true, otherwise false
valid
=
valid
&&
is_same_v
<
RefLayout
,
DLayout
>
;
});
return
valid
;
}
static
bool
IsSupportedArgument
(
const
Argument
&
arg
)
static
bool
IsSupportedArgument
(
const
Argument
&
arg
)
{
{
if
(
!
(
ck
::
get_device_name
()
==
"gfx908"
||
ck
::
get_device_name
()
==
"gfx90a"
))
if
(
!
(
ck
::
get_device_name
()
==
"gfx908"
||
ck
::
get_device_name
()
==
"gfx90a"
))
...
@@ -795,6 +809,24 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
...
@@ -795,6 +809,24 @@ struct DeviceBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
return
false
;
return
false
;
}
}
// Check supported layouts
// A0 - Row
// B0 - Col
// D0s - Rows
// B1 - Row or Col
// D1s - Rows
// E1 - Row
if
(
!
(
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
A0Layout
>
&&
is_same_v
<
tensor_layout
::
gemm
::
ColumnMajor
,
B0Layout
>
&&
CheckDLayout
<
tensor_layout
::
gemm
::
RowMajor
,
D0sLayout
,
NumD0Tensor
>
()
&&
(
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
B1Layout
>
||
is_same_v
<
tensor_layout
::
gemm
::
ColumnMajor
,
B1Layout
>
)
&&
CheckDLayout
<
tensor_layout
::
gemm
::
RowMajor
,
D1sLayout
,
NumD1Tensor
>
()
&&
is_same_v
<
tensor_layout
::
gemm
::
RowMajor
,
E1Layout
>
))
{
return
false
;
}
return
GridwiseGemm
::
CheckValidity
(
arg
.
a0_grid_desc_m_k_
,
return
GridwiseGemm
::
CheckValidity
(
arg
.
a0_grid_desc_m_k_
,
arg
.
b0_grid_desc_n_k_
,
arg
.
b0_grid_desc_n_k_
,
arg
.
b1_grid_desc_n_k_
,
arg
.
b1_grid_desc_n_k_
,
...
...
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