Skip to content

Commit 3ef4894

Browse files
committed
fix(v4): device-aware Path-B status reason (CUDA EAGER bridge vs Metal MSL)
_path_b_status() for GDN/KDA now reports the TileLang-CUDA EAGER bridge on a CUDA host instead of the misleading 'via mx.fast.metal_kernel' string (which only describes the Apple Metal path). Apple still reports the hand-MSL path unchanged. Mirrors the device-aware switch in the forward (4a62306).
1 parent 2041f34 commit 3ef4894

2 files changed

Lines changed: 48 additions & 18 deletions

File tree

cppmega_v4/_tilelang/kda_paths.py

Lines changed: 21 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -59,19 +59,32 @@ def _path_a_call(*args, **kwargs):
5959

6060
def _path_b_status() -> PathStatus:
6161
try:
62-
importlib.import_module("cppmega_v4._tilelang.kda_path_b")
62+
b_mod = importlib.import_module("cppmega_v4._tilelang.kda_path_b")
6363
import mlx.core as mx
64-
if not hasattr(mx, "fast") or not hasattr(mx.fast, "metal_kernel"):
64+
# Device-aware: Apple -> hand-MSL mx.fast.metal_kernel; CUDA host ->
65+
# host _cuda_eager bridge (see kda_forward_path_b._device_can_run_metal).
66+
if b_mod._device_can_run_metal():
67+
if not hasattr(mx, "fast") or not hasattr(mx.fast, "metal_kernel"):
68+
return PathStatus(
69+
path="path_b", available=False,
70+
reason="mx.fast.metal_kernel not available on this build",
71+
)
6572
return PathStatus(
66-
path="path_b", available=False,
67-
reason="mx.fast.metal_kernel not available on this build",
73+
path="path_b", available=True,
74+
reason=(
75+
"hand-MSL KDA forward via mx.fast.metal_kernel; the "
76+
"autograd-aware wrapper kda_apply_path_b in kda_path_b_bwd.py "
77+
"also provides a real Metal backward (V <= 256)"
78+
),
6879
)
80+
from cppmega_mlx.nn._tilelang._cuda_eager import cuda_eager_available
81+
cuda_ok, cuda_reason = cuda_eager_available()
6982
return PathStatus(
70-
path="path_b", available=True,
83+
path="path_b", available=cuda_ok,
7184
reason=(
72-
"hand-MSL KDA forward via mx.fast.metal_kernel; the "
73-
"autograd-aware wrapper kda_apply_path_b in kda_path_b_bwd.py "
74-
"also provides a real Metal backward (V <= 256)"
85+
"KDA Path B forward via TileLang-CUDA EAGER bridge "
86+
"(kda_fwd_cuda_eager; Metal unavailable on this CUDA host). "
87+
f"{cuda_reason}"
7588
),
7689
)
7790
except Exception as exc:

cppmega_v4/_tilelang/linear_attention_paths.py

Lines changed: 27 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -70,20 +70,37 @@ def _path_a_call(*args, **kwargs):
7070

7171
def _path_b_status() -> PathStatus:
7272
try:
73-
importlib.import_module("cppmega_v4._tilelang.linear_attention_path_b")
74-
# Confirm mx.fast.metal_kernel is callable (Metal available).
75-
if not hasattr(mx, "fast") or not hasattr(mx.fast, "metal_kernel"):
73+
b_mod = importlib.import_module(
74+
"cppmega_v4._tilelang.linear_attention_path_b"
75+
)
76+
# Device-aware: on Apple the forward runs the hand-MSL
77+
# mx.fast.metal_kernel; on a CUDA host it routes the same GDN
78+
# recurrence through the host _cuda_eager bridge (see
79+
# gdn_forward_path_b._device_can_run_metal).
80+
if b_mod._device_can_run_metal():
81+
if not hasattr(mx, "fast") or not hasattr(mx.fast, "metal_kernel"):
82+
return PathStatus(
83+
path="path_b", available=False,
84+
reason="mx.fast.metal_kernel not available on this build",
85+
)
7686
return PathStatus(
77-
path="path_b", available=False,
78-
reason="mx.fast.metal_kernel not available on this build",
87+
path="path_b", available=True,
88+
reason=(
89+
"hand-MSL GDN forward via mx.fast.metal_kernel; the "
90+
"autograd-aware wrapper gdn_apply_path_b in "
91+
"linear_attention_path_b_bwd.py also provides a real Metal "
92+
"backward (max(K, V) <= 256)"
93+
),
7994
)
95+
# CUDA host: route through the TileLang-CUDA EAGER bridge.
96+
from cppmega_mlx.nn._tilelang._cuda_eager import cuda_eager_available
97+
cuda_ok, cuda_reason = cuda_eager_available()
8098
return PathStatus(
81-
path="path_b", available=True,
99+
path="path_b", available=cuda_ok,
82100
reason=(
83-
"hand-MSL GDN forward via mx.fast.metal_kernel; the "
84-
"autograd-aware wrapper gdn_apply_path_b in "
85-
"linear_attention_path_b_bwd.py also provides a real Metal "
86-
"backward (max(K, V) <= 256)"
101+
"GDN Path B forward via TileLang-CUDA EAGER bridge "
102+
"(gdn_fwd_cuda_eager; Metal unavailable on this CUDA host). "
103+
f"{cuda_reason}"
87104
),
88105
)
89106
except Exception as exc:

0 commit comments

Comments
 (0)