ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-16 08:44:55 +00:00
Code Issues Packages Projects Releases Wiki Activity

added stream-k policies for 3 tile and 4 tile

Browse Source
This commit is contained in:
Harisankar Sadasivan
2024-06-25 22:37:55 +00:00
parent 03c252551c
commit 6cc03335e3
2 changed files with 11 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
View File

@@ -1471,6 +1471,16 @@ struct BlockToCTileMap_GemmStreamK_v2
{
sk_tiles = bigEnough ? (grid_size + num_tiles % grid_size) : num_tiles;
}
else if(streamk_sel == 3)
{
sk_tiles = (num_tiles > (2 * grid_size)) ? (2 * grid_size + num_tiles % grid_size)
: num_tiles;
}
else if(streamk_sel == 4)
{
sk_tiles = (num_tiles > (3 * grid_size)) ? (3 * grid_size + num_tiles % grid_size)
: num_tiles;
}
sk_num_blocks = sk_tiles;
// if(sk_tiles < sk_num_blocks)
// {

2
profiler/include/profiler/profile_gemm_universal_streamk_impl.hpp
View File

@@ -153,7 +153,7 @@ bool profile_gemm_universal_streamk_impl(int do_verification,
for(auto& op_ptr : op_ptrs)
{
std::vector<int> grid_size_list = {38, 76, 114, 152, 190, 228, 266, 304, 342, 380};
std::vector<int> streamk_sel_list = {0, 1, 2};
std::vector<int> streamk_sel_list = {0, 1, 2, 3, 4};
if(Grid_size == -1)
{