CK-Tile Grouped GEMM refactor and post PR fixes #1756
Conversation
mozga-amd
requested review from
junliume,
illsilin,
carlushuang,
qianfengz,
aosewski,
poyenc,
geyyer,
bartekxk and
andriy-ca
as code owners
aosewski
changed the title
aosewski
dismissed
their stale review
Will be off from tomorrow till 2nd Jan. Please pass on review to ie. @bartekxk
|
|
||
| static constexpr index_t MPerBlock = BlockGemmShape::kM; | ||
| static constexpr index_t NPerBlock = BlockGemmShape::kN; | ||
| static constexpr index_t KPerBlock = BlockGemmShape::kK; | ||
|
|
||
| CK_TILE_HOST static constexpr auto GridSize(index_t M, index_t N) | ||
| CK_TILE_HOST static constexpr auto | ||
| GridSize(index_t M, index_t N) noexcept(noexcept(MPerBlock != 0 && NPerBlock != 0)) -> dim3 |
There was a problem hiding this comment.
I wonder why not making static_assert right after those constexpr variables ?
There was a problem hiding this comment.
Mainly, we additionally provide info to compiler that we do not expect the function to throw an exception in any cases, except in this specific case...
The result will be identical (in fn-level). For class-lvl... the difference lies only in when the exception will be caught and what is caught—either at the class scope level or the function scope level.
In our case, the function will catch the exception at the function call level, not at the class call level. So, it will happen slightly later - but still in compile-time. Additionally, it's obvious that the function will never throw ~ and it's good to mark it as noexcept. Additionally. it could reduce the binary size by removing the handling code
Just let me know, If you want I can use either static_assert in class-scope or fn-scope
| const index_t NBlock = GetNBlock(N_); | ||
|
|
||
| const index_t iM = __builtin_amdgcn_readfirstlane(blockIdx / NBlock); | ||
| const index_t iN = __builtin_amdgcn_readfirstlane(modulo(blockIdx, NBlock)); |
This pull-request contains changes as following:
get_stridefunction, etc., to a different location,