Skip to content

Commit 82aa315

Browse files
committed
fix(v4): device-aware Path-C status reason (CUDA EAGER bridge vs Metal DSL)
Mirrors the Path-B status fix (3ef4894). _path_c_status() for GDN/KDA reported a hardcoded "tilelang.compile(target='metal')" even on a CUDA host where Path-C actually routes through the host _cuda_eager bridge (target='cuda'). Now uses the path_c module's _device_can_run_metal() to report truthfully (RULE #1: honest reporting). Metal host text unchanged (verified non-regression).
1 parent b6a6177 commit 82aa315

2 files changed

Lines changed: 32 additions & 9 deletions

File tree

cppmega_v4/_tilelang/kda_paths.py

Lines changed: 17 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -141,19 +141,31 @@ def _path_b_call(*args, allow_fallback: bool = True, **kwargs):
141141

142142
def _path_c_status() -> PathStatus:
143143
try:
144-
from cppmega_v4._tilelang.kda_path_c import _path_c_runtime_status
144+
from cppmega_v4._tilelang.kda_path_c import (
145+
_path_c_runtime_status,
146+
_device_can_run_metal,
147+
)
145148
except Exception as exc:
146149
return PathStatus(
147150
path="path_c", available=False,
148151
reason=f"path_c module not importable: {exc}",
149152
)
150153
ok, reason = _path_c_runtime_status()
154+
# Device-aware (mirrors _path_b_status): Apple compiles target='metal';
155+
# CUDA host routes the same recurrence via the host _cuda_eager bridge.
156+
if _device_can_run_metal():
157+
detail = (
158+
"KDA Path C: TileLang DSL @T.prim_func → tilelang.compile("
159+
"target='metal', execution_backend='tvm_ffi')."
160+
)
161+
else:
162+
detail = (
163+
"KDA Path C via TileLang-CUDA EAGER bridge (kda_fwd_cuda_eager; "
164+
"target='cuda', Metal unavailable on this CUDA host)."
165+
)
151166
return PathStatus(
152167
path="path_c", available=ok,
153-
reason=(
154-
f"KDA Path C: TileLang DSL @T.prim_func → tilelang.compile("
155-
f"target='metal', execution_backend='tvm_ffi'). {reason}"
156-
),
168+
reason=f"{detail} {reason}",
157169
)
158170

159171

cppmega_v4/_tilelang/linear_attention_paths.py

Lines changed: 15 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -161,19 +161,30 @@ def _path_c_status() -> PathStatus:
161161
try:
162162
from cppmega_v4._tilelang.linear_attention_path_c import (
163163
_path_c_runtime_status,
164+
_device_can_run_metal,
164165
)
165166
except Exception as exc:
166167
return PathStatus(
167168
path="path_c", available=False,
168169
reason=f"path_c module not importable: {exc}",
169170
)
170171
ok, reason = _path_c_runtime_status()
172+
# Device-aware (mirrors _path_b_status): on Apple Path C compiles the
173+
# TileLang DSL for target='metal' (tvm_ffi); on a CUDA host it routes the
174+
# same recurrence through the host _cuda_eager bridge (target='cuda').
175+
if _device_can_run_metal():
176+
detail = (
177+
"GDN Path C: TileLang DSL @T.prim_func → tilelang.compile("
178+
"target='metal', execution_backend='tvm_ffi')."
179+
)
180+
else:
181+
detail = (
182+
"GDN Path C via TileLang-CUDA EAGER bridge (gdn_fwd_cuda_eager; "
183+
"target='cuda', Metal unavailable on this CUDA host)."
184+
)
171185
return PathStatus(
172186
path="path_c", available=ok,
173-
reason=(
174-
f"GDN Path C: TileLang DSL @T.prim_func → tilelang.compile("
175-
f"target='metal', execution_backend='tvm_ffi'). {reason}"
176-
),
187+
reason=f"{detail} {reason}",
177188
)
178189

179190

0 commit comments

Comments
 (0)