From adb5822c26c6a038a406a11692aaf64f53e62ce3 Mon Sep 17 00:00:00 2001 From: Joseph Macaranas Date: Mon, 14 Sep 2026 00:51:30 -0400 Subject: [PATCH] fix(tensilelite): revert Stream-K launch limit changes Revert the Stream-K launch limit changes from #11518. TheRock Multi-Arch CI showed that the rocm-libraries #8183 bump failed the standard TensileLite test job on both gfx94X and gfx950. Restoring the three files changed by #11518 on top of that failing pin made both targeted jobs pass. Fixes #12021 --- .../tensilelite/src/ContractionSolution.cpp | 19 +- .../tensilelite/tests/CuCount_test.cpp | 346 ++++++------------ .../tests/StreamKLaunchSummary_test.cpp | 7 +- 3 files changed, 120 insertions(+), 252 deletions(-) diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index dd60daa9c4c2..553efbfed0c4 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -5454,9 +5454,7 @@ namespace TensileLite { // Limit workgroups per CU to 3 // TODO Verify this limit is best - const size_t occupancy - = std::max(self.sizeMapping.CUOccupancy, static_cast(1)); - auto kernelOccupancy = std::min(occupancy, size_t{3}); + auto kernelOccupancy = std::min(self.sizeMapping.CUOccupancy, 3); auto maxGrid = cuCount * kernelOccupancy; if(pAMDGPU->skMaxCUs > 0) { @@ -5680,22 +5678,11 @@ namespace TensileLite { size_t itersPerTile = problem.getItersPerTile(self.sizeMapping); size_t itersPerWG = tiles * itersPerTile / skGrid; - // numWorkItems = skGrid * threadsPerWorkGroup - // the largest grid this fallback can express is - // UINT32_MAX / threadsPerWorkGroup -- 2^24 tiles at 256 threads - // but 2^25 at 128 and 2^22 at 1024. - const size_t threadsPerWorkGroup = self.sizeMapping.workGroupSize.x - * self.sizeMapping.workGroupSize.y - * self.sizeMapping.workGroupSize.z; - const size_t maxTiles = (size_t{std::numeric_limits::max()} + 1) - / std::max(threadsPerWorkGroup, size_t{1}); if(itersPerTile >= 65536 || itersPerWG >= 65536 - || (tiles * itersPerTile) >= maxTiles) + || (tiles * itersPerTile) >= 16777216) { - const size_t occupancy - = std::max(self.sizeMapping.CUOccupancy, static_cast(1)); - skGrid = (tiles >= maxTiles) ? cuCount * occupancy : tiles; + skGrid = tiles; if(outTreeBoundsFallback) *outTreeBoundsFallback = true; } diff --git a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp index 1630871595ab..b25156b1bd9a 100644 --- a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp @@ -33,30 +33,30 @@ using namespace TensileLite::testing; TEST(CuCountPredicateTest, MatchesSPX) { - auto pred = std::make_shared(_SPX_CU); - AMDGPU spx = makeDevice(_MI350_CHIP_ID, _SPX_CU, "spx"); - AMDGPU cpx = makeDevice(_MI350_CHIP_ID, _CPX_CU, "cpx"); + auto pred = std::make_shared(_SPX_CU); + AMDGPU spx = makeDevice(_MI350_CHIP_ID, _SPX_CU, "spx"); + AMDGPU cpx = makeDevice(_MI350_CHIP_ID, _CPX_CU, "cpx"); - EXPECT_TRUE((*pred)(spx)) << "CUCountEqual(256) should match SPX (CU=256)"; + EXPECT_TRUE((*pred)(spx)) << "CUCountEqual(256) should match SPX (CU=256)"; EXPECT_FALSE((*pred)(cpx)) << "CUCountEqual(256) should NOT match CPX (CU=64)"; } TEST(CuCountPredicateTest, MatchesCPX) { - auto pred = std::make_shared(_CPX_CU); - AMDGPU spx = makeDevice(_MI350_CHIP_ID, _SPX_CU, "spx"); - AMDGPU cpx = makeDevice(_MI350_CHIP_ID, _CPX_CU, "cpx"); + auto pred = std::make_shared(_CPX_CU); + AMDGPU spx = makeDevice(_MI350_CHIP_ID, _SPX_CU, "spx"); + AMDGPU cpx = makeDevice(_MI350_CHIP_ID, _CPX_CU, "cpx"); - EXPECT_TRUE((*pred)(cpx)) << "CUCountEqual(64) should match CPX (CU=64)"; + EXPECT_TRUE((*pred)(cpx)) << "CUCountEqual(64) should match CPX (CU=64)"; EXPECT_FALSE((*pred)(spx)) << "CUCountEqual(64) should NOT match SPX (CU=256)"; } TEST(CuCountPredicateTest, NoCuCheckMatchesBoth) { // A hardware predicate with no CUCountEqual accepts any CU configuration. - auto hwPred = makeHwPred(AMDGPU::Processor::gfx950, _MI350_CHIP_ID); - AMDGPU spx = makeDevice(_MI350_CHIP_ID, _SPX_CU, "spx"); - AMDGPU cpx = makeDevice(_MI350_CHIP_ID, _CPX_CU, "cpx"); + auto hwPred = makeHwPred(AMDGPU::Processor::gfx950, _MI350_CHIP_ID); + AMDGPU spx = makeDevice(_MI350_CHIP_ID, _SPX_CU, "spx"); + AMDGPU cpx = makeDevice(_MI350_CHIP_ID, _CPX_CU, "cpx"); EXPECT_TRUE((*hwPred.value)(spx)) << "Predicate without CU check should match SPX"; EXPECT_TRUE((*hwPred.value)(cpx)) << "Predicate without CU check should match CPX"; @@ -89,9 +89,10 @@ class CuCountFallbackTest : public ::testing::Test { std::string got = selectSolution(lib, device, device.deviceName); EXPECT_EQ(got, expectedName) - << "Device " << device.deviceName << " (chip=" << hexChipId(device.pciChipId().value()) - << ", CU=" << device.computeUnitCount << "): expected \"" << expectedName - << "\", got \"" << got << "\""; + << "Device " << device.deviceName + << " (chip=" << hexChipId(device.pciChipId().value()) + << ", CU=" << device.computeUnitCount + << "): expected \"" << expectedName << "\", got \"" << got << "\""; } }; @@ -114,10 +115,12 @@ TEST_F(CuCountFallbackTest, CpxFallsBackToSpxOob) buildProblemLib(singleLib(spx_eq), singleLib(spx_oob))}, // Row 2: mi350, any CU -- oob only (catch-all for CPX) - {makeHwPred(gfx950, _MI350_CHIP_ID), buildProblemLib(singleLib(spx_oob))}, + {makeHwPred(gfx950, _MI350_CHIP_ID), + buildProblemLib(singleLib(spx_oob))}, // Row 3: gfx950 catch-all - {makeHwPred(gfx950), buildProblemLib(singleLib(spx_oob))}, + {makeHwPred(gfx950), + buildProblemLib(singleLib(spx_oob))}, }); expectSelected(*lib, mi350spx, "mi350spx_eq"); @@ -145,7 +148,8 @@ TEST_F(CuCountFallbackTest, CpxWithOwnEq) buildProblemLib(singleLib(cpx_eq), singleLib(spx_oob))}, // Row 3: mi350, any CU -- oob - {makeHwPred(gfx950, _MI350_CHIP_ID), buildProblemLib(singleLib(spx_oob))}, + {makeHwPred(gfx950, _MI350_CHIP_ID), + buildProblemLib(singleLib(spx_oob))}, }); expectSelected(*lib, mi350spx, "mi350spx_eq"); @@ -177,13 +181,16 @@ TEST_F(CuCountFallbackTest, CpxFallsToSameChipOob) buildProblemLib(singleLib(mi350spx_eq), singleLib(mi350spx_oob))}, // Row 3: mi355, any CU -- mi355 oob - {makeHwPred(gfx950, _MI355_CHIP_ID), buildProblemLib(singleLib(mi355spx_oob))}, + {makeHwPred(gfx950, _MI355_CHIP_ID), + buildProblemLib(singleLib(mi355spx_oob))}, // Row 4: mi350, any CU -- mi350 oob - {makeHwPred(gfx950, _MI350_CHIP_ID), buildProblemLib(singleLib(mi350spx_oob))}, + {makeHwPred(gfx950, _MI350_CHIP_ID), + buildProblemLib(singleLib(mi350spx_oob))}, // Row 5: catch-all - {makeHwPred(gfx950), buildProblemLib(singleLib(mi350spx_oob))}, + {makeHwPred(gfx950), + buildProblemLib(singleLib(mi350spx_oob))}, }); // CPX devices skip the CU=256 rows, then hit their chip-specific no-CU row. @@ -217,7 +224,8 @@ TEST_F(CuCountFallbackTest, CpxAndSpxIndependent) buildProblemLib(singleLib(mi350cpx_eq), singleLib(mi350spx_oob))}, // Row 3: mi350, any CU -- oob - {makeHwPred(gfx950, _MI350_CHIP_ID), buildProblemLib(singleLib(mi350spx_oob))}, + {makeHwPred(gfx950, _MI350_CHIP_ID), + buildProblemLib(singleLib(mi350spx_oob))}, }); expectSelected(*lib, mi350spx, "mi350spx_eq"); @@ -225,7 +233,7 @@ TEST_F(CuCountFallbackTest, CpxAndSpxIndependent) // Verify each mode did NOT cross-select. // An mi350spx device should not get the CPX solution and vice versa. - auto problem = dummyProblem(); + auto problem = dummyProblem(); auto spxResult = lib->findBestSolution(problem, mi350spx); auto cpxResult = lib->findBestSolution(problem, mi350cpx); @@ -238,17 +246,17 @@ TEST_F(CuCountFallbackTest, CpxAndSpxIndependent) TEST(StreamKForceDPOnlyTest, UsesHardwareCuCount) { ContractionSolution solution; - solution.sizeMapping.streamK = 3; - solution.sizeMapping.streamKForceDPOnly = 1; - solution.sizeMapping.macroTile = TensileLite::dim3(128, 128, 1); - solution.sizeMapping.depthU = 64; - solution.sizeMapping.matrixInstruction = {16, 16, 32, 1}; - solution.sizeMapping.CUOccupancy = 1; - - auto problem = dummyProblem(); - auto device = makeDevice(_MI350_CHIP_ID, _CPX_CU, "mi350cpx"); + solution.sizeMapping.streamK = 3; + solution.sizeMapping.streamKForceDPOnly = 1; + solution.sizeMapping.macroTile = TensileLite::dim3(128, 128, 1); + solution.sizeMapping.depthU = 64; + solution.sizeMapping.matrixInstruction = {16, 16, 32, 1}; + solution.sizeMapping.CUOccupancy = 1; + + auto problem = dummyProblem(); + auto device = makeDevice(_MI350_CHIP_ID, _CPX_CU, "mi350cpx"); device.skDynamicGrid = 0; - auto tiles = problem.getNumTiles(solution.sizeMapping, 1); + auto tiles = problem.getNumTiles(solution.sizeMapping, 1); EXPECT_EQ(solution.getSKReduction(problem, device), origami::reduction_t::tree); EXPECT_EQ(solution.getSKGrid(problem, device, tiles, origami::reduction_t::tree), _CPX_CU); @@ -257,18 +265,18 @@ TEST(StreamKForceDPOnlyTest, UsesHardwareCuCount) TEST(StreamKForceDPOnlyTest, FixedGridOverridesForceDPOnlyGrid) { ContractionSolution solution; - solution.sizeMapping.streamK = 3; - solution.sizeMapping.streamKForceDPOnly = 1; - solution.sizeMapping.macroTile = TensileLite::dim3(128, 128, 1); - solution.sizeMapping.depthU = 64; - solution.sizeMapping.matrixInstruction = {16, 16, 32, 1}; - solution.sizeMapping.CUOccupancy = 1; - - auto problem = dummyProblem(); - auto device = makeDevice(_MI350_CHIP_ID, _CPX_CU, "mi350cpx"); + solution.sizeMapping.streamK = 3; + solution.sizeMapping.streamKForceDPOnly = 1; + solution.sizeMapping.macroTile = TensileLite::dim3(128, 128, 1); + solution.sizeMapping.depthU = 64; + solution.sizeMapping.matrixInstruction = {16, 16, 32, 1}; + solution.sizeMapping.CUOccupancy = 1; + + auto problem = dummyProblem(); + auto device = makeDevice(_MI350_CHIP_ID, _CPX_CU, "mi350cpx"); device.skDynamicGrid = 0; - device.skFixedGrid = 17; - auto tiles = problem.getNumTiles(solution.sizeMapping, 1); + device.skFixedGrid = 17; + auto tiles = problem.getNumTiles(solution.sizeMapping, 1); EXPECT_EQ(solution.getSKGrid(problem, device, tiles, origami::reduction_t::tree), device.skFixedGrid); @@ -278,7 +286,7 @@ TEST(StreamKForceDPOnlyTest, DoesNotRequestPartialWorkspace) { ContractionSolution solution; solution.sizeMapping.streamK = 3; - solution.sizeMapping.streamKForceDPOnly = 1; + solution.sizeMapping.streamKForceDPOnly = 1; solution.sizeMapping.streamKAtomic = 0; solution.sizeMapping.macroTile = TensileLite::dim3(256, 256, 1); solution.sizeMapping.depthU = 64; @@ -286,10 +294,10 @@ TEST(StreamKForceDPOnlyTest, DoesNotRequestPartialWorkspace) solution.sizeMapping.CUOccupancy = 1; solution.sizeMapping.workspaceSizePerElemC = 4; - auto problem = dummyProblem(); - auto device = makeDevice(_MI350_CHIP_ID, _CPX_CU, "mi350cpx"); + auto problem = dummyProblem(); + auto device = makeDevice(_MI350_CHIP_ID, _CPX_CU, "mi350cpx"); device.skDynamicGrid = 0; - auto tiles = problem.getNumTiles(solution.sizeMapping, 1); + auto tiles = problem.getNumTiles(solution.sizeMapping, 1); ASSERT_NE(tiles % _CPX_CU, 0); EXPECT_EQ(solution.requiredWorkspaceSize(problem, device), 0); @@ -321,7 +329,7 @@ namespace return origami::hardware_t(arch_t::gfx950, kGfx950AnalyticalCuCount, 163840, - 262144, // rf_capacity: 65536 regs * 4 bytes + 262144, // rf_capacity: 65536 regs * 4 bytes 8, 1.0, 1.0, @@ -376,48 +384,28 @@ namespace ContractionProblemGemm makeGemmProblem(size_t m, size_t n, size_t k) { - auto problem = ContractionProblemGemm::GEMM(false, false, m, n, k, m, n, m, 1.0, false, 1); + auto problem = ContractionProblemGemm::GEMM( + false, false, m, n, k, m, n, m, 1.0, false, 1); problem.setComputeInputTypeA(rocisa::DataType::Float); problem.setComputeInputTypeB(rocisa::DataType::Float); return problem; } - void initBenchStreamK5Solution(ContractionSolution& solution, - TensileLite::dim3 macroTile, - size_t depthU) - { - solution.sizeMapping.streamK = 5; - solution.sizeMapping.macroTile = macroTile; - solution.sizeMapping.depthU = depthU; - solution.sizeMapping.workGroupSize = TensileLite::dim3(16, 16, 1); - solution.sizeMapping.matrixInstruction = {16, 16, 32, 1}; - solution.sizeMapping.CUOccupancy = 1; - } - - size_t threadsPerWorkGroup(ContractionSolution const& solution) - { - return solution.sizeMapping.workGroupSize.x * solution.sizeMapping.workGroupSize.y - * solution.sizeMapping.workGroupSize.z; - } - - uint64_t totalWorkItems(size_t skGrid, size_t threadsPerWG) - { - return static_cast(skGrid) * threadsPerWG; - } - StreamKHostPack computeStreamKHostPack(ContractionSolution const& solution, ContractionProblemGemm& problem, Hardware const& hardware) { StreamKHostPack pack{}; - pack.tiles = problem.getNumTiles(solution.sizeMapping, 1); - pack.itersPerTile = std::max(size_t{1}, problem.getItersPerTile(solution.sizeMapping)); + pack.tiles = problem.getNumTiles(solution.sizeMapping, 1); + pack.itersPerTile + = std::max(size_t{1}, problem.getItersPerTile(solution.sizeMapping)); if(solution.sizeMapping.streamK == 5) { pack.effectiveDynamic = solution.streamK5EffectiveDynamic(problem, hardware); - pack.reduction = pack.effectiveDynamic ? origami::reduction_t::tree - : solution.getSKReduction(problem, hardware); + pack.reduction = pack.effectiveDynamic + ? origami::reduction_t::tree + : solution.getSKReduction(problem, hardware); } else { @@ -429,9 +417,9 @@ namespace if(pack.reduction == origami::reduction_t::parallel) { - uint32_t skSplit = static_cast(pack.grid / pack.tiles); - pack.skItersPerWG = static_cast(pack.itersPerTile) / skSplit; - pack.skTiles = skSplit; + uint32_t skSplit = static_cast(pack.grid / pack.tiles); + pack.skItersPerWG = static_cast(pack.itersPerTile) / skSplit; + pack.skTiles = skSplit; } else { @@ -495,11 +483,11 @@ class StreamK5ExplicitModeTest : public ::testing::TestWithParam const& info) { - return info.param.mode == 0 ? "OffStatic" : "OnDynamic"; - }); +INSTANTIATE_TEST_SUITE_P( + StreamK5HybridModeTest, + StreamK5ExplicitModeTest, + ::testing::Values(StreamK5ExplicitModeParam{0, false}, StreamK5ExplicitModeParam{1, true}), + [](::testing::TestParamInfo const& info) { + return info.param.mode == 0 ? "OffStatic" : "OnDynamic"; + }); TEST(StreamK5HybridModeTest, TriStateAutoRequiresAnalyticalHardware) { ContractionSolution solution; initStreamK5Solution(solution); - auto device = makeHipDeviceWithoutAnalytical(); - auto problem = makeGemmProblem(4096, 4096, 64); + auto device = makeHipDeviceWithoutAnalytical(); + auto problem = makeGemmProblem(4096, 4096, 64); problem.setParams().setStreamKTileSchedulingMode(2); EXPECT_THROW(solution.streamK5EffectiveDynamic(problem, device), std::runtime_error) @@ -541,7 +529,7 @@ class StreamK5AutoOrigamiTest : public ::testing::TestWithParam const& info) { return info.param.suffix; }); @@ -612,7 +602,7 @@ TEST(StreamKSmCountTargetTest, SmCountTargetChangesReductionAndGrid) // streamK=3 on the gfx950 analytical device (256 CUs), k_split_aware selector. StreamK5AnalyticalEnv env; env.solution.sizeMapping.streamK = 3; - env.device.skDynamicGrid = static_cast(origami::grid_selection_t::k_split_aware); + env.device.skDynamicGrid = static_cast(origami::grid_selection_t::k_split_aware); // Make smCountTarget the sole grid budget source (AMDGPU defaults, explicit). env.device.skFixedGrid = 0; @@ -625,10 +615,10 @@ TEST(StreamKSmCountTargetTest, SmCountTargetChangesReductionAndGrid) { auto problem = makeGemmProblem(512, 512, 8192); - problem.setParams().setSmCountTarget(0); // use all device CUs (256) + problem.setParams().setSmCountTarget(0); // use all device CUs (256) const auto reductionAllCUs = env.solution.getSKReduction(problem, env.device); - problem.setParams().setSmCountTarget(32); // tight CU budget + problem.setParams().setSmCountTarget(32); // tight CU budget const auto reductionCapped = env.solution.getSKReduction(problem, env.device); EXPECT_EQ(reductionAllCUs, origami::reduction_t::parallel) @@ -684,7 +674,7 @@ TEST(StreamK5WorkspaceRegressionTest, QueryAndLaunchAgreeForDynamicMode) { StreamK5AnalyticalEnv env; env.solution.sizeMapping.workspaceSizePerElemC = 4; - env.solution.sizeMapping.streamKAtomic = 0; + env.solution.sizeMapping.streamKAtomic = 0; // Pick M/N so tiles % grid != 0 (partial tiles exist). // macroTile=128x128 → tiles = ceil(M/128)*ceil(N/128). @@ -718,7 +708,7 @@ TEST(StreamK5WorkspaceRegressionTest, StaticModeWorkspaceIsPartialTilesOnly) { StreamK5AnalyticalEnv env; env.solution.sizeMapping.workspaceSizePerElemC = 4; - env.solution.sizeMapping.streamKAtomic = 0; + env.solution.sizeMapping.streamKAtomic = 0; auto problem = makeGemmProblem(4096, 4224, 64); problem.setParams().setStreamKTileSchedulingMode(0); // OFF (static) @@ -743,14 +733,14 @@ TEST(StreamK5WorkspaceRegressionTest, SufficientWorkspacePreventsDPFallback) { StreamK5AnalyticalEnv env; env.solution.sizeMapping.workspaceSizePerElemC = 4; - env.solution.sizeMapping.streamKAtomic = 0; + env.solution.sizeMapping.streamKAtomic = 0; // Use dimensions that produce partial tiles. auto problem = makeGemmProblem(4096, 4224, 64); problem.setParams().setStreamKTileSchedulingMode(1); // ON (dynamic) size_t wsNeeded = env.solution.requiredWorkspaceSize( - [&] { + [&]{ auto p = problem; p.setWorkspaceSize(std::numeric_limits::max()); return p; @@ -765,14 +755,16 @@ TEST(StreamK5WorkspaceRegressionTest, SufficientWorkspacePreventsDPFallback) size_t wsActual = env.solution.requiredWorkspaceSize(problem, env.device); EXPECT_EQ(wsActual, wsNeeded) << "With workspace >= required, requiredWorkspaceSize must return the " - << "full amount (not 0 from DP fallback). ws=" << wsActual << " needed=" << wsNeeded; + << "full amount (not 0 from DP fallback). ws=" << wsActual + << " needed=" << wsNeeded; // Providing one byte less must trigger fallback (returns 0 for SK partial). auto problemShort = makeGemmProblem(4096, 4224, 64); problemShort.setParams().setStreamKTileSchedulingMode(1); problemShort.setWorkspaceSize(wsNeeded - 1); size_t wsShort = env.solution.requiredWorkspaceSize(problemShort, env.device); - EXPECT_EQ(wsShort, 0u) << "With workspace < required, must fall back (return 0 for partial)"; + EXPECT_EQ(wsShort, 0u) + << "With workspace < required, must fall back (return 0 for partial)"; } TEST(Sk3Sk5OffPartition512Test, NativeSk3MatchesSk5OffHostPack) @@ -796,7 +788,7 @@ TEST(Sk3Sk5OffPartition512Test, NativeSk3MatchesSk5OffHostPack) auto problemSk5 = makeGemmProblem(512, 512, 512); problemSk5.setParams().setStreamKTileSchedulingMode(0); // SK5-off - auto sk3Pack = computeStreamKHostPack(sk3Solution, problemSk3, *hardware); + auto sk3Pack = computeStreamKHostPack(sk3Solution, problemSk3, *hardware); auto sk5OffPack = computeStreamKHostPack(sk5Solution, problemSk5, *hardware); EXPECT_FALSE(sk5OffPack.effectiveDynamic); @@ -841,7 +833,8 @@ namespace return 0; try { - return origami::hardware_t::get_default_num_xcds(hipAMDGPU->analyticalHardware->arch); + return origami::hardware_t::get_default_num_xcds( + hipAMDGPU->analyticalHardware->arch); } catch(std::exception const&) { @@ -861,7 +854,8 @@ namespace return true; size_t baked = streamKBakedQueueCountRef(hardware); size_t numXCD = hipAMDGPU->analyticalHardware->NUM_XCD; - return baked == 0 || numXCD == 0 || (numXCD & (numXCD - 1)) != 0 || numXCD != baked; + return baked == 0 || numXCD == 0 || (numXCD & (numXCD - 1)) != 0 + || numXCD != baked; } // Mirror of ContractionSolution::streamKDynamicQueueSupported(). Returns @@ -903,11 +897,11 @@ namespace hip::HipAMDGPU makeGfx942DeviceWithXcd(size_t numXCD) { hip::HipAMDGPU device; - device.processor = AMDGPU::Processor::gfx942; - device.computeUnitCount = 304; - device.deviceName = "test-gfx942-xcd"; - device.analyticalHardware - = std::make_shared(makeGfx942HardwareWithXcd(numXCD)); + device.processor = AMDGPU::Processor::gfx942; + device.computeUnitCount = 304; + device.deviceName = "test-gfx942-xcd"; + device.analyticalHardware = std::make_shared( + makeGfx942HardwareWithXcd(numXCD)); return device; } @@ -942,8 +936,8 @@ namespace TEST(StreamKDynamicQueueXcdGateTest, RejectsMi300aSixXcd) { - hip::HipAMDGPU mi300a = makeGfx942DeviceWithXcd(6); - Hardware const& hw = mi300a; + hip::HipAMDGPU mi300a = makeGfx942DeviceWithXcd(6); + Hardware const& hw = mi300a; EXPECT_TRUE(streamKDynamicQueueUnsupportedRef(hw)) << "MI300A (NUM_XCD=6, not a power of two) must flag the dynamic-queue " "work-stealing path as unsupported"; @@ -951,8 +945,8 @@ TEST(StreamKDynamicQueueXcdGateTest, RejectsMi300aSixXcd) TEST(StreamKDynamicQueueXcdGateTest, AllowsMi300xEightXcd) { - hip::HipAMDGPU mi300x = makeGfx942DeviceWithXcd(8); - Hardware const& hw = mi300x; + hip::HipAMDGPU mi300x = makeGfx942DeviceWithXcd(8); + Hardware const& hw = mi300x; EXPECT_FALSE(streamKDynamicQueueUnsupportedRef(hw)) << "MI300X (NUM_XCD=8, power of two) must keep the dynamic-queue path"; } @@ -973,8 +967,8 @@ TEST(StreamKDynamicQueueXcdGateTest, RejectsGfx942FourXcdPowerOfTwoButMismatched TEST(StreamKDynamicQueueXcdGateTest, AllowsGfx950EightXcd) { // gfx950 (local MI355X) analytical hardware advertises 8 XCDs. - hip::HipAMDGPU gfx950 = makeHipDeviceWithAnalytical(makeGfx950AnalyticalHardware()); - Hardware const& hw = gfx950; + hip::HipAMDGPU gfx950 = makeHipDeviceWithAnalytical(makeGfx950AnalyticalHardware()); + Hardware const& hw = gfx950; EXPECT_FALSE(streamKDynamicQueueUnsupportedRef(hw)) << "gfx950 (NUM_XCD=8) must keep the dynamic-queue work-stealing path"; } @@ -1113,7 +1107,8 @@ namespace size_t clampedGrid(StreamK5AnalyticalEnv& env, ContractionProblemGemm& problem) { auto tiles = problem.getNumTiles(env.solution.sizeMapping, 1); - EXPECT_NE(tiles % StreamKFlagElements, 0u) << "grid must leave partial tiles to fix up"; + EXPECT_NE(tiles % StreamKFlagElements, 0u) + << "grid must leave partial tiles to fix up"; return env.solution.getSKGrid(problem, env.device, tiles, origami::reduction_t::tree); } } // namespace @@ -1177,116 +1172,3 @@ TEST(StreamKFlagBound, StreamK5StaticSubPathKeepsTheWholeBlock) << "SK5 on its static sub-path indexes from offset 0, so it keeps the " "whole block"; } -// SKLaunchGridLimitsTest -- tree-fixup must not force skGrid=tiles when that -// would overflow the uint32_t work-item limit (M=524288, N=98304, K=128 bench). - -TEST(SKLaunchGridLimitsTest, CapsGridWhenTilesReach2Pow24) -{ - ContractionSolution solution; - initBenchStreamK5Solution(solution, TensileLite::dim3(32, 96, 1), 32); - const size_t tpg = threadsPerWorkGroup(solution); - const size_t maxTilesBeforeCap = (size_t{std::numeric_limits::max()} + 1) / tpg; - - auto problem = makeGemmProblem(524288, 98304, 128); - AMDGPU device = makeDevice(_MI350_CHIP_ID, _SPX_CU, "mi350spx"); - device.skDynamicGrid = 0; - - auto tiles = problem.getNumTiles(solution.sizeMapping, 1); - ASSERT_EQ(tiles, maxTilesBeforeCap); - - size_t grid = solution.getSKGrid(problem, device, tiles, origami::reduction_t::tree); - EXPECT_EQ(grid, static_cast(_SPX_CU)) - << "Tree-fixup must cap skGrid to cuCount*occupancy, not full tile count"; - EXPECT_NE(grid, tiles); - EXPECT_LE(totalWorkItems(grid, tpg), - static_cast(std::numeric_limits::max())); -} - -TEST(SKLaunchGridLimitsTest, CapsGridWithAnalyticalOrigamiPath) -{ - origami::hardware_t hw = makeGfx950AnalyticalHardware(); - hip::HipAMDGPU device = makeHipDeviceWithAnalytical(hw); - - ContractionSolution solution; - initBenchStreamK5Solution(solution, TensileLite::dim3(32, 96, 1), 32); - const size_t tpg = threadsPerWorkGroup(solution); - const size_t maxTilesBeforeCap = (size_t{std::numeric_limits::max()} + 1) / tpg; - auto problem = makeGemmProblem(524288, 98304, 128); - auto tiles = problem.getNumTiles(solution.sizeMapping, 1); - ASSERT_EQ(tiles, maxTilesBeforeCap); - - size_t grid = solution.getSKGrid(problem, device, tiles, origami::reduction_t::tree); - EXPECT_EQ(grid, kGfx950AnalyticalCuCount) - << "Origami dynamic-grid path must still cap when tiles >= 2^24"; - EXPECT_NE(grid, tiles); -} - -TEST(SKLaunchGridLimitsTest, StillUsesDpFallbackBelowTileThreshold) -{ - ContractionSolution solution; - solution.sizeMapping.workGroupSize = TensileLite::dim3(128, 1, 1); - initBenchStreamK5Solution(solution, TensileLite::dim3(16, 16, 1), 128); - const size_t tpg = threadsPerWorkGroup(solution); - const size_t maxTilesBeforeCap = (size_t{std::numeric_limits::max()} + 1) / tpg; - - auto problem = makeGemmProblem(65536, 65552, 128); - AMDGPU device = makeDevice(_MI350_CHIP_ID, _SPX_CU, "mi350spx"); - device.skDynamicGrid = 0; - - auto tiles = problem.getNumTiles(solution.sizeMapping, 1); - ASSERT_EQ(tiles, 16781312u); - ASSERT_LT(tiles, maxTilesBeforeCap); - - size_t itersPerTile = problem.getItersPerTile(solution.sizeMapping); - ASSERT_GE(tiles * itersPerTile, maxTilesBeforeCap) - << "Tree-fixup trigger must fire even below the tile cap threshold"; - - size_t grid = solution.getSKGrid(problem, device, tiles, origami::reduction_t::tree); - EXPECT_EQ(grid, tiles) << "Non-overflow tile count must keep the DP fallback (skGrid=tiles)"; -} - -struct BenchLaunchLimitCase -{ - const char* label; - size_t mt0; - size_t mt1; - size_t depthU; - size_t workGroupSizeX; -}; - -class SKLaunchGridLimitsParamTest : public ::testing::TestWithParam -{ -}; - -TEST_P(SKLaunchGridLimitsParamTest, FailingMacroTilesStayWithinWorkItemLimit) -{ - auto const& param = GetParam(); - - ContractionSolution solution; - solution.sizeMapping.workGroupSize = TensileLite::dim3(param.workGroupSizeX, 1, 1); - initBenchStreamK5Solution(solution, TensileLite::dim3(param.mt0, param.mt1, 1), param.depthU); - - origami::hardware_t hw = makeGfx950AnalyticalHardware(); - hip::HipAMDGPU device = makeHipDeviceWithAnalytical(hw); - - auto problem = makeGemmProblem(524288, 98304, 128); - auto tiles = problem.getNumTiles(solution.sizeMapping, 1); - - size_t grid = solution.getSKGrid(problem, device, tiles, origami::reduction_t::tree); - EXPECT_LT(grid, tiles) << param.label << ": skGrid must not equal full tile count"; - - const size_t tpg = threadsPerWorkGroup(solution); - EXPECT_LE(totalWorkItems(grid, tpg), - static_cast(std::numeric_limits::max())) - << param.label << ": total work items must fit in uint32_t"; -} - -INSTANTIATE_TEST_SUITE_P(BenchSweep, - SKLaunchGridLimitsParamTest, - ::testing::Values(BenchLaunchLimitCase{"MT32x96x32", 32, 96, 32, 256}, - BenchLaunchLimitCase{"MT32x96x64", 32, 96, 64, 256}, - BenchLaunchLimitCase{"MT256x256x32", 256, 256, 32, 256}, - BenchLaunchLimitCase{"MT16x16x128", 16, 16, 128, 128}), - [](::testing::TestParamInfo const& info) { - return info.param.label; - }); diff --git a/projects/hipblaslt/tensilelite/tests/StreamKLaunchSummary_test.cpp b/projects/hipblaslt/tensilelite/tests/StreamKLaunchSummary_test.cpp index 88b8487583c0..48c5cb9f811d 100644 --- a/projects/hipblaslt/tensilelite/tests/StreamKLaunchSummary_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/StreamKLaunchSummary_test.cpp @@ -15,8 +15,8 @@ // pin that relationship between dynamicPartialsSlots, tiles%grid divisibility, // and whether a partials workspace is reserved. -#include #include +#include #include #include #include @@ -80,7 +80,6 @@ namespace solution.sizeMapping.macroTile = TensileLite::dim3(128, 128, 1); solution.sizeMapping.depthU = 64; solution.sizeMapping.matrixInstruction = {16, 16, 32, 1}; - solution.sizeMapping.workGroupSize = TensileLite::dim3(256, 1, 1); solution.sizeMapping.CUOccupancy = 1; solution.sizeMapping.workspaceSizePerElemC = 4; } @@ -424,7 +423,7 @@ TEST(StreamKLaunchSummaryTest, DpOnlySourceDistinguishesParamVsRuntime) ContractionSolution paramSol; initStreamKSolution(paramSol, 4); paramSol.sizeMapping.streamKForceDPOnly = 1; - auto paramProblem = makeGemmProblem(4096, 4224, 64); + auto paramProblem = makeGemmProblem(4096, 4224, 64); paramProblem.setWorkspaceSize(std::numeric_limits::max()); auto pd = paramSol.computeStreamKDecisions(paramProblem, env.device); @@ -537,7 +536,7 @@ TEST(StreamKLaunchSummaryTest, DynamicSlotsPositiveButDivisible_NoWorkspace) auto device = makeDevice(_MI350_CHIP_ID, _CPX_CU, "mi350cpx"); device.skDynamicGrid = 0; device.skTiles = 256; // override: number of split stream-k tiles - device.skSplit = 4; // override: k-split factor per tile + device.skSplit = 4; // override: k-split factor per tile auto d = solution.computeStreamKDecisions(problem, device);