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
2a16c61c
Commit
2a16c61c
authored
Jun 05, 2024
by
Adam Osewski
Browse files
Fix GetNextKTiles.
parent
3f74017c
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
4 additions
and
18 deletions
+4
-18
include/ck/utility/work_scheduling.hpp
include/ck/utility/work_scheduling.hpp
+4
-18
No files found.
include/ck/utility/work_scheduling.hpp
View file @
2a16c61c
...
@@ -59,26 +59,12 @@ class StridedReductionTileLoop
...
@@ -59,26 +59,12 @@ class StridedReductionTileLoop
/// @return The number of next k-tiles to process.
/// @return The number of next k-tiles to process.
__device__
index_t
GetNextKTiles
(
index_t
k_tiles
,
index_t
k_tile_idx
)
__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
k_tiles_left
=
k_tiles
-
k_tile_idx
;
index_t
next_k_tiles
=
index_t
block_tiles_left
=
tiles_per_block_
-
block_tile_idx_
;
k_tiles_left
<=
tiles_per_block_
?
k_tiles_left
:
tiles_per_block_
-
block_tile_idx_
;
index_t
next_k_tiles
=
k_tiles_left
<=
block_tiles_left
?
k_tiles_left
:
block_tiles_left
;
tile_id_
+=
next_k_tiles
;
tile_id_
+=
next_k_tiles
;
block_tile_idx_
+=
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
;
return
next_k_tiles
;
}
}
...
...
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