From dff42623c6ddee567aa57a145531ec628c311453 Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Wed, 30 Sep 2026 19:10:15 +0000 Subject: [PATCH 1/2] fix(hipblaslt): launch every GSU slice of hand-written split-K kernels Since #11155, hand-written custom kernels launch through generateCustomCall, and for those without an explicit CustomKernel block the grid is inferred. Kernels that take no NumWorkGroups argument were given [TilesX, TilesY, Batch], which drops the GSU factor generateSingleCall multiplies into Y, so each tile launched one GSU slice. On gfx942 this hits the nine CustomGSUs_* kernels and the two fp8 x fp16 Equality shortname kernels: MultipleBufferSingleKernel kernels never reach their last-workgroup reduction and leave D unwritten, and MultipleBuffer kernels reduce a single valid slice into D. Both return HIPBLAS_STATUS_SUCCESS. Since #12599 fixed the zero macro tile, the default heuristic reaches this at the two tuned fp8 x fp16 shapes instead of crashing. - Add CustomGridSize::TilesYGSU (tiles.y * GSU) and infer [TilesX, TilesYGSU, Batch] for these kernels. - Fail library generation when a hand-written kernel can run with a GlobalSplitU other than 0 or 1 but its grid has no GSU term, and clear SupportUserGSU for such grids so a runtime GSU override is rejected during solution selection. - Throw from generateCustomCall when a split-K launch would use a grid with no GSU term, and log exceptions caught in runContractionProblem and makeArgument, so a library built before this fix returns an error instead of success. Tested on MI300X with a library built from the five affected gfx942 logic files. All 188 fp8 CustomGSUs_* solutions and 6 of the 8 fp16 ones pass CPU verification when run by index, and the default heuristic passes at both tuned fp8 x fp16 shapes. The two fp16 GSU 17 entries fail the same way on the pre-#11155 launch path; that is a separate kernel issue. TensileLite Python unit tests and tensilelite-tests pass. JIRA ID: ROCM-31692 Co-authored-by: Cursor --- .../amd_detail/rocblaslt/src/tensile_host.cpp | 2 + .../tensilelite/Tensile/CustomKernels.py | 5 +- .../Tensile/SolutionStructs/Solution.py | 21 ++++++++- .../Tests/unit/test_CustomKernelMetadata.py | 3 +- .../Tests/unit/test_custom_kernel_host.py | 47 +++++++++++++++++++ .../include/Tensile/ContractionSolution.hpp | 1 + .../tensilelite/src/ContractionSolution.cpp | 27 +++++++++++ .../tensilelite/tests/CustomKernel_test.cpp | 42 ++++++++++++++++- 8 files changed, 142 insertions(+), 6 deletions(-) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index 93fd141d4e57..be8b40fdd25c 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -3728,6 +3728,7 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle } catch(const std::exception& e) { + log_error(__func__, e.what()); } catch(...) { @@ -4110,6 +4111,7 @@ rocblaslt_status makeArgument(rocblaslt_handle handle, } catch(const std::exception& e) { + log_error(__func__, e.what()); } catch(...) { diff --git a/projects/hipblaslt/tensilelite/Tensile/CustomKernels.py b/projects/hipblaslt/tensilelite/Tensile/CustomKernels.py index c7929da99528..da5d0953ae14 100644 --- a/projects/hipblaslt/tensilelite/Tensile/CustomKernels.py +++ b/projects/hipblaslt/tensilelite/Tensile/CustomKernels.py @@ -351,8 +351,9 @@ def _buildCustomKernelFromMetadata(kernelName, fullYaml, kernelConfig): grid = ["TilesXYBatchGSU", "One", "One"] else: # Version 0 kernels rely on hardware gridDim for tile decomposition, - # so the grid must be multi-dimensional. - grid = ["TilesX", "TilesY", "Batch"] + # so the grid must be multi-dimensional. Like generated kernels, they + # take their GSU slices along gridDim.y. + grid = ["TilesX", "TilesYGSU", "Batch"] # Workspace is left unset here and derived in # Solution._assignCustomKernelParameters, which sees the logic file's diff --git a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py index d9f676e71eed..20d29af2280a 100644 --- a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py +++ b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py @@ -1873,7 +1873,9 @@ def _assignCustomKernelParameters(state: dict) -> None: Raises: RuntimeError: If a macrotile component is supplied by neither the - CustomKernel block nor the consuming logic file. + CustomKernel block nor the consuming logic file, or if the solution + can split K (GlobalSplitU other than 0 or 1) but its grid launches one + GSU slice per tile. """ ck = state["CustomKernel"] @@ -1901,6 +1903,23 @@ def _assignCustomKernelParameters(state: dict) -> None: # tile and workspace sizing. ck["macrotile"] = macrotile + # A split-K kernel reduces into D only once every GSU slice of a tile has + # arrived, so a grid without a GSU term would return success with D unwritten. + # GlobalSplitU -1 lets the runtime pick a split above 1. + gsu = state.get("GlobalSplitU", 1) + grid = ck.get("grid", []) + gridSplitsK = state.get("StreamK", 0) \ + or any(g in ("TilesYGSU", "TilesXYBatchGSU") for g in grid) + if gsu not in (0, 1) and not gridSplitsK: + raise RuntimeError( + f"Custom kernel '{ck.get('name', '?')}' runs with GlobalSplitU {gsu}, but its " + f"CustomKernel grid {grid} launches one GSU slice per tile; use TilesYGSU " + f"or TilesXYBatchGSU.") + # The same grid cannot honor a runtime GSU override either; clearing the flag + # rejects one during solution selection instead of failing at launch. + if not gridSplitsK: + state.setdefault("InternalSupportParams", {})["SupportUserGSU"] = False + # Derive _GlobalAccumulation from GlobalSplitUAlgorithm so the C++ # runtime sees a non-zero sizeMapping.globalAccumulation for GSU>1 # solutions. Without this the legacy beta-only kernel diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_CustomKernelMetadata.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_CustomKernelMetadata.py index ef24ca41acb4..aac2bf5711fa 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_CustomKernelMetadata.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_CustomKernelMetadata.py @@ -748,10 +748,11 @@ def test_build_from_metadata_numworkgroups_grid(): def test_build_from_metadata_default_multidim_grid(): + # Version 0 split-K kernels need one work group per GSU slice along Y. ck = _buildCustomKernelFromMetadata( "k", _kernel_yaml([_D_ARG]), {"MatrixInstruction": [16, 16, 16, 1]} ) - assert ck["grid"] == ["TilesX", "TilesY", "Batch"] + assert ck["grid"] == ["TilesX", "TilesYGSU", "Batch"] def test_build_from_metadata_macrotile_name_fallback(): diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py index dc2e7699c01a..08c1e46788e3 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py @@ -114,6 +114,53 @@ def test_assign_custom_kernel_params_default_depthu_is_not_a_tile(): Solution._assignCustomKernelParameters(state) +@pytest.mark.parametrize("grid", [ + ["TilesX", "TilesYGSU", "Batch"], + ["TilesXYBatchGSU", "One", "One"], +]) +def test_assign_custom_kernel_params_split_k_grid_is_accepted(grid): + state = _ck_state( + GlobalSplitU=16, + GlobalSplitUAlgorithm="MultipleBufferSingleKernel", + InternalSupportParams={"SupportUserGSU": True}, + ) + state["CustomKernel"]["grid"] = grid + Solution._assignCustomKernelParameters(state) + assert state["_GlobalAccumulation"] == "MultipleBufferSingleKernel" + assert state["InternalSupportParams"]["SupportUserGSU"] is True + + +@pytest.mark.parametrize("gsu", [16, -1]) +def test_assign_custom_kernel_params_split_k_without_gsu_grid_raises(gsu): + # A split-K kernel reduces into D only once every GSU slice has arrived, so a + # grid without a GSU term would launch one slice and leave D unwritten. + # GlobalSplitU -1 lets the runtime pick a split above 1. + state = _ck_state(GlobalSplitU=gsu, GlobalSplitUAlgorithm="MultipleBufferSingleKernel") + state["CustomKernel"]["grid"] = ["TilesX", "TilesY", "Batch"] + with pytest.raises(RuntimeError, match="launches one GSU slice per tile"): + Solution._assignCustomKernelParameters(state) + + +@pytest.mark.parametrize("gsu", [1, 0]) # 0: GSU disabled +def test_assign_custom_kernel_params_grid_without_gsu_term_rejects_user_gsu(gsu): + # Such a grid launches one GSU slice per tile, so a runtime GSU override has to + # be turned away during solution selection rather than fail at launch. + state = _ck_state(GlobalSplitU=gsu, InternalSupportParams={"SupportUserGSU": True}) + state["CustomKernel"]["grid"] = ["TilesX", "TilesY", "Batch"] + Solution._assignCustomKernelParameters(state) + assert state["InternalSupportParams"]["SupportUserGSU"] is False + + +def test_assign_custom_kernel_params_streamk_keeps_user_gsu(): + # Stream-K splits K through its own grid, so neither the GSU check nor the + # override flag applies to it. + state = _ck_state( + GlobalSplitU=16, StreamK=2, InternalSupportParams={"SupportUserGSU": True}) + state["CustomKernel"]["grid"] = ["TilesX", "TilesY", "Batch"] + Solution._assignCustomKernelParameters(state) + assert state["InternalSupportParams"]["SupportUserGSU"] is True + + def test_assign_custom_kernel_params_enable_mi_sets_wave_params(): state = _ck_state() Solution._assignCustomKernelParameters(state) diff --git a/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp b/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp index 296eb21503f2..7b7d6f2cd359 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp @@ -187,6 +187,7 @@ namespace TensileLite StreamKWithBatch, StreamKNoBatch, TilesXYBatchGSU, + TilesYGSU, CustomGridSize_Count, }; diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 088c24b0944a..befd2f92fb6b 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -156,6 +156,7 @@ namespace TensileLite case CustomGridSize::StreamKWithBatch: return "StreamKWithBatch"; case CustomGridSize::StreamKNoBatch: return "StreamKNoBatch"; case CustomGridSize::TilesXYBatchGSU: return "TilesXYBatchGSU"; + case CustomGridSize::TilesYGSU: return "TilesYGSU"; case CustomGridSize::CustomGridSize_Count: break; } @@ -175,6 +176,7 @@ namespace TensileLite {"StreamKWithBatch", CustomGridSize::StreamKWithBatch}, {"StreamKNoBatch", CustomGridSize::StreamKNoBatch}, {"TilesXYBatchGSU", CustomGridSize::TilesXYBatchGSU}, + {"TilesYGSU", CustomGridSize::TilesYGSU}, }; auto it = lookup.find(str); @@ -2898,6 +2900,9 @@ namespace TensileLite case CustomGridSize::TilesXYBatchGSU: dim = tiles.x * tiles.y * tiles.z * (gsu > 0 ? gsu : 1); break; + case CustomGridSize::TilesYGSU: + dim = tiles.y * (gsu > 0 ? gsu : 1); + break; case CustomGridSize::StreamKWithBatch: // generateCustomCall is only used for handwritten/external // custom kernels; Tensile-generated kernels are routed to @@ -2919,6 +2924,28 @@ namespace TensileLite assignGridSize(rv.numWorkGroups.y, customKernel.grid.y); assignGridSize(rv.numWorkGroups.z, customKernel.grid.z); + // A split-K kernel reduces into D only once every GSU slice of a tile has + // arrived, so a grid that launches one slice per tile would return success + // with D unwritten. + auto splitsK = [](CustomGridSize size) { + return size == CustomGridSize::TilesYGSU || size == CustomGridSize::TilesXYBatchGSU + || size == CustomGridSize::StreamKWithBatch + || size == CustomGridSize::StreamKNoBatch; + }; + if(gsu > 1 && !splitsK(customKernel.grid.x) && !splitsK(customKernel.grid.y) + && !splitsK(customKernel.grid.z)) + throw std::runtime_error(concatenate("Solution ", + kernelName, + " runs with GSU ", + gsu, + " but its custom-kernel grid [", + customKernel.grid.x, + ", ", + customKernel.grid.y, + ", ", + customKernel.grid.z, + "] launches one GSU slice per tile")); + bool enableCluster = (sizeMapping.clusterDim.x > 1 || sizeMapping.clusterDim.y > 1); bool hasNumWorkGroupsArg = std::any_of( customKernel.args.begin(), customKernel.args.end(), diff --git a/projects/hipblaslt/tensilelite/tests/CustomKernel_test.cpp b/projects/hipblaslt/tensilelite/tests/CustomKernel_test.cpp index 37bdebee436d..430d099696fa 100644 --- a/projects/hipblaslt/tensilelite/tests/CustomKernel_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/CustomKernel_test.cpp @@ -77,10 +77,13 @@ namespace "mi350spx"); } - uint32_t emittedSplitK(int16_t compiledGsu, int16_t runtimeGsu) + KernelInvocation splitKProbeCall(int16_t compiledGsu, + int16_t runtimeGsu, + CustomGridSize gridY = CustomGridSize::TilesYGSU) { ContractionSolution solution; configureProbeKernel(solution, {CustomArgType::uint32, CustomArgSemantic::SplitK}); + solution.customKernel.grid.y = gridY; solution.sizeMapping.globalSplitU = compiledGsu; auto problem = dummyProblem(); @@ -90,7 +93,12 @@ namespace ContractionInputs inputs; StreamKSettings sk; - auto invocation = solution.generateCustomCall(problem, inputs, device, sk); + return solution.generateCustomCall(problem, inputs, device, sk); + } + + uint32_t emittedSplitK(int16_t compiledGsu, int16_t runtimeGsu) + { + auto invocation = splitKProbeCall(compiledGsu, runtimeGsu); EXPECT_EQ(invocation.args.size(), sizeof(uint32_t)); uint32_t splitK = 0; @@ -141,6 +149,36 @@ TEST(CustomKernelTest, SplitKFallsBackToTheCompiledGsu) EXPECT_EQ(emittedSplitK(/*compiled*/ 16, /*runtime*/ 0), 4u); } +// A split-K kernel reduces into D only once every GSU slice of a tile has arrived, +// so TilesYGSU has to launch one row of work groups per slice, following the +// runtime-effective GSU like the SplitK argument does. +TEST(CustomKernelTest, TilesYGsuLaunchesEveryGsuSlice) +{ + auto const unsplit = splitKProbeCall(/*compiled*/ 1, /*runtime*/ 0).numWorkGroups; + + auto const compiled = splitKProbeCall(/*compiled*/ 16, /*runtime*/ 0).numWorkGroups; + EXPECT_EQ(compiled.x, unsplit.x); + EXPECT_EQ(compiled.y, 16 * unsplit.y); + EXPECT_EQ(compiled.z, unsplit.z); + + EXPECT_EQ(splitKProbeCall(/*compiled*/ 1, /*runtime*/ 4).numWorkGroups.y, 4 * unsplit.y); + EXPECT_EQ(splitKProbeCall(/*compiled*/ 0, /*runtime*/ 0).numWorkGroups.y, unsplit.y); + EXPECT_EQ( + splitKProbeCall(/*compiled*/ 1, /*runtime*/ 0, CustomGridSize::TilesY).numWorkGroups.y, + unsplit.y); +} + +// A grid with no GSU term would launch one slice per tile and report success with +// D unwritten, so a split-K launch on one has to fail instead. +TEST(CustomKernelTest, GridWithoutGsuTermRejectsSplitK) +{ + EXPECT_THROW(splitKProbeCall(/*compiled*/ 16, /*runtime*/ 0, CustomGridSize::TilesY), + std::runtime_error); + EXPECT_THROW(splitKProbeCall(/*compiled*/ 1, /*runtime*/ 4, CustomGridSize::TilesY), + std::runtime_error); + EXPECT_NO_THROW(splitKProbeCall(/*compiled*/ 1, /*runtime*/ 0, CustomGridSize::TilesY)); +} + // A custom kernel declares alpha and beta as 32-bit slots, so a narrower compute // type has to be widened before it is written or every argument after it shifts. // KernelArguments only widens an argument spelled exactly "alpha" or "beta". From 08f2e017304d8ea8775243566e31583b70b8573a Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Mon, 5 Oct 2026 14:58:30 +0000 Subject: [PATCH 2/2] fix(hipblaslt): judge split-K grids the same way at build time and launch The build-time GSU guard also exempted persistent kernels by execution policy, while generateCustomCall looks only at the grid, so a persistent kernel declared with a tile-count grid and a GlobalSplitU above 1 would build and then fail at launch. Every persistent grid already fails the tile-count test, so both guards now judge the grid alone. Co-authored-by: Cursor --- .../Tensile/SolutionStructs/Solution.py | 3 +-- .../Tensile/Tests/unit/test_custom_kernel_host.py | 15 ++++++++++----- 2 files changed, 11 insertions(+), 7 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py index ba98d40cb01d..d6f868c18a90 100644 --- a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py +++ b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py @@ -1918,9 +1918,8 @@ def _assignCustomKernelParameters(state: dict) -> None: # GlobalSplitU -1 lets the runtime pick a split above 1. gsu = state.get("GlobalSplitU", 1) grid = ck.get("grid", []) - tileCountGrid = all( + oneSlicePerTile = all( g in ("One", "TilesX", "TilesY", "Batch", "TilesXY", "TilesXYBatch") for g in grid) - oneSlicePerTile = tileCountGrid and not isPersistent(state) if gsu not in (0, 1) and oneSlicePerTile: raise RuntimeError( f"Custom kernel '{ck.get('name', '?')}' runs with GlobalSplitU {gsu}, but its " diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py index 8eaaa45da2cb..60e5c36f126b 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py @@ -131,12 +131,17 @@ def test_assign_custom_kernel_params_split_k_grid_is_accepted(grid): assert state["InternalSupportParams"]["SupportUserGSU"] is True -@pytest.mark.parametrize("gsu", [16, -1]) -def test_assign_custom_kernel_params_split_k_without_gsu_grid_raises(gsu): +@pytest.mark.parametrize("over", [ + {"GlobalSplitU": 16}, + {"GlobalSplitU": -1}, # lets the runtime pick a split above 1 + # generateCustomCall judges the grid alone, so a persistent kernel with a + # tile-count grid has to be rejected here too rather than at launch. + {"GlobalSplitU": 16, "TileProcessingStrategy": "StreamK"}, +]) +def test_assign_custom_kernel_params_split_k_without_gsu_grid_raises(over): # A split-K kernel reduces into D only once every GSU slice has arrived, so a # grid without a GSU term would launch one slice and leave D unwritten. - # GlobalSplitU -1 lets the runtime pick a split above 1. - state = _ck_state(GlobalSplitU=gsu, GlobalSplitUAlgorithm="MultipleBufferSingleKernel") + state = _ck_state(GlobalSplitUAlgorithm="MultipleBufferSingleKernel", **over) state["CustomKernel"]["grid"] = ["TilesX", "TilesY", "Batch"] with pytest.raises(RuntimeError, match="launches one GSU slice per tile"): Solution._assignCustomKernelParameters(state) @@ -153,7 +158,7 @@ def test_assign_custom_kernel_params_grid_without_gsu_term_rejects_user_gsu(gsu) @pytest.mark.parametrize("strategy,grid", [ - ("StreamK", ["TilesX", "TilesY", "Batch"]), + ("StreamK", ["StreamKWithBatch", "One", "One"]), ("DataParallel", ["PersistentGrid", "One", "One"]), ("None", ["PersistentNoBatch", "One", "One"]), ])