Remove NDGridTileStrategy

This can use the standard `hl.tile()` codegen.

stack-info: PR: #209, branch: jansel/stack/60

Author: jansel

helion/_compiler/compile_environment.py

@@ -454,8 +454,10 @@ def from_config(self, config: Config, block_id: int) -> int | torch.SymInt:
 
 @dataclasses.dataclass
 class GridBlockSizeSource(BlockSizeSource):
+    """Block size source for grid indices - always has block_size=1 but marks as grid for special indexing"""
+
     def from_config(self, config: Config, block_id: int) -> int:
-        raise NotImplementedError
+        return 1
 
     def is_grid(self) -> bool:
         return True

helion/_compiler/tile_dispatch.py

@@ -15,7 +15,6 @@
 from helion._compiler.tile_strategy import CompactedShape
 from helion._compiler.tile_strategy import DeviceLoopState
 from helion._compiler.tile_strategy import FlattenedTileStrategy
-from helion._compiler.tile_strategy import NDGridTileStrategy
 from helion._compiler.tile_strategy import NDTileStrategy
 from helion._compiler.tile_strategy import TileStrategy
 
@@ -67,13 +66,7 @@ def _add_loop_strategy(
             config.l2_groupings, block_ids[0], 1
         )
 
-        if block_size_infos[0].is_grid():
-            strategy: TileStrategy = NDGridTileStrategy(
-                fn,
-                block_ids,
-                loop_order=loop_order,
-            )
-        elif block_size_infos[0].is_flattened(config):
+        if block_size_infos[0].is_flattened(config):
             block_size = functools.reduce(
                 operator.mul, [bs.from_config_assert(config) for bs in block_size_infos]
             )

helion/_compiler/tile_strategy.py

@@ -633,31 +633,6 @@ def select_pid_strategy(self) -> ProgramIDs:
         return super().select_pid_strategy()
 
 
-class NDGridTileStrategy(_BaseNDTileStrategy):
-    def __init__(
-        self,
-        fn: DeviceFunction,
-        block_ids: list[int],
-        loop_order: list[int],
-    ) -> None:
-        super().__init__(
-            fn=fn,
-            block_ids=block_ids,
-            block_size=[1] * len(block_ids),  # pyre-ignore[6]
-            loop_order=loop_order,
-        )
-
-    def mask_var(self, block_idx: int) -> str | None:
-        return None
-
-    def _setup_mask(
-        self,
-        *args: object,
-        **kwargs: object,
-    ) -> None:
-        return None
-
-
 class CompactedShape(NamedTuple):
     size_str: str
     user_indices: list[int]