ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-04-19 22:39:03 +00:00
Code Issues Packages Projects Releases Wiki Activity

add the check of granularity for atomic add (#2959)

Browse Source
This commit is contained in:
Thomas Ning
2025-10-02 11:15:24 -07:00
committed by GitHub
parent 6fc28ab493
commit cadafde722
2 changed files with 8 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3.hpp
View File

@@ -682,6 +682,10 @@ struct DeviceGemmMultiD_Xdl_CShuffle_V3 : public DeviceGemmMultipleDSplitK<ALayo
return GridwiseGemm64::CheckValidity(arg);
}
}
if(CDEShuffleBlockTransferScalarPerVectors{}[Number<0>{}] <= 1 && (arg.KBatch > 1))
{
return false;
}
else
{
if constexpr(NXdlPerWave32 > 0)

8
test/ck_tile/gemm/test_gemm_pipeline_ut_cases.inc
View File

@@ -11,7 +11,7 @@ TYPED_TEST(TEST_SUITE_NAME, SmallM)
std::vector<int> Ms{1, 2, 3, 4, 5, 6};
constexpr int N = 1024;
std::vector<int> Ks;
for (auto K_count: {2, 3, 4, 10, 11})
for(auto K_count : {2, 3, 4, 10, 11})
{
Ks.push_back(K_count * TestFixture::K_Tile);
}
@@ -36,10 +36,10 @@ TYPED_TEST(TEST_SUITE_NAME, SmallM)
TYPED_TEST(TEST_SUITE_NAME, MidLargeM)
{
std::vector<int> Ms{127, 255, 312, 799, 1573};
constexpr int N = 1024;
constexpr int N = 1024;
std::vector<int> Ks;
for (auto K_count: {2, 3, 4, 10, 11})
for(auto K_count : {2, 3, 4, 10, 11})
{
Ks.push_back(K_count * TestFixture::K_Tile);
}
@@ -51,7 +51,7 @@ TYPED_TEST(TEST_SUITE_NAME, MidLargeM)
for(int M : Ms)
{
for (int K: Ks)
for(int K : Ks)
{
if constexpr(std::is_same_v<typename TestFixture::ALayout,
ck_tile::tensor_layout::gemm::ColumnMajor>)