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
5581dc00
Commit
5581dc00
authored
May 04, 2023
by
Po-Yen, Chen
Browse files
Reuse the existing implementation
parent
b250bbb8
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
9 additions
and
7 deletions
+9
-7
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
+8
-5
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
...nsor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
+1
-2
No files found.
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
View file @
5581dc00
...
...
@@ -125,14 +125,17 @@ struct BlockToCTileMap_M00_N0_M01Adapt
{
}
__host__
constexpr
index_t
CalculateGridSize
(
const
CGridDesc_M_N
&
c_grid_desc_m_n
)
const
__host__
static
constexpr
index_t
CalculateGridSize
(
index_t
M
,
index_t
N
)
{
const
auto
M0
=
math
::
integer_divide_ceil
(
c_grid_desc_m_n
.
GetLength
(
I0
)
,
MPerBlock
);
const
auto
N0
=
math
::
integer_divide_ceil
(
c_grid_desc_m_n
.
GetLength
(
I1
)
,
NPerBlock
);
const
auto
M0
=
math
::
integer_divide_ceil
(
M
,
MPerBlock
);
const
auto
N0
=
math
::
integer_divide_ceil
(
N
,
NPerBlock
);
const
index_t
grid_size
=
M0
*
N0
;
return
M0
*
N0
;
}
return
grid_size
;
__host__
constexpr
index_t
CalculateGridSize
(
const
CGridDesc_M_N
&
c_grid_desc_m_n
)
const
{
return
CalculateGridSize
(
c_grid_desc_m_n
.
GetLength
(
I0
),
c_grid_desc_m_n
.
GetLength
(
I1
));
}
template
<
typename
TopIdx
>
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
View file @
5581dc00
...
...
@@ -144,8 +144,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
__host__
__device__
static
auto
CalculateGridSize
(
index_t
M
,
index_t
N
)
{
// reference the implementation of class 'BlockToCTileMap_M00_N0_M01Adapt'
return
std
::
make_tuple
(
INTEGER_DIVIDE_CEIL
(
M
,
MPerBlock
)
*
INTEGER_DIVIDE_CEIL
(
N
,
NPerBlock
),
1
,
1
);
return
std
::
make_tuple
(
DefaultBlock2CTileMap
::
CalculateGridSize
(
M
,
N
),
1
,
1
);
}
__host__
__device__
static
auto
CalculateMPadded
(
index_t
M
)
...
...
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