[Metal] FP8 vector cast lanes 2/3/4 (extends storage-only FP8)

[apstenku123]

Summary

Extends the storage-only Metal FP8 codegen to handle vectorized casts at IR
lanes 2 / 3 / 4. Without this change, every T.Cast("float16x4", fp8_x4)
emitted by upstream TileLang DSL programs raises
LOG(FATAL): Vector FP8 casts (lanes=4) are not yet supported, forcing
callers to manually scalarise the cast and giving up the IR-level vector
type for any subsequent pass (vectorize, fragment-to-simdgroup, etc.).

This PR is the TileLang/tvm submodule half that mirrors the same
change into the vendored TVM codegen at
src/target/source/codegen_metal.{cc,h}. The companion TileLang
supermodule half is filed at
https://github.com/tile-ai/tilelang/pulls (search
cppmega/metal-fp8-vector-cast). Both halves only share helper names;
they can land independently but should be merged in tandem so the
TileLang supermodule and its vendored 3rdparty/tvm checkout stay in sync
with this codepath.

What this changes

Adds an enable_fp8_vector_ codegen flag and a new
PrintFP8VectorPrelude(...) that emits inline MSL helpers that wrap the
existing scalar helpers (__tvm_fp8_e4m3_to_half, etc.) per lane:

inline half4 __tvm_fp8_e4m3_to_half_v4(uchar4 x) {
  return half4(__tvm_fp8_e4m3_to_half(x.x), __tvm_fp8_e4m3_to_half(x.y),
               __tvm_fp8_e4m3_to_half(x.z), __tvm_fp8_e4m3_to_half(x.w));
}

Mirrors are emitted for _v2 / _v3, plus the reverse direction
(half -> fp8) and the e5m2 variant. The compiler is free to scalarise
back into per-lane calls; the goal here is to preserve the IR-level
vector type so subsequent passes can keep their vector loads and stores
and the downstream MSL is uchar4-typed instead of uchar arrays.

Finish() is updated to splice the vector prelude after the scalar
prelude when at least one vector FP8 cast is encountered during codegen.

Wider lanes (8 / 16) keep the existing LOG(FATAL) with a sharper
message — those widths print as uint2 / uint4 packed storage and
need an out-pointer ABI to be wired through; callers should lower them
to scalar casts upstream.

Why Apple Silicon needs software FP8 emulation

Apple Silicon (M1 through M4 Max, including the M5 NAX which is
FP16/INT8 only) has no native FP8 ALU. FP8 is realised as uchar
storage with explicit dequantize-on-load / quantize-on-store; the
encoding mirrors the OCP "OFP8 Formats for Deep Learning" v1.0 spec
(E4M3 finite-only, E5M2 IEEE-style with NaN/Inf).

The vector helpers in this PR are inline-trivial wrappers around the
scalar helpers that landed in the storage-only PR — no new conversion
math. Their value is purely codegen: the IR-level vector type is
preserved so the rest of the lowering pipeline can vectorise.

Path C consumer evidence (vector lanes matter)

The downstream cppmega.mlx project's Path C TileLang FP8 vecmat kernel
(cppmega_mlx/nn/_tilelang/fp8_vecmat_path_c.py) explicitly uses
T.alloc_local((4,), "float8_e4m3") and a T.vectorized(4) inner loop
over packed FP8 weights. Without this PR, that kernel cannot be lowered
on Metal — the FP8 cast inside the K-loop hits the lanes=4 FATAL.
With this PR, the cast lowers and the resulting MSL preserves
uchar4-typed loads through the K-loop.

Dependency

This PR stacks on two prereqs:

1. tilelang_metal_fp8 storage-only patch (parallel
   [Metal] FP8 storage-only emulation (uchar storage + LUT decode helpers)
   PR being filed against this same repo). That patch adds
   PrintFP8Prelude, enable_fp8_, and the scalar
   __tvm_fp8_*_to_half / __tvm_half_to_fp8_* helpers that the
   vector helpers in this PR call. Reviewers will need that patch
   applied first; the branch in this PR includes it as the first commit
   [Metal] FP8 storage-only emulation ... [prereq] for self-contained
   review.
2. Base commit 0e15b274bce8b46f971abf5ac390e844aa6acee5 (the
   submodule pin used by tile-ai/tilelang's
   metal-gemm-upstream-rebase branch).

When the storage-only PR merges, the prereq commit on this branch
should be rebased away. Before that, this branch is reviewable as
2-commits stacked.

Test plan

• git apply --check clean against base
  0e15b274bce8b46f971abf5ac390e844aa6acee5 with the storage-only
  prereq applied first
• git apply --reverse --check clean for both commits in sequence
  (round-trip verified)
• xcrun --sdk macosx metal -c compile of any prim_func with vector
  FP8 cast (lanes 2/3/4) lowers to MSL using the new vector helpers,
  not scalar fallback
• Direct probe /tmp/test_fp8_vector_cast.py: lanes=4 cast lowers
  and the resulting MSL contains __tvm_fp8_e4m3_to_half_v4 with
  uchar4 typed loads
• cppmega.mlx tilelang test suite: 134 passed, 0 regressions

[Copilot]

Copilot wasn't able to review this pull request because it exceeds the maximum number of files (300). Try reducing the number of changed files and requesting a review from Copilot again.