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
6230f4ec
Commit
6230f4ec
authored
Nov 08, 2023
by
Harisankar Sadasivan
Browse files
block to ctile map for tall and skinny gemm
parent
49facb91
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
8 additions
and
0 deletions
+8
-0
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
+8
-0
No files found.
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
View file @
6230f4ec
...
...
@@ -658,6 +658,14 @@ struct BlockToCTileMap_3DGrid_KSplit
return
make_tuple
(
blockIdx
.
z
,
blockIdx
.
y
,
blockIdx
.
x
);
}
__host__
__device__
inline
constexpr
auto
convert_1D_block_idx_to_3D_tuple
(
const
index_t
&
block_1d_id
,
const
index_t
&
N
,
const
index_t
&
k_batch
)
const
{
const
auto
Ndim
=
math
::
integer_divide_ceil
(
N
,
NPerBlock
);
return
make_tuple
(((
block_1d_id
)
/
(
k_batch
*
Ndim
)),
(((
block_1d_id
)
/
k_batch
)
%
Ndim
),
(
block_1d_id
)
%
k_batch
);
// returns 3D tuple as (Mid,Nid,Kid)
}
template
<
typename
CTileIdx
,
typename
CTileDim
>
__host__
__device__
bool
ValidCTileIndex
(
const
CTileIdx
&
/* c_tile_idx */
,
const
CTileDim
&
/* c_tile_dim */
)
const
...
...
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