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
060d171f
Commit
060d171f
authored
Feb 13, 2025
by
ozturkosu
Browse files
changed StreamKReductionStrategy from Atomic to Reduction as hardcoded
parent
3c7fef7f
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
25 additions
and
3 deletions
+25
-3
include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp
...n/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp
+0
-0
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
+7
-1
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp
...ration/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp
+18
-2
No files found.
include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp
100644 → 100755
View file @
060d171f
File mode changed from 100644 to 100755
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp
100644 → 100755
View file @
060d171f
...
...
@@ -1010,10 +1010,16 @@ enum StreamKReductionStrategy
Reduction
,
// let some workgroup responsible for doing the reduction operation
};
// template <uint32_t MPerBlock_,
// uint32_t NPerBlock_,
// uint32_t KPerBlock_,
// StreamKReductionStrategy ReductionStrategy_ = StreamKReductionStrategy::Atomic,
// uint32_t TileSwizzleSubM_ = 8>
template
<
uint32_t
MPerBlock_
,
uint32_t
NPerBlock_
,
uint32_t
KPerBlock_
,
StreamKReductionStrategy
ReductionStrategy_
=
StreamKReductionStrategy
::
Atomic
,
StreamKReductionStrategy
ReductionStrategy_
=
StreamKReductionStrategy
::
Reduction
,
uint32_t
TileSwizzleSubM_
=
8
>
struct
BlockToCTileMap_GemmStreamK
{
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp
View file @
060d171f
...
...
@@ -539,10 +539,18 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3
const
ADataType
*
p_a_grid
;
const
BDataType
*
p_b_grid
;
CDataType
*
p_c_grid
;
// BlockToCTileMap_GemmStreamK_v2<MPerBlock,
// NPerBlock,
// KPerBlock,
// StreamKReductionStrategy::Atomic,
// 8,
// 4>
// block_2_ctile_map_streamk;
BlockToCTileMap_GemmStreamK_v2
<
MPerBlock
,
NPerBlock
,
KPerBlock
,
StreamKReductionStrategy
::
Atomic
,
StreamKReductionStrategy
::
Reduction
,
8
,
4
>
block_2_ctile_map_streamk
;
...
...
@@ -1176,10 +1184,18 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3
}();
return
c_partial_acc_block_m_n
;
}
// using Block2CTileMap_streamk = BlockToCTileMap_GemmStreamK_v2<MPerBlock,
// NPerBlock,
// KPerBlock,
// StreamKReductionStrategy::Atomic,
// 8,
// 4>;
using
Block2CTileMap_streamk
=
BlockToCTileMap_GemmStreamK_v2
<
MPerBlock
,
NPerBlock
,
KPerBlock
,
StreamKReductionStrategy
::
Atomic
,
StreamKReductionStrategy
::
Reduction
,
8
,
4
>
;
...
...
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