fix(tensilelite): Do not overwrite logic-file tiles with a zero custom-kernel macrotile

projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py

@@ -1790,16 +1790,48 @@ def getDivisorName(state, tC):
     return divisorName
 
   @staticmethod
-  def _assignCustomKernelParameters(state):
+  def _assignCustomKernelParameters(state: dict) -> None:
     """Minimal parameter setup for handwritten custom kernels.
 
     These kernels carry their own argument layout and don't go through the
     full assignDerivedParameters validation (which would reject them for
-    missing MatrixInstruction, etc.)."""
+    missing MatrixInstruction, etc.).
+
+    Args:
+      state: Solution state carrying a "CustomKernel" block. Mutated in place.
+
+    Returns:
+      None.
+
+    Raises:
+      RuntimeError: If a macrotile component is supplied by neither the
+        CustomKernel block nor the consuming logic file.
+    """
     ck = state["CustomKernel"]
-    state["MacroTile0"] = ck["macrotile"][0]
-    state["MacroTile1"] = ck["macrotile"][1]
-    state["DepthU"]     = ck["macrotile"][2]
+
+    # CustomKernels.py derives the macrotile from MatrixInstruction / MIWaveTile
+    # in custom.config, then from an MTxxx token in the kernel name. Any kernel
+    # that has neither infers 0 and must not overwrite
+    # the tile the consuming logic file already states. MacroTile0 == 0 reaches
+    # ContractionProblemGemm::getNumTiles() as an integer divide by zero, so
+    # dense enumeration (--algo_method 1) dies with SIGFPE as soon as that
+    # solution is costed. Kernels that already encode MT in the name
+    # or declare MI in custom.config are unchanged.
+    macrotile = list(ck.get("macrotile") or [])
+    macrotile += [0] * (3 - len(macrotile))
+    for i, key in enumerate(("MacroTile0", "MacroTile1", "DepthU")):
+      value = macrotile[i] if macrotile[i] > 0 else state.get(key, 0)
+      if not isinstance(value, int) or value <= 0:
+        raise RuntimeError(
+          f"Custom kernel '{ck.get('name', '?')}' has no usable {key}: it is absent "
+          f"from the kernel's custom.config, the kernel name encodes no MTxxx token, "
+          f"and the consuming logic file does not supply one.")
+      macrotile[i] = value
+      state[key]   = value
+    # Keep the block the C++ runtime deserializes in step with the solution:
+    # ContractionSolution reads customKernel.macrotile directly for custom-kernel
+    # tile and workspace sizing.
+    ck["macrotile"] = macrotile
 
     # Derive _GlobalAccumulation from GlobalSplitUAlgorithm so the C++
     # runtime sees a non-zero sizeMapping.globalAccumulation for GSU>1

projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_custom_kernel_host.py

@@ -68,6 +68,46 @@ def test_assign_custom_kernel_params_basic_derivation():
     assert state["_GlobalAccumulation"] is None  # GlobalSplitUAlgorithm == ""
 
 
+def test_assign_custom_kernel_params_unset_macrotile_falls_back_to_logic_file():
+    # Any handwritten kernel without MI fields and without an MTxxx name token
+    # reaches here with macrotile [0, 0, 0]. Taking that zero over the logic-file tile put
+    # a 0 in sizeMapping.macroTile, and getNumTiles() then divided by it.
+    state = _ck_state(MacroTile0=256, MacroTile1=256, DepthU=128)
+    state["CustomKernel"]["macrotile"] = [0, 0, 0]
+    Solution._assignCustomKernelParameters(state)
+
+    assert state["MacroTile0"] == 256
+    assert state["MacroTile1"] == 256
+    assert state["DepthU"] == 128
+    # The C++ runtime reads this block directly, so it must agree with the state.
+    assert state["CustomKernel"]["macrotile"] == [256, 256, 128]
+
+
+def test_assign_custom_kernel_params_macrotile_wins_over_logic_file():
+    # When the kernel does declare a tile it stays authoritative.
+    state = _ck_state(MacroTile0=64, MacroTile1=64, DepthU=16)
+    Solution._assignCustomKernelParameters(state)
+
+    assert [state["MacroTile0"], state["MacroTile1"], state["DepthU"]] == [128, 256, 64]
+
+
+def test_assign_custom_kernel_params_unresolvable_macrotile_raises():
+    # Neither source supplies a tile: fail the build rather than emit a library
+    # that divides by zero at solution-selection time.
+    state = _ck_state()
+    state["CustomKernel"]["macrotile"] = [0, 0, 0]
+    with pytest.raises(RuntimeError, match="no usable MacroTile0"):
+        Solution._assignCustomKernelParameters(state)
+
+
+def test_assign_custom_kernel_params_default_depthu_is_not_a_tile():
+    # Logic files default DepthU to -1; that must not be accepted as a tile.
+    state = _ck_state(MacroTile0=256, MacroTile1=256, DepthU=-1)
+    state["CustomKernel"]["macrotile"] = [0, 0, 0]
+    with pytest.raises(RuntimeError, match="no usable DepthU"):
+        Solution._assignCustomKernelParameters(state)
+
+
 def test_assign_custom_kernel_params_enable_mi_sets_wave_params():
     state = _ck_state()
     Solution._assignCustomKernelParameters(state)