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
d8f97e5b
Commit
d8f97e5b
authored
Jul 19, 2023
by
turneram
Browse files
Fix
parent
9472bce3
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
8 additions
and
6 deletions
+8
-6
include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp
...ration/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp
+8
-6
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp
View file @
d8f97e5b
...
@@ -273,7 +273,7 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
...
@@ -273,7 +273,7 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
// check consistency of desc
// check consistency of desc
if
(
!
(
M
==
e_grid_desc_m_n
.
GetLength
(
I0
)
&&
N
==
e_grid_desc_m_n
.
GetLength
(
I1
)))
if
(
!
(
M
==
e_grid_desc_m_n
.
GetLength
(
I0
)
&&
N
==
e_grid_desc_m_n
.
GetLength
(
I1
)))
{
{
static_assert
(
false
,
"e_grid_desc invalid
\n
"
);
static_assert
(
(
M
==
e_grid_desc_m_n
.
GetLength
(
I0
)
&&
N
==
e_grid_desc_m_n
.
GetLength
(
I1
))
,
"e_grid_desc invalid
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -286,14 +286,14 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
...
@@ -286,14 +286,14 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
if
(
!
valid
)
if
(
!
valid
)
{
{
static_assert
(
f
al
se
,
"ds_grid_desc invalid
\n
"
);
static_assert
(
v
al
id
,
"ds_grid_desc invalid
\n
"
);
return
false
;
return
false
;
}
}
// check tile size
// check tile size
if
(
!
(
M
%
MPerBlock
==
0
&&
N
%
NPerBlock
==
0
&&
K
%
KPerBlock
==
0
))
if
(
!
(
M
%
MPerBlock
==
0
&&
N
%
NPerBlock
==
0
&&
K
%
KPerBlock
==
0
))
{
{
static_assert
(
false
,
"tile size invalid
\n
"
);
static_assert
(
(
M
%
MPerBlock
==
0
&&
N
%
NPerBlock
==
0
&&
K
%
KPerBlock
==
0
)
,
"tile size invalid
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -302,14 +302,14 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
...
@@ -302,14 +302,14 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
if
(
!
GridwiseGemmPipe
::
IsSupported
(
num_k_loop
))
if
(
!
GridwiseGemmPipe
::
IsSupported
(
num_k_loop
))
{
{
static_assert
(
false
,
"num_k_loop invalid
\n
"
);
static_assert
(
GridwiseGemmPipe
::
IsSupported
(
num_k_loop
)
,
"num_k_loop invalid
\n
"
);
return
false
;
return
false
;
}
}
// check block-to-E-tile
// check block-to-E-tile
if
(
!
block_2_etile_map
.
CheckValidity
(
e_grid_desc_m_n
))
if
(
!
block_2_etile_map
.
CheckValidity
(
e_grid_desc_m_n
))
{
{
static_assert
(
false
,
"block_2_etile_map invalid
\n
"
);
static_assert
(
block_2_etile_map
.
CheckValidity
(
e_grid_desc_m_n
)
,
"block_2_etile_map invalid
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -321,7 +321,9 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
...
@@ -321,7 +321,9 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
b_grid_desc_n_k
.
GetElementSpaceSize
()
*
sizeof
(
ABDataType
)
<=
TwoGB
&&
b_grid_desc_n_k
.
GetElementSpaceSize
()
*
sizeof
(
ABDataType
)
<=
TwoGB
&&
e_grid_desc_m_n
.
GetElementSpaceSize
()
*
sizeof
(
EDataType
)
<=
TwoGB
))
e_grid_desc_m_n
.
GetElementSpaceSize
()
*
sizeof
(
EDataType
)
<=
TwoGB
))
{
{
static_assert
(
false
,
"invalid tensor (> 2GB)
\n
"
);
static_assert
((
a_grid_desc_m_k
.
GetElementSpaceSize
()
*
sizeof
(
ABDataType
)
<=
TwoGB
&&
b_grid_desc_n_k
.
GetElementSpaceSize
()
*
sizeof
(
ABDataType
)
<=
TwoGB
&&
e_grid_desc_m_n
.
GetElementSpaceSize
()
*
sizeof
(
EDataType
)
<=
TwoGB
),
"invalid tensor (> 2GB)
\n
"
);
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