Add epilogue subtiling

stack-info: PR: #948, branch: PaulZhang12/stack/14

Author: PaulZhang12

examples/matmul.py

@@ -28,6 +28,19 @@
 @helion.kernel(
     # static_shapes=True gives a performance boost for matmuls
     static_shapes=True,
+    config=helion.Config(
+        block_sizes=[64, 64, 64],
+        loop_orders=[[0, 1]],
+        l2_groupings=[4],
+        range_unroll_factors=[0, 1],
+        range_num_stages=[0, 3],
+        range_multi_buffers=[None, False],
+        range_flattens=[None, None],
+        num_warps=8,
+        num_stages=6,
+        indexing='tensor_descriptor',
+        pid_type='flat'
+    )
 )
 def matmul(
     x: Tensor,
@@ -44,6 +57,7 @@ def matmul(
     Returns:
         Tensor: Resulting matrix of shape [m, n].
     """
+
     m, k = x.size()
     k2, n = y.size()
     assert k == k2, f"size mismatch {k} != {k2}"

helion/_compiler/device_function.py

@@ -415,9 +415,14 @@ def tensor_arg(
     def tensor_descriptor_arg(
         self, fake_value: torch.Tensor, block_size: list[int | torch.SymInt]
     ) -> TensorDescriptorArg:
+        import re
         host_function = HostFunction.current()
         block_size_expr = ", ".join(map(self.literal_expr, block_size))
+        pattern = r'triton_helpers\.div_floor_integer\(([^,]+),\s*(\d+)\)'
+        replacement = r'\1 // \2'
+        block_size_expr = re.sub(pattern, replacement, block_size_expr)
         key = (fake_value, block_size_expr)
+
         if key not in self._tensor_descriptor_args:
             origin = host_function.tensor_to_origin[fake_value]
             desc_name = self.new_var(origin.suggest_var_name() + "_desc")

helion/_compiler/indexing_strategy.py

@@ -15,6 +15,7 @@
 from .. import exc
 from .._compat import get_tensor_descriptor_fn_name
 from .ast_extension import expr_from_string
+from .ast_extension import statement_from_string
 from .compile_environment import CompileEnvironment
 from .device_function import DeviceFunction
 from .host_function import HostFunction
@@ -385,21 +386,118 @@ def codegen_store(
         indexing = BlockedSubscriptIndexing.create(state, fake_tensor, subscript)
 
         # Apply permutation to the value being stored if needed
-        desc_arg = indexing.tensor_descriptor_arg(state)
+        # desc_arg = indexing.tensor_descriptor_arg(state, subtile=True)
         store_value = indexing.reshape_store(state, value)
 
-        if desc_arg.permutation is not None:
-            # Apply permutation to the value
-            store_value = expr_from_string(
-                f"tl.permute({{store_val}}, {desc_arg.permutation!r})",
-                store_val=store_value,
+        # if desc_arg.permutation is not None:
+        #     # Apply permutation to the value
+        #     store_value = expr_from_string(
+        #         f"tl.permute({{store_val}}, {desc_arg.permutation!r})",
+        #         store_val=store_value,
+        #     )
+        
+        if (
+            subtile_store := self._codegen_epilogue_subtile_store(
+                state, fake_tensor, indexing, store_value
             )
-
+        ) is not None:
+            return subtile_store
+        
         return expr_from_string(
             f"{indexing.tensor_descriptor(state)}.store({indexing.offsets_str_permuted(state)}, {{value}})",
             value=store_value,
         )
 
+    def _codegen_epilogue_subtile_store(
+        self,
+        state: CodegenState,
+        fake_tensor: torch.Tensor,
+        indexing: BlockedSubscriptIndexing,
+        store_value: ast.AST,
+    ) -> ast.AST | None:
+       # Currently support 2D tiles without permutations
+        if len(indexing.block_shape) != 2 or len(indexing.offsets) != 2:
+            return None
+
+        env = CompileEnvironment.current()
+        block_m, block_n = indexing.block_shape
+        try:
+            block_n_hint = env.size_hint(block_n)
+        except Exception:
+            return None
+
+        if block_n_hint % 2 != 0:
+            return None
+
+        device_fn = state.device_function
+        codegen = state.codegen
+
+        block_m_str = device_fn.literal_expr(block_m)
+        block_n_str = device_fn.literal_expr(block_n)
+        indexing.block_shape[1] //= 2
+        desc_arg = indexing.tensor_descriptor_arg(state)
+
+        if desc_arg.permutation is not None:
+            return None
+
+
+        block_n_half_str = f"({block_n_str} // 2)"
+
+        # Lift the store value into a temporary variable for reuse
+        acc_var = codegen.lift(store_value, prefix="acc")
+
+        reshape_expr = expr_from_string(
+            "tl.reshape({acc}, [{dim_m}, 2, {dim_half}])",
+            acc=acc_var,
+            dim_m=expr_from_string(block_m_str),
+            dim_half=expr_from_string(block_n_half_str),
+        )
+        reshape_var = codegen.lift(reshape_expr, prefix="acc")
+
+        permute_expr = expr_from_string(
+            "tl.permute({acc}, [0, 2, 1])",
+            acc=reshape_var,
+        )
+        permute_var = codegen.lift(permute_expr, prefix="acc")
+
+        acc0_name = codegen.tmpvar(prefix="acc")
+        acc1_name = codegen.tmpvar(prefix="acc")
+        codegen.add_statement(
+            statement_from_string(
+                f"{acc0_name}, {acc1_name} = tl.split({{acc}})",
+                acc=permute_var,
+            )
+        )
+        acc0 = expr_from_string(acc0_name)
+        acc1 = expr_from_string(acc1_name)
+
+        desc_name = indexing.tensor_descriptor(state)
+        offset0 = expr_from_string(indexing.offsets[0])
+        offset1 = expr_from_string(indexing.offsets[1])
+
+        # First subtile store
+        codegen.add_statement(
+            statement_from_string(
+                f"{desc_name}.store([{{off0}}, {{off1}}], {{value}})",
+                off0=offset0,
+                off1=offset1,
+                value=acc0,
+            )
+        )
+
+        offset1_shifted = expr_from_string(
+            "({offset} + {half})",
+            offset=expr_from_string(indexing.offsets[1]),
+            half=expr_from_string(block_n_half_str),
+        )
+
+        # Emit second subtile store as the expression returned to the caller
+        return expr_from_string(
+            f"{desc_name}.store([{{off0}}, {{off1}}], {{value}})",
+            off0=offset0,
+            off1=offset1_shifted,
+            value=acc1,
+        )
 
 class StackIndexingStrategy:
     """