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_ROCM
Commits
be48abdb
Commit
be48abdb
authored
May 21, 2024
by
Adam Osewski
Browse files
Add function returing # of next K tiles.
parent
821f5dff
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
29 additions
and
0 deletions
+29
-0
include/ck/utility/work_scheduling.hpp
include/ck/utility/work_scheduling.hpp
+29
-0
No files found.
include/ck/utility/work_scheduling.hpp
View file @
be48abdb
...
...
@@ -53,6 +53,35 @@ class StridedReductionTileLoop
return
HasTile
();
}
/// @brief Returns the number of next k-tiles to process.
/// @param[in] k_tiles The number of tiles in the reduced dimension.
/// @param[in] k_tile_idx Current k-tile index.
/// @return The number of next k-tiles to process.
__device__
index_t
GetNextKTiles
(
index_t
k_tiles
,
index_t
k_tile_idx
)
{
index_t
k_tiles_left
=
k_tiles
-
k_tile_idx
;
index_t
next_k_tiles
=
k_tiles_left
<=
tiles_per_block_
?
k_tiles_left
:
tiles_per_block_
-
block_tile_idx_
;
tile_id_
+=
next_k_tiles
;
block_tile_idx_
+=
next_k_tiles
;
if
(
blockIdx
.
x
<
4
&&
ck
::
debug
::
is_thread_local_1d_id_idx
<
0
>
())
{
printf
(
"[GetNextKTiles] bid: %d, k_tiles: %d, k_idx:%d, next_k_tiles: %d, "
"k_tiles_left: %d,"
" tile_id: %d, block_tile_idx: %d
\n
"
,
static_cast
<
index_t
>
(
blockIdx
.
x
),
k_tiles
,
k_tile_idx
,
next_k_tiles
,
k_tiles_left
,
tile_id_
,
block_tile_idx_
);
}
return
next_k_tiles
;
}
__device__
index_t
GetFlagCount
()
const
{
return
get_grid_size
();
}
///
...
...
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