Fix gfx1250 TDM illegal memory access

projects/hipblaslt/tensilelite/Tensile/Components/TensorDataMover.py

@@ -524,6 +524,9 @@ def setIterations(self, group2: int | str, sgprNumIters: int | str) -> Module:
         mod.add(SLShiftLeftB32(sgpr(f"{group2}+3"), hex(16), sgpr(f"{group2}+3")))
         return mod
 
+    # pad_interval is log2(LdsBlockSizePerPad//4)-1 and must fit in 3 bits (<=7).
+    MAX_NON_ITERATE_LBSPP_BYTES: int = 1024
+
     @staticmethod
     def calPadInterval(ldsBlockSizePerPad: int) -> int:
         ldsBlockDwordsPerPad = ldsBlockSizePerPad // 4 # bytes to dwords
@@ -537,3 +540,29 @@ def calPadAmount(ldsPadSize: int) -> int:
         ldsPadDwords =  ldsPadSize // 4 # bytes to dwords
         assert ldsPadDwords > 0
         return ldsPadDwords - 1
+
+    @staticmethod
+    def needsIterateModeForPad(ldsBlockSizePerPad: int, depth_u: int, bpe: float) -> bool:
+        """Return True when pad_interval non-iterate TDM must not be used.
+
+        When the natural VW=8 LBSPP would exceed the 1024 B pad_interval cap, prefer
+        iterate mode instead of halving VectorWidth — halved-VW non-iterate TDM on
+        gfx1250 produces illegal global accesses (BBS_NN_426 class).
+
+        Args:
+            ldsBlockSizePerPad: Lds block size per pad in bytes (256-aligned).
+            depth_u: DepthU unroll depth for this tensor.
+            bpe: Bytes per element for the tensor mac/data type.
+
+        Returns:
+            True if iterate_enable should be set for this tensor.
+        """
+        if ldsBlockSizePerPad == 0:
+            return False
+        if ldsBlockSizePerPad > TensorDataMoverLoad.MAX_NON_ITERATE_LBSPP_BYTES:
+            return True
+        full_vw_bytes = int(depth_u * bpe * 8)
+        if full_vw_bytes <= 0:
+            return False
+        full_vw_lbspp = ((full_vw_bytes + 255) // 256) * 256
+        return full_vw_lbspp > TensorDataMoverLoad.MAX_NON_ITERATE_LBSPP_BYTES

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

@@ -3223,8 +3223,12 @@ def depthUIteration(
               continue
             vw = state["VectorWidth%s" % tc]
             bpe_tc = state["ProblemType"]["MacDataType%s" % tc].numBytes()
-            lbspp = roundUpToNearestMultiple(int(state["_DepthU%s" % tc] * bpe_tc * vw), 256)
-            if lbspp > 1024:
+            lbspp_cfg = state["LdsBlockSizePerPad%s" % tc]
+            if lbspp_cfg == -1:
+              lbspp = roundUpToNearestMultiple(int(state["_DepthU%s" % tc] * bpe_tc * vw), 256)
+            else:
+              lbspp = lbspp_cfg
+            if TensorDataMoverLoad.needsIterateModeForPad(lbspp, state["_DepthU%s" % tc], bpe_tc):
               state["_TDMIterateMode%s" % tc] = True
           # Resolve -1 into the concrete mask for kernel naming. backupValues
           # restores -1 before each DepthU candidate, so this overwrite is safe.
@@ -3251,8 +3255,10 @@ def depthUIteration(
                    "TDMIterateMode %s requires TLU%s to be False" % (tc, tc))
             return
 
-        # Stage 2: for non-iterate tensors, halve auto-derived VW until LBSPP
-        # fits the pad_interval 1024 B limit.
+        # Stage 2: for non-iterate tensors with auto LBSPP, halve auto-derived VW until
+        # LBSPP fits pad_interval encoding. Skip when LdsBlockSizePerPad is explicit
+        # (halving VW there does not change LBSPP). When full VW=8 LBSPP would exceed
+        # the non-iterate cap, prefer iterate (needsIterateModeForPad) over halving.
         multiple = 256
         for tc in ["A", "B"]:
           if state.get("_TDMIterateMode%s" % tc, False):
@@ -3261,11 +3267,20 @@ def depthUIteration(
             continue
           if not state["UnrollMajorLDS%s" % tc]:
             continue
+          if state["LdsBlockSizePerPad%s" % tc] != -1:
+            continue
           vw = state["VectorWidth%s" % tc]
           if vw <= 1:
             continue
           tmpBpe = state["ProblemType"]["MacDataType%s" % tc].numBytes()
           depthU_tc = state["_DepthU%s" % tc]
+          lbspp = roundUpToNearestMultiple(int(depthU_tc * tmpBpe * vw), multiple)
+          # When TDMIterateMode is explicit (not -1), Stage 1 does not auto-select
+          # iterate; block VW halving and force iterate for the gfx1250 broken path.
+          if iterModeMask != -1 and TensorDataMoverLoad.needsIterateModeForPad(lbspp, depthU_tc, tmpBpe):
+            state["_TDMIterateMode%s" % tc] = True
+            state["TDMIterateMode"] = state["TDMIterateMode"] | (1 if tc == "A" else 2)
+            continue
           origVw = vw
           while vw > 1:
             candidate = roundUpToNearestMultiple(int(depthU_tc * tmpBpe * vw), multiple)
@@ -4360,7 +4375,8 @@ def calSwizzlePackK(state, tc):
                 extraComment = ": DepthU(%u) < Min-DU for swizzleB + LSU(%u)"%(depthUB, state["LocalSplitU"])
 
           # Subtile bypasses the 1024B pad_interval limit via iterate mode;
-          # non-subtile TDM handles it earlier via _TDMIterateMode + VW halving.
+          # non-subtile TDM handles it earlier via _TDMIterateMode and VW halving
+          # (iterate preferred over halving when needsIterateModeForPad applies).
         # this depthU is valid, done unless user wants to double (for TN)
         if validDepthU:
           state["DepthU"] = depthU

projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_streamk_cluster_multicast_gfx1250_char.ambr

@@ -6,31 +6,31 @@
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg9gjgV9hsSndadJjhuMT7as1DTv6SwnrqLiCclIme5ls=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgBEgEr51tsEQXpadk0WONxD0PUCUYDEfZDnEm0rDBwEY=',
       'err': 0,
     }),
     dict({
       'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgB_rKUEOGTAXS6Zuu59JqBp-HJwL9j2MqiovnkI8OE28=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgIEhOJNkP7_Vk84oQKwf_0Q0P5WloFcxaaItngYRl8HY=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgTbWHZF76AK958G-H8kRhEDEXNmXZ9WXhAMN9ng2IKw4=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgTHJzvzP0uCnEdasjFwMrj6gKqWmTtkOd3P1FLLc8ChQ=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgifyUzyOWqkpx7ZJNhDkgpv6dIQP346P8AZEkBTvzE7w=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgaLxF2UAjzelMh40A0tG4hW-W2gRuQN9s-rIk4CYuig0=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgnpeJWEWXMDPXZSJEyfY8Yr4r4Ym_97Q7zIcos1c6gwo=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgnpeJWEWXMDPXZSJEyfY8Yr4r4Ym_97Q7zIcos1c6gwo=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgodt8t83TM-N12cDOnoKg__albvlwU9LY1JgKr8Mlp_o=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgodt8t83TM-N12cDOnoKg__albvlwU9LY1JgKr8Mlp_o=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgq9X6V0EaISEguFhJQDJP8fB2AQg0ZULrLLPJirzC5tA=',
       'err': 0,
     }),
   ])
@@ -46,27 +46,27 @@
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgL4W_yBY6jyTQxjkDMj-qNvV-0opjiHHCbkah8KwCs7s=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgUAlVg601_FOYM70fMZ3Ps80tvg9dibaEthSREanF6go=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgLgwKXBVCRvRXB5RN9ISkynm8bYcurB0p3gE8kI1dkE4=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgUxlusQzSNiB_QUiHLNLLgF8CF1liXpPu1xGVQigzp50=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgUxlusQzSNiB_QUiHLNLLgF8CF1liXpPu1xGVQigzp50=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgkVP3jqcsmTuvqiAaPpYmnT9F8yRESQ5si6VCMGf9jR4=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgYPnMjn8WTWA_Oo4GDTlywtTbd1lGy2xsuJYAuu2IZWw=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgu7S-HBSgwf98qbOiQdZGF5TA5cYQb8FLAKejQZl6hcI=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg_axr_wF9qcNCkXRHjWethcUq18jXigESmgfl1jXI26U=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgz25_rye8DebQd4dFLSA1n1FuXm56CppIvGQRrS0_4lM=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgkVP3jqcsmTuvqiAaPpYmnT9F8yRESQ5si6VCMGf9jR4=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgzZJke1r3gd7HQkdpPDV4kxWSYHFeoPK0WZCgITJuauM=',
       'err': 0,
     }),
   ])
@@ -77,10 +77,6 @@
       'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg5N8ZqbnwzUt-4GDI-oxs5UmqLzHe6HSFSlTShem8_SE=',
       'err': 0,
     }),
-    dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg6AaPSexOhERiASBqm1R1RWMEMAaebVQwhxIrxc7h918=',
-      'err': 0,
-    }),
     dict({
       'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgFhDpgOA40F3oTAhrrTh3CalXjw0_GGF8k56a5TQNqq4=',
       'err': 0,
@@ -90,19 +86,23 @@
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgLwXm_Ol7CLqvhUBISj6DNLfmY90pQtGeBPxlanLRrWU=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgPLgVBBk_XClg8z_BnXM3rDyFpeGDzo3JiBIrUwoVO-g=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgW_K9FMjpgO7YBNPE_3mGc3QTJVQvjJWXmIGvAZREMKo=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgY4ujb2yHCxfWaVBqkOAc_zVM4BoFdLDh3od8fxmHZNw=',
       'err': 0,
     }),
     dict({
       'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgp25tUTG4Mb2s-i9zC1LhVO5OUQx1BzAtNGynkowUr48=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArguUYynrZQUDgsy5KQwKyP2VqK0ksczV1IclPPhqmqDBQ=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgs1SHiIGbN3pHluEIGGEqP4O4u7UQ6nIUxebtPSUV6BE=',
+      'err': 0,
+    }),
+    dict({
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgslDLnx1Zh3A87YOFQKrEt_Y46XzFpzZrfvh12LZAYc0=',
       'err': 0,
     }),
   ])
@@ -114,27 +114,27 @@
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgFhtrlQbgj2FkFkP6HnoNq6eal7qNe66_bB2saXIBob0=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg4CDEYU0bHYb948UAiz8E2MkBIOE_sFZaVt6qpav2-m8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgMo3wr5V0tQSUTRX-M5NLKgo5n9vLW6Ov5XjDouNE3BU=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgCSgxLcrg2AGV16Ie4RpOdL6jR5ifR7aMJGBy2_ANnIA=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgWUPoRPPVJtBObNcUqdQcOfcO7HLEkyBuU5FOsqrFtIE=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgFhtrlQbgj2FkFkP6HnoNq6eal7qNe66_bB2saXIBob0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg_FkRXm_Y5SbwkwbKL_8CnsWMoqcv1JhEDDl892UpE1I=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgMhYk1BsKySFPTYGbQnHw83LJmmkmvg2BcDwwnW3NDcA=',
       'err': 0,
     }),
     dict({
       'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArggFKAU1D9Qf-6DPzGzQA7TgiLT_pprtt3it1WQZCgzhQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArggQtuuf3-krRBnMi4jlHJ3YqVkxuaTXxqxGA2KLJe4go=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgn-T_M0dnRoowXvkNRTgXzGrgcmvKjV5t-Xg6jJPgdOE=',
       'err': 0,
     }),
     dict({

projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_tdm_prefetchgl2.yaml

@@ -63,7 +63,7 @@ BenchmarkProblems:
         - ForceDisableShadowInit: [True]
         - MatrixInstruction:
           - [16, 16, 128, 1, 1, 4, 4, 2, 2]
-        - DepthU: [256]
+        - DepthU: [128]  # VW=8 at LBSPP=1024 non-iterate; DU=256 forces TDM iterate and overflows SGPR budget
         - ClusterLocalRead: [0]
         - TransposeLDS: [-1]
         - LdsPadA: [-1]

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

@@ -0,0 +1,105 @@
+################################################################################
+#
+# Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+#
+# SPDX-License-Identifier: MIT
+################################################################################
+import pytest
+
+from Tensile.Components.TensorDataMover import TensorDataMoverLoad
+
+pytestmark = pytest.mark.unit
+
+
+class TestTensorDataMoverPadHelpers:
+    """Unit tests for TDM pad / iterate-mode selection helpers."""
+
+    def test_cal_pad_interval_1024(self) -> None:
+        """LBSPP=1024 encodes as pad_interval=7 (the hardware maximum).
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.calPadInterval(1024) == 7
+
+    def test_needs_iterate_du128_vwb8(self) -> None:
+        """DU=128 BF16 VWB=8 exceeds the non-iterate pad cap.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(2048, 128, 2.0) is True
+
+    def test_needs_iterate_du128_vwb4_halved(self) -> None:
+        """Halved VWB=4 at DU=128 must use iterate when full VW=8 LBSPP exceeds cap.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(1024, 128, 2.0) is True
+
+    def test_no_iterate_du64_vwb8(self) -> None:
+        """DU=64 BF16 VWB=8 stays on the non-iterate path.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(1024, 64, 2.0) is False
+
+    def test_no_iterate_du64_vwb4_when_full_vw_fits(self) -> None:
+        """VWB=4 at DU=64 is safe when full VW=8 LBSPP still fits the 1024 B cap.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(512, 64, 2.0) is False
+
+    def test_no_iterate_lbspp_zero(self) -> None:
+        """LBSPP=0 disables padding; iterate mode must not be forced.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(0, 64, 2.0) is False
+
+    def test_needs_iterate_fp4_non_integer_row_bytes(self) -> None:
+        """fp4 (bpe=0.5) with non-integer DepthU*bpe must not under-estimate VW=8 LBSPP.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(1024, 257, 0.5) is True
+
+    def test_needs_iterate_bf6_halved_path(self) -> None:
+        """bf6/fp6 (bpe=0.75) halved-VW path must detect full VW=8 LBSPP over cap.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(768, 171, 0.75) is True
+
+    def test_no_iterate_fp4_at_cap(self) -> None:
+        """fp4 at exactly 1024 B VW=8 LBSPP stays non-iterate.
+
+        Returns:
+            None
+        """
+        assert TensorDataMoverLoad.needsIterateModeForPad(1024, 256, 0.5) is False