sycl: int8 prompt experts default to auto on Battlemage (bmg-g21 / bmg-g31 by architecture id), FP16 elsewhere; STRATA_PREFILL_MMQ=0 off, =1 forced; launch failure or missing kernels fall back to the FP16 chain with one log line; load-time log line; mode-choice unit test; docs

docs/INTEL.md

@@ -900,13 +900,13 @@ difference 1.0 (the chunk boundary moves where the QSA block selection is made,
 two runs of the same size are, KL 0). With the 32K context (own prompt buffers, no borrowing) a bigger chunk is +20% on a 16K
 prompt (1,236 -> 1,478-1,482) but takes 1,397 cache slots (10,926 -> 9,529) and decode after it falls 9-16%, so 4,096 stays there.
 
-## Arc Pro B70: the fused int8 expert path for the prompt (2026-10-18, branch f-151-b70, opt-in)
+## Arc Pro B70: the fused int8 expert path for the prompt (2026-10-18, branch f-151-b70, default on for Battlemage)
 
-`STRATA_PREFILL_MMQ=1` (or `auto`) computes the routed experts of the prompt with ESIMD kernels that decode the quantized weights in
+On a Battlemage GPU (B70, B580, B570: detected by the device architecture id, `bmg-g21` and `bmg-g31`, not by name) the engine computes the routed experts of the prompt with ESIMD kernels that decode the quantized weights in
 registers straight into the s8 DPAS B operand (no FP16 copy of the expert, 2.1 MB of weights read instead of 9.8 MB written and read
 back) and multiply them by q8_1-quantized activations; `sycl/src/prefill/esimd_moe.dp.cpp`, behaviour note and switches in
 [INTEL_INT8_EXPERTS.md](INTEL_INT8_EXPERTS.md). Gate/up IQ3_S, IQ3_XXS, IQ2_S and down IQ4_NL, Q2_0 are covered (47 of the 48 layers of
-the mixed IQ3_S pack; the other layer, IQ4_XS gate/up, and any other format keep the FP16 path). Off by default: the activations are
+the mixed IQ3_S pack; the other layer, IQ4_XS gate/up, and any other format keep the FP16 path). Default (STRATA_PREFILL_MMQ unset) is `auto` on Battlemage and the FP16 chain on every other GPU (Alchemist A750 / A770, Lunar Lake, other vendors); `=0` turns it off, `=1` forces it, `=auto` is the default. The activations are
 quantized, so the bits differ from the FP16 path. `auto` times both chains on this GPU at the first layer of each format and takes int8
 below the measured crossover (130-250 rows per expert here), per group of experts.
 

docs/INTEL_INT8_EXPERTS.md

@@ -13,9 +13,9 @@ never exists and the weights are read once, compressed (2.1 MB per expert instea
 ## Switches
 
 * `STRATA_PREFILL_MMQ=1` (alias `STRATA_PF_I8=1`): every group of experts whose formats are supported takes the int8 chain.
-* `STRATA_PREFILL_MMQ=0` (or unset, the default): the FP16 path everywhere. The name is MMQ's because the other backends' int8
+* `STRATA_PREFILL_MMQ=0`: the FP16 path everywhere. The name is MMQ's because the other backends' int8
   prompt-expert path is MMQ; here it is this kernel.
-* `STRATA_PREFILL_MMQ=auto`: per group, int8 while the group's mean rows per expert are under the crossover measured on this GPU.
+* `STRATA_PREFILL_MMQ=auto` (the Battlemage default): per group, int8 while the group's mean rows per expert are under the crossover measured on this GPU.
   At the first layer of each expert format the engine times both chains on synthetic experts (random blobs of the layer's format and
   shape, uneven row counts like a real group) at 16, 32, 64, 128 and 256 mean rows per expert, interpolates where the FP16 chain
   becomes the faster one, flags the result "clear" when the cost ratio falls steadily and the bracketing samples are 5% away from
@@ -24,7 +24,7 @@ never exists and the weights are read once, compressed (2.1 MB per expert instea
   prompt. A measurement that fails (no VRAM) falls back to 64 rows.
 * `STRATA_I8_FORCE=0` hides the device from the kernels, `STRATA_I8_MT=2|4|8` forces the tile height, `STRATA_I8_WG=0` the
   non-work-group kernels, `STRATA_I8_GRF=128` the 128-register mode, `STRATA_I8_DBG` bring-up switches (skip decode / loads / stores).
-* Default off until the owner decides. The activations are quantized, so the bits differ from the FP16 path (see "Quality").
+* Default (unset): `auto` on a Battlemage GPU (bmg-g21 or bmg-g31 by the SYCL architecture id, not by name; the kernels must be in the binary and pass `available()`), the FP16 path on every other GPU. The engine prints one line at load ("int8 prompt experts are on ... STRATA_PREFILL_MMQ=0 turns them off"). If the kernels are not in the binary for the device (an AOT build for other devices) or a launch throws, the group and the rest of the process take the FP16 chain with one log line. The activations are quantized, so the bits differ from the FP16 path (see "Quality").
 
 ## Tiles and thresholds
 

sycl/CMakeLists.txt

@@ -394,6 +394,9 @@ if(STRATA_SYCL_PARITY)
   add_executable(esimd_moe_parity ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/esimd_moe_parity.cpp)   # 151-b70: fused int8 expert products vs dequant+oneMKL and FP64
   target_link_libraries(esimd_moe_parity PRIVATE strata_kernels strata_prefill)
   add_test(NAME esimd_moe_parity COMMAND esimd_moe_parity 60 8 2)   # 151-b70: skips (exit 0) on a device without the s8 8x16x32 DPAS
+  add_executable(esimd_moe_mode_test ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/esimd_moe_mode_test.cpp)   # 151-b70: the default choice of the expert chain (no GPU)
+  target_include_directories(esimd_moe_mode_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
+  add_test(NAME esimd_moe_mode_test COMMAND esimd_moe_mode_test)
   if(NOT STRATA_SYCL_AOT OR STRATA_SYCL_AOT MATCHES "^(bmg|lnl)")   # DPAS s8 8x16x32 (Xe2): not for an Alchemist / Xe-HPC AOT build
     add_executable(esimd_dpas_probe ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/esimd_dpas_probe.cpp)   # ESIMD DPAS peak (registers only)
     add_executable(esimd_i8_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/esimd_i8_bench.cpp)   # ESIMD int8 DPAS GEMM on pre-decoded weights

sycl/include/strata/prefill/esimd_moe.hpp

@@ -29,6 +29,26 @@ void make_layout(Layout& lo, const int32_t* cnt, int n);
 
 /// True when this device can run the kernels (a Battlemage-class device: sub-group 16 and the s8 DPAS 8x16x32).
 bool available(void* stream);
+/// True when this device is a Battlemage GPU by its architecture id (bmg-g21: B580 / B570, bmg-g31: B70), not by its name.  The default
+/// choice below uses it: Lunar Lake and the other Xe2 parts can run the kernels (`available`) but are not defaulted.  False in a build
+/// without the ESIMD kernels (STRATA_I8_NO_ESIMD).
+bool battlemage(void* stream);
+/// True when any GPU on this machine is a Battlemage one (sizes the staging ring before a queue exists).
+bool any_battlemage();
+
+/// Which chain the prompt experts take.  `env` is STRATA_PREFILL_MMQ (else STRATA_PF_I8), null or empty when unset.  Returns 0 = the
+/// FP16 chain (dequantize + oneMKL), 1 = the int8 chain everywhere, 2 = auto (int8 per group while the group's rows per expert are under
+/// the crossover measured on this GPU).  Set: "0" -> 0, "auto" -> 2, any other number -> 1.  Unset: auto on a Battlemage device whose
+/// kernels are in the binary and run, else 0 (Alchemist and every other GPU keep the FP16 chain).
+inline int choose_mode(const char* env, bool battlemage_device, bool kernels_available) {
+    if (env != nullptr && env[0] != 0) {
+        if (env[0] == 'a' || env[0] == 'A') return 2;
+        long v = 0;
+        for (const char* c = env; *c >= '0' && *c <= '9' && v < 1000000; ++c) v = v * 10 + (*c - '0');
+        return v != 0 ? 1 : 0;
+    }
+    return (battlemage_device && kernels_available) ? 2 : 0;
+}
 
 size_t xq_bytes(const Layout& lo, int K);   // int8 activations, padded tiles
 size_t dx_bytes(const Layout& lo, int K);   // their scales (floats)

sycl/src/kernels/esimd_moe_mode_test.cpp

@@ -0,0 +1,30 @@
+// 151-b70: the default choice of the prompt-expert chain (strata::prefill::i8::choose_mode), no GPU needed.
+//   unset            : auto on a Battlemage device whose kernels are available, else the FP16 chain
+//   STRATA_PREFILL_MMQ=0 / 1 / auto : off / forced / auto, whatever the device
+#include <cstdio>
+#include "strata/prefill/esimd_moe.hpp"
+
+using strata::prefill::i8::choose_mode;
+static int bad = 0;
+#define CHECK(cond, what) do { if (!(cond)) { std::printf("FAIL: %s\n", what); ++bad; } } while (0)
+
+int main() {
+    // unset or empty
+    CHECK(choose_mode(nullptr, true, true) == 2, "unset on Battlemage -> auto");
+    CHECK(choose_mode("", true, true) == 2, "empty on Battlemage -> auto");
+    CHECK(choose_mode(nullptr, true, false) == 0, "unset on Battlemage without kernels (AOT stub, failed probe) -> FP16");
+    CHECK(choose_mode(nullptr, false, true) == 0, "unset on Alchemist / any non-Battlemage device that can run the kernels -> FP16");
+    CHECK(choose_mode(nullptr, false, false) == 0, "unset elsewhere -> FP16");
+    // explicit
+    CHECK(choose_mode("0", true, true) == 0, "=0 on Battlemage turns it off");
+    CHECK(choose_mode("1", true, true) == 1, "=1 forces int8");
+    CHECK(choose_mode("1", false, true) == 1, "=1 forces it on a non-Battlemage device too (the caller still needs available())");
+    CHECK(choose_mode("auto", false, true) == 2, "=auto");
+    CHECK(choose_mode("Auto", true, true) == 2, "=Auto");
+    CHECK(choose_mode("0", false, false) == 0, "=0 elsewhere");
+    CHECK(choose_mode("00", true, true) == 0, "=00 is 0");
+    CHECK(choose_mode("2", true, true) == 1, "any other number forces int8");
+    CHECK(choose_mode("off", true, true) == 0, "an unparsable word is 0, as atoi had it");
+    if (bad == 0) std::printf("esimd_moe_mode_test: all OK\n");
+    return bad == 0 ? 0 : 1;
+}

sycl/src/prefill/esimd_moe.dp.cpp

@@ -48,6 +48,8 @@ void make_layout(Layout& lo, const int32_t* cnt, int n) {
     lo.tile0[n] = t;
 }
 bool available(void*) { return false; }
+bool battlemage(void*) { return false; }
+bool any_battlemage() { return false; }
 size_t xq_bytes(const Layout& lo, int K) { return (size_t) lo.tile0[lo.n] * 8 * K; }
 size_t dx_bytes(const Layout& lo, int K) { return (size_t) lo.tile0[lo.n] * (K / 32) * 8 * sizeof(float); }
 bool workspace(const Layout&, int, int, Workspace&, void*) { return false; }
@@ -805,6 +807,31 @@ void make_layout(Layout& lo, const int32_t* cnt, int n) {
     lo.tile0[n] = t;
 }
 
+namespace {
+bool bmg_arch(const sycl::device& d) {
+    namespace ex = sycl::ext::oneapi::experimental;
+    if (!d.is_gpu()) return false;
+    const auto a = d.get_info<ex::info::device::architecture>();
+    return a == ex::architecture::intel_gpu_bmg_g21 || a == ex::architecture::intel_gpu_bmg_g31;
+}
+}  // namespace
+
+bool battlemage(void* stream) {
+    sycl::queue& q = *strata::q_of(stream);
+    return bmg_arch(q.get_device());
+}
+
+bool any_battlemage() {
+    static const bool v = [] {
+        try {
+            for (const auto& d : sycl::device::get_devices(sycl::info::device_type::gpu))
+                if (bmg_arch(d)) return true;
+        } catch (...) {}
+        return false;
+    }();
+    return v;
+}
+
 bool available(void* stream) {
     sycl::queue& q = *strata::q_of(stream);
     static std::mutex mu;

sycl/src/prefill/prefill.cpp

@@ -238,21 +238,31 @@ inline int ring_slots(size_t T) {
 }
 // A770 port: STRATA_PF_XMX=2 holds up to a group of staged experts until its launch, so routed-only staging needs
 // that many more slots for the lookahead to keep copying (the slots are only reused once the group's event is in)
-// 151-b70: the fused int8 DPAS expert products (prefill/esimd_moe.dp.cpp, Battlemage).  STRATA_PREFILL_MMQ=1 (or STRATA_PF_I8=1) turns
-// them on for the layers whose experts are IQ3_S gate/up + IQ4_NL down, STRATA_PREFILL_MMQ=0 keeps the FP16 path; off by default
-// until the owner has the quality check (the activations are quantized, so the bits differ from the FP16 path).
+// 151-b70: the fused int8 DPAS expert products (prefill/esimd_moe.dp.cpp, Battlemage).  Unset STRATA_PREFILL_MMQ (alias STRATA_PF_I8) is
+// auto on a Battlemage GPU (bmg-g21 / bmg-g31 by architecture id) whose kernels are in the binary, the FP16 chain on every other GPU
+// (Alchemist included); =0 forces the FP16 chain, =1 the int8 chain everywhere it applies (IQ3_S / IQ3_XXS / IQ2_S gate/up with IQ4_NL /
+// Q2_0 down, the 640 x 2560 expert), =auto the per-group choice from the crossover measured on this GPU.  The activations are quantized,
+// so the bits differ from the FP16 path (docs/INTEL_INT8_EXPERTS.md: KL 0.0002 nats, top-1 99.7%).
 static_assert(strata::kernels::kXmxGroupMax <= strata::prefill::i8::kMaxGroup, "the int8 path holds the grouped XMX path's experts: its Layout must be at least as big");
-inline int i8_mode() {
-    static const int v = [] {
-        const char* e = std::getenv("STRATA_PREFILL_MMQ");
-        if (e == nullptr || e[0] == 0) e = std::getenv("STRATA_PF_I8");
-        if (e == nullptr || e[0] == 0) return 0;
-        if (e[0] == 'a' || e[0] == 'A') return 2;
-        return std::atoi(e) != 0 ? 1 : 0;
-    }();
-    return v;
+inline const char* i8_env() {
+    const char* e = std::getenv("STRATA_PREFILL_MMQ");
+    if (e == nullptr || e[0] == 0) e = std::getenv("STRATA_PF_I8");
+    return (e != nullptr && e[0] != 0) ? e : nullptr;
+}
+std::atomic<bool> g_i8_failed{false};    // an int8 launch threw: the FP16 chain for the rest of the process
+// The chain this queue's prompt experts take: 0 FP16, 1 int8, 2 auto.
+inline int i8_mode(void* cs) {
+    if (g_i8_failed.load(std::memory_order_relaxed)) return 0;
+    const char* e = i8_env();
+    if (e != nullptr) return strata::prefill::i8::choose_mode(e, false, false);
+    return strata::prefill::i8::choose_mode(nullptr, strata::prefill::i8::battlemage(cs), strata::prefill::i8::available(cs));
+}
+// ring sizing happens before any queue is at hand: room for the int8 groups whenever the choice could be on
+inline int stage_slots() {
+    const char* e = i8_env();
+    const bool i8 = e != nullptr ? strata::prefill::i8::choose_mode(e, false, false) != 0 : strata::prefill::i8::any_battlemage();
+    return (strata::prefill::xmx::mode() == 2 || i8) ? STAGE + strata::prefill::xmx::kMaxGroup : STAGE;
 }
-inline int stage_slots() { return (strata::prefill::xmx::mode() == 2 || i8_mode() != 0) ? STAGE + strata::prefill::xmx::kMaxGroup : STAGE; }
 constexpr int DQ = 2;              // dequantized-expert ring (FP16 gate/up + down)
 // A770 port: STRATA_PF_XMX=1 dequantizes a whole group of experts, then one grouped XMX launch per product
 inline int dq_slots() { return strata::prefill::xmx::mode() == 1 ? strata::prefill::xmx::kMaxGroup : DQ; }
@@ -1151,6 +1161,11 @@ bool Prefill::init(const core::WeightTable &wt, const core::ModelGeometry &g,
         err = "prefill: geometry differs from the artifact's"; return false;
     }
     m.device = dpct::get_current_device_id();
+    if (static bool said_i8 = false; !said_i8 && i8_mode(m.cs) != 0) {
+        said_i8 = true;
+        std::fprintf(stderr, "strata: int8 prompt experts are on (%s); STRATA_PREFILL_MMQ=0 turns them off\n",
+                     i8_env() == nullptr ? "Battlemage default, auto" : "STRATA_PREFILL_MMQ set");
+    }
     if (stage_le_ < 0) stage_le_ = g.n_layers;
     if (stage_lb_ < 0 || stage_lb_ >= stage_le_ || stage_le_ > g.n_layers || (stage_le_ < g.n_layers) != (next_ != nullptr)) {
         err = "prefill: the stage's layer range is wrong";
@@ -4155,13 +4170,14 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
                         const bool xmx2_only = strata::prefill::xmx::mode() == 2 && !use_mmq && xmx2_fmt();
                         // 151-b70: the fused int8 DPAS products (IQ3_S gate/up, IQ4_NL down, the 640 x 2560 expert), the same
                         // held group of staged experts as above
-                        const bool i8on = !xmx2_only && i8_mode() != 0 && !use_mmq && lay.native &&
+                        const int i8m = i8_mode(m.cs);
+                        const bool i8on = !xmx2_only && i8m != 0 && !use_mmq && lay.native &&
                                           (lay.fmt[(size_t) l].gu_type == 21 || lay.fmt[(size_t) l].gu_type == 18 || lay.fmt[(size_t) l].gu_type == 22) &&
                                           (lay.fmt[(size_t) l].d_type == 20 || lay.fmt[(size_t) l].d_type == 42) &&
                                           lay.fmt[(size_t) l].n_ff == 640 && lay.fmt[(size_t) l].n_embd == 2560 &&
                                           strata::prefill::i8::available(m.cs);
                         const bool xmx2 = xmx2_only || i8on;
-                        if (i8_mode() != 0 && lay.native) {
+                        if (i8_env() != nullptr && i8m != 0 && lay.native) {
                             static std::set<std::tuple<int, int, int, int, int>> said_fmt;
                             const auto& ff = lay.fmt[(size_t) l];
                             if (said_fmt.insert({ff.gu_type, ff.d_type, (int) ff.n_ff, (int) ff.n_embd, (int) i8on}).second)
@@ -4193,8 +4209,10 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
                             if (i8on) {
                                 // STRATA_PREFILL_MMQ=1: always; =auto: while the group's mean rows per expert are under the crossover
                                 // measured on this GPU (i8_calibrate); a group above it takes the FP16 chain below
+                                bool done8 = false;
+                                try {
                                 bool use_i8 = true;
-                                if (i8_mode() == 2) {
+                                if (i8m == 2) {
                                     static std::map<std::tuple<const void*, int, int>, I8Cross> xover;   // per queue (a layer split has a card per queue)
                                     auto it = xover.find({(const void*) m.cs, f.gu_type, f.d_type});
                                     if (it == xover.end()) it = xover.emplace(std::make_tuple((const void*) m.cs, f.gu_type, f.d_type), i8_calibrate(m, f)).first;
@@ -4206,12 +4224,11 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
                                         std::fprintf(stderr, "strata: int8 experts (auto): %ld groups took the int8 chain, %ld the FP16 chain\n", groups_i8, groups_fp);
                                     }
                                 }
-                                bool done8 = false;
                                 if (use_i8) {
                                     if (static bool said = false; !said) {
                                         said = true;
                                         std::fprintf(stderr, "strata: prompt experts on the fused int8 DPAS products (STRATA_PREFILL_MMQ=%s)\n",
-                                                     i8_mode() == 2 ? "auto" : "1");
+                                                     i8m == 2 ? "auto" : "1");
                                     }
                                     namespace i8 = strata::prefill::i8;
                                     i8::Layout lo8;
@@ -4233,6 +4250,11 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
                                         }
                                     }
                                 }
+                                } catch (const std::exception& ex) {   // the kernels are not in this binary for this device, or a launch failed: the FP16 chain, for good
+                                    done8 = false;
+                                    if (!g_i8_failed.exchange(true))
+                                        std::fprintf(stderr, "strata: the int8 expert kernels failed to launch (%s): the FP16 expert chain is used from here\n", ex.what());
+                                }
                                 if (!done8) {   // the FP16 chain: dequantize, two oneMKL GEMMs, SwiGLU, per expert
                                     int64_t o = xg_r0;
                                     for (int i = 0; i < xg_n; ++i) {