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
MIGraphX
Commits
f5d3f160
Commit
f5d3f160
authored
Oct 20, 2022
by
Paul
Browse files
Fix assert
parent
f903e858
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
8 additions
and
7 deletions
+8
-7
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
+2
-1
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm_includes.hpp
...gpu/kernels/include/migraphx/kernels/ck_gemm_includes.hpp
+6
-6
No files found.
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
View file @
f5d3f160
...
...
@@ -30,6 +30,7 @@
#include <migraphx/kernels/tensor_view.hpp>
#include <migraphx/kernels/ck.hpp>
#include <migraphx/kernels/ck_gemm_includes.hpp>
#include <migraphx/kernels/print.hpp>
namespace
migraphx
{
...
...
@@ -59,7 +60,7 @@ __device__ void ck_gemm(E e, A a, B b, Ds... ds)
constexpr
auto
e_grid_desc_mblock_mperblock_nblock_nperblock
=
GridwiseGemm
::
MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
(
e_grid_desc_m_n
);
MIGRAPHX_CK_STATIC_ASSERT
(
G
ridwiseGemm
::
CheckValidity
(
MIGRAPHX_CK_STATIC_ASSERT
(
G
::
CheckValidity
(
a_grid_desc_m_k
,
b_grid_desc_n_k
,
ds_grid_desc_m_n
,
e_grid_desc_m_n
,
block_2_etile_map
));
__shared__
char
p_shared_block
[
GridwiseGemm
::
GetSharedMemoryNumberOfByte
()];
...
...
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm_includes.hpp
View file @
f5d3f160
...
...
@@ -233,18 +233,18 @@ struct CK_DeviceGemmMultipleD
const
EGridDesc_M_N
&
e_grid_desc_m_n
,
const
Block2ETileMap
&
block_2_etile_map
)
{
const
expr
auto
M
=
a_grid_desc_m_k
.
GetLength
(
I0
);
const
expr
auto
N
=
b_grid_desc_n_k
.
GetLength
(
I0
);
const
expr
auto
K
=
a_grid_desc_m_k
.
GetLength
(
I1
);
const
auto
M
=
a_grid_desc_m_k
.
GetLength
(
I0
);
const
auto
N
=
b_grid_desc_n_k
.
GetLength
(
I0
);
const
auto
K
=
a_grid_desc_m_k
.
GetLength
(
I1
);
// check consistency of desc
static_assert
(
M
==
e_grid_desc_m_n
.
GetLength
(
I0
)
&&
N
==
e_grid_desc_m_n
.
GetLength
(
I1
));
MIGRAPHX_CHECK
(
M
==
e_grid_desc_m_n
.
GetLength
(
I0
)
&&
N
==
e_grid_desc_m_n
.
GetLength
(
I1
));
// check tile size
static_assert
(
M
%
MPerBlock
==
0
&&
N
%
NPerBlock
==
0
&&
K
%
KPerBlock
==
0
);
MIGRAPHX_CHECK
(
M
%
MPerBlock
==
0
&&
N
%
NPerBlock
==
0
&&
K
%
KPerBlock
==
0
);
// check block-to-E-tile
static_assert
(
block_2_etile_map
.
CheckValidity
(
e_grid_desc_m_n
));
MIGRAPHX_CHECK
(
block_2_etile_map
.
CheckValidity
(
e_grid_desc_m_n
));
return
GridwiseGemm
::
CheckValidity
(
a_grid_desc_m_k
,
b_grid_desc_n_k
,
ds_grid_desc_m_n
,
e_grid_desc_m_n
,
block_2_etile_map
);
...
...
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