diff --git a/Cargo.toml b/Cargo.toml
index ac6299e4b..00322e481 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -329,3 +329,8 @@ required-features = ["xla-diagnostics"]
 [[example]]
 name = "xla_aux_smoke"
 required-features = ["xla-iree"]
+
+# Pinned #874 MLX-vs-IREE Phi4MM Conformer and speech/vision projection parity.
+[[example]]
+name = "xla_phi4_audio_check"
+required-features = ["xla-iree"]
diff --git a/examples/xla_phi4_audio_check.rs b/examples/xla_phi4_audio_check.rs
new file mode 100644
index 000000000..5a8c74105
--- /dev/null
+++ b/examples/xla_phi4_audio_check.rs
@@ -0,0 +1,281 @@
+// Copyright 2025-2026 Lablup Inc. and Jeongkyu Shin
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+//! Real-checkpoint #874 projection parity for the IREE Phi4MM audio module.
+//!
+//! This diagnostic intentionally loads the MLX model only to capture the
+//! qualified reference, drops it, and then loads the independent IREE-only
+//! audio runtime. Production XLA execution never constructs the MLX decoder.
+//! Full intermediate comparisons remain diagnostic; the pinned #874 projection
+//! prefix is the qualified acceptance boundary.
+
+use std::path::PathBuf;
+
+use mlxcel::LoadedModel;
+use mlxcel_xla::{
+    PHI4MM_AUDIO_CHECKPOINT_REVISION, Phi4AudioDiagnosticRuntime, phi4_audio_bucket_for_frames,
+};
+
+fn mlx_f32(array: &mlxcel_core::MlxArray) -> Result<Vec<f32>, String> {
+    let array = mlxcel_core::astype(array, mlxcel_core::dtype::FLOAT32);
+    let bytes = mlxcel_core::try_array_to_raw_bytes(&array)
+        .map_err(|error| format!("export MLX reference: {error}"))?;
+    Ok(bytes
+        .chunks_exact(4)
+        .map(|chunk| f32::from_ne_bytes(chunk.try_into().expect("four-byte f32")))
+        .collect())
+}
+
+struct Comparison {
+    max: f32,
+    max_index: usize,
+    actual_at_max: f32,
+    expected_at_max: f32,
+    rms: f32,
+}
+
+struct ReferenceCheckpoint {
+    name: String,
+    shape: Vec<usize>,
+    values: Vec<f32>,
+}
+
+fn mlx_checkpoint(
+    name: impl Into<String>,
+    array: &mlxcel_core::MlxArray,
+) -> Result<ReferenceCheckpoint, String> {
+    let name = name.into();
+    let shape = mlxcel_core::array_shape(array)
+        .into_iter()
+        .map(|dimension| {
+            usize::try_from(dimension)
+                .map_err(|_| format!("negative MLX diagnostic dimension {dimension}"))
+        })
+        .collect::<Result<Vec<_>, String>>()?;
+    Ok(ReferenceCheckpoint {
+        name,
+        shape,
+        values: mlx_f32(array)?,
+    })
+}
+
+fn compare(actual: &[f32], expected: &[f32], label: &str) -> Result<Comparison, String> {
+    if actual.len() != expected.len() {
+        return Err(format!(
+            "{label} length {}, expected {}",
+            actual.len(),
+            expected.len()
+        ));
+    }
+    let mut max = 0.0f32;
+    let mut max_index = 0usize;
+    let mut squared = 0.0f64;
+    for (index, (&actual, &expected)) in actual.iter().zip(expected).enumerate() {
+        let difference = (actual - expected).abs();
+        if difference > max {
+            max = difference;
+            max_index = index;
+        }
+        squared += f64::from(difference) * f64::from(difference);
+    }
+    let rms = (squared / actual.len() as f64).sqrt() as f32;
+    Ok(Comparison {
+        max,
+        max_index,
+        actual_at_max: actual[max_index],
+        expected_at_max: expected[max_index],
+        rms,
+    })
+}
+
+fn print_comparison(label: &str, comparison: &Comparison, width: usize) {
+    println!(
+        "{label} max_abs={:.8} at row={} dim={} (actual={:.8}, expected={:.8}) rms={:.8}",
+        comparison.max,
+        comparison.max_index / width,
+        comparison.max_index % width,
+        comparison.actual_at_max,
+        comparison.expected_at_max,
+        comparison.rms,
+    );
+}
+
+fn main() -> Result<(), String> {
+    let model_dir = PathBuf::from(
+        std::env::var_os("MLXCEL_PHI4MM_MODEL")
+            .ok_or("set MLXCEL_PHI4MM_MODEL to the pinned official checkpoint")?,
+    );
+    let revision =
+        std::env::var("MLXCEL_PHI4MM_REVISION").map_err(|_| "set MLXCEL_PHI4MM_REVISION")?;
+    if revision != PHI4MM_AUDIO_CHECKPOINT_REVISION {
+        return Err(format!(
+            "checkpoint revision {revision} does not match pinned {PHI4MM_AUDIO_CHECKPOINT_REVISION}"
+        ));
+    }
+    let device = std::env::args()
+        .nth(1)
+        .unwrap_or_else(|| "local-task".to_string());
+    let fixture: serde_json::Value =
+        serde_json::from_str(include_str!("../tests/fixtures/phi4mm_audio_parity.json"))
+            .map_err(|error| format!("parse #874 fixture: {error}"))?;
+    let audio = model_dir.join(
+        fixture["audio"]
+            .as_str()
+            .ok_or("#874 fixture audio path is not a string")?,
+    );
+
+    let (loaded, _) = mlxcel::load_model(&model_dir)
+        .map_err(|error| format!("load #874 MLX reference: {error}"))?;
+    let LoadedModel::Phi4MMVLM(model) = loaded else {
+        return Err("pinned checkpoint did not load as Phi4MMVLM".to_string());
+    };
+    let (samples, sample_rate) =
+        mlxcel::audio::load_wav_file(&audio).map_err(|error| error.to_string())?;
+    let batch = model.extract_audio(&[(samples, sample_rate)])?;
+    let shape = mlxcel_core::array_shape(&batch.clips[0]);
+    if shape.len() != 3 || shape[0] != 1 || shape[2] != 80 {
+        return Err(format!("unexpected #874 feature shape {shape:?}"));
+    }
+    let frame_len =
+        usize::try_from(shape[1]).map_err(|_| format!("negative frame length {}", shape[1]))?;
+    let features = mlx_f32(&batch.clips[0])?;
+    let encoder = model.audio_encoder.forward_diagnostics(&batch.clips[0])?;
+    let speech = model
+        .audio_projection
+        .forward_diagnostics(&encoder.encoded, false);
+    let vision = model
+        .audio_projection
+        .forward_diagnostics(&encoder.encoded, true);
+    let mut references = vec![
+        mlx_checkpoint("subsample.conv0", &encoder.subsample.conv0)?,
+        mlx_checkpoint(
+            "subsample.conv1.depthwise",
+            &encoder.subsample.conv1_depthwise,
+        )?,
+        mlx_checkpoint(
+            "subsample.conv1.pointwise",
+            &encoder.subsample.conv1_pointwise,
+        )?,
+        mlx_checkpoint("subsample.conv1", &encoder.subsample.conv1)?,
+        mlx_checkpoint(
+            "subsample.conv2.depthwise",
+            &encoder.subsample.conv2_depthwise,
+        )?,
+        mlx_checkpoint(
+            "subsample.conv2.pointwise",
+            &encoder.subsample.conv2_pointwise,
+        )?,
+        mlx_checkpoint("subsample.conv2", &encoder.subsample.conv2)?,
+        mlx_checkpoint("subsample.projected", &encoder.subsample.projected)?,
+        mlx_checkpoint("block0.after_ff_in", &encoder.block0.after_ff_in)?,
+        mlx_checkpoint("block0.attention", &encoder.block0.attention)?,
+        mlx_checkpoint("block0.after_attention", &encoder.block0.after_attention)?,
+        mlx_checkpoint("block0.convolution", &encoder.block0.convolution)?,
+        mlx_checkpoint(
+            "block0.after_convolution",
+            &encoder.block0.after_convolution,
+        )?,
+        mlx_checkpoint("block0.ff_out", &encoder.block0.ff_out)?,
+        mlx_checkpoint("block0.output", &encoder.block0.output)?,
+    ];
+    for (index, values) in &encoder.selected_blocks {
+        references.push(mlx_checkpoint(format!("block{index}.output"), values)?);
+    }
+    references.extend([
+        mlx_checkpoint("encoder.output", &encoder.encoded)?,
+        mlx_checkpoint("projection.speech.first", &speech.first)?,
+        mlx_checkpoint("projection.speech.output", &speech.output)?,
+        mlx_checkpoint("projection.vision.first", &vision.first)?,
+        mlx_checkpoint("projection.vision.output", &vision.output)?,
+    ]);
+    drop(model);
+
+    let frame_bucket = phi4_audio_bucket_for_frames(frame_len)
+        .ok_or_else(|| format!("{frame_len} feature frames exceed the Phi4MM XLA policy"))?;
+    let mut runtime = Phi4AudioDiagnosticRuntime::load(&model_dir, &device, frame_bucket)?;
+    let diagnostics = runtime.capture(&features, frame_len)?;
+    let expected_rows = fixture["audio_embed_size"]
+        .as_u64()
+        .ok_or("#874 fixture audio_embed_size is not an integer")? as usize;
+    if diagnostics.valid_rows != expected_rows {
+        return Err(format!(
+            "projected length {}, expected {expected_rows}",
+            diagnostics.valid_rows
+        ));
+    }
+    if diagnostics.checkpoints.len() != references.len() {
+        return Err(format!(
+            "IREE returned {} diagnostic checkpoints, MLX captured {}",
+            diagnostics.checkpoints.len(),
+            references.len()
+        ));
+    }
+    for (actual, expected) in diagnostics.checkpoints.iter().zip(&references) {
+        if actual.name != expected.name {
+            return Err(format!(
+                "diagnostic checkpoint order drift: IREE {}, MLX {}",
+                actual.name, expected.name
+            ));
+        }
+        if actual.shape.len() != expected.shape.len()
+            || actual.shape[0] != expected.shape[0]
+            || actual.shape[2..] != expected.shape[2..]
+            || actual.shape[1] < expected.shape[1]
+        {
+            return Err(format!(
+                "{} shape mismatch: IREE {:?}, MLX {:?}",
+                actual.name, actual.shape, expected.shape
+            ));
+        }
+        let comparison = compare(
+            &actual.values[..expected.values.len()],
+            &expected.values,
+            actual.name,
+        )?;
+        print_comparison(
+            actual.name,
+            &comparison,
+            *expected.shape.last().unwrap_or(&1),
+        );
+    }
+    let oracle_prefix = fixture["projection_first"]
+        .as_array()
+        .ok_or("#874 projection_first is not an array")?
+        .iter()
+        .map(|value| value.as_f64().map(|value| value as f32))
+        .collect::<Option<Vec<_>>>()
+        .ok_or("#874 projection_first contains a non-number")?;
+    let speech_output = diagnostics
+        .checkpoints
+        .iter()
+        .find(|checkpoint| checkpoint.name == "projection.speech.output")
+        .ok_or("IREE diagnostic omitted projection.speech.output")?;
+    let prefix_comparison = compare(
+        &speech_output.values[..oracle_prefix.len()],
+        &oracle_prefix,
+        "#874 speech prefix",
+    )?;
+    println!(
+        "device={device} bucket={frame_bucket} frames={frame_len} projected_rows={}",
+        diagnostics.valid_rows,
+    );
+    print_comparison("#874 speech prefix", &prefix_comparison, 3072);
+    if prefix_comparison.max > 0.025 {
+        return Err(format!(
+            "#874 speech projection prefix exceeded tolerance: max_abs={} at flat index {}",
+            prefix_comparison.max, prefix_comparison.max_index
+        ));
+    }
+    Ok(())
+}
diff --git a/src/audio/phi4mm/encoder.rs b/src/audio/phi4mm/encoder.rs
index 0e71854e6..f67a552d6 100644
--- a/src/audio/phi4mm/encoder.rs
+++ b/src/audio/phi4mm/encoder.rs
@@ -170,6 +170,18 @@ struct Subsampler {
     out: UnifiedLinear,
 }
 
+#[doc(hidden)]
+pub struct Phi4MMAudioSubsampleDiagnostics {
+    pub conv0: UniquePtr<MlxArray>,
+    pub conv1_depthwise: UniquePtr<MlxArray>,
+    pub conv1_pointwise: UniquePtr<MlxArray>,
+    pub conv1: UniquePtr<MlxArray>,
+    pub conv2_depthwise: UniquePtr<MlxArray>,
+    pub conv2_pointwise: UniquePtr<MlxArray>,
+    pub conv2: UniquePtr<MlxArray>,
+    pub projected: UniquePtr<MlxArray>,
+}
+
 impl Subsampler {
     fn load(weights: &WeightMap, prefix: &str, channels: i32) -> Result<Self, String> {
         Ok(Self {
@@ -228,23 +240,36 @@ impl Subsampler {
     }
 
     fn forward(&self, x: &MlxArray) -> Result<UniquePtr<MlxArray>, String> {
+        Ok(self.forward_diagnostics(x)?.projected)
+    }
+
+    fn forward_diagnostics(&self, x: &MlxArray) -> Result<Phi4MMAudioSubsampleDiagnostics, String> {
         let shape = mlxcel_core::array_shape(x);
         let x = mlxcel_core::reshape(x, &[shape[0], shape[1], shape[2], 1]);
         let x = self.conv0.forward(&x)?;
-        let x = relu(&x);
-        let x = self.conv1_depthwise.forward(&x)?;
-        let x = self.conv1_pointwise.forward(&x)?;
-        let x = relu(&x);
-        let x = self.conv2_depthwise.forward(&x)?;
-        let x = self.conv2_pointwise.forward(&x)?;
-        let x = relu(&x);
+        let conv0 = relu(&x);
+        let conv1_depthwise = self.conv1_depthwise.forward(&conv0)?;
+        let conv1_pointwise = self.conv1_pointwise.forward(&conv1_depthwise)?;
+        let conv1 = relu(&conv1_pointwise);
+        let conv2_depthwise = self.conv2_depthwise.forward(&conv1)?;
+        let conv2_pointwise = self.conv2_pointwise.forward(&conv2_depthwise)?;
+        let conv2 = relu(&conv2_pointwise);
         // PyTorch keeps convolution output as [B,C,T,F] and flattens after
         // transpose(1,2), so channel must precede frequency for embed.out.
-        let x = mlxcel_core::transpose_axes(&x, &[0, 1, 3, 2]);
+        let x = mlxcel_core::transpose_axes(&conv2, &[0, 1, 3, 2]);
         let shape = mlxcel_core::array_shape(&x);
         let x = mlxcel_core::reshape(&x, &[shape[0], shape[1], shape[2] * shape[3]]);
-        let x = self.out.forward(&x);
-        Ok(x)
+        let projected = self.out.forward(&x);
+        Ok(Phi4MMAudioSubsampleDiagnostics {
+            conv0,
+            conv1_depthwise,
+            conv1_pointwise,
+            conv1,
+            conv2_depthwise,
+            conv2_pointwise,
+            conv2,
+            projected,
+        })
     }
 }
 
@@ -415,6 +440,17 @@ struct ConformerBlock {
     output_norm: LayerNorm,
 }
 
+#[doc(hidden)]
+pub struct Phi4MMAudioBlockDiagnostics {
+    pub after_ff_in: UniquePtr<MlxArray>,
+    pub attention: UniquePtr<MlxArray>,
+    pub after_attention: UniquePtr<MlxArray>,
+    pub convolution: UniquePtr<MlxArray>,
+    pub after_convolution: UniquePtr<MlxArray>,
+    pub ff_out: UniquePtr<MlxArray>,
+    pub output: UniquePtr<MlxArray>,
+}
+
 impl ConformerBlock {
     fn load(
         weights: &WeightMap,
@@ -443,22 +479,46 @@ impl ConformerBlock {
         x: &MlxArray,
         relative_bias: &MlxArray,
     ) -> Result<UniquePtr<MlxArray>, String> {
-        let h = mlxcel_core::add(
+        Ok(self.forward_diagnostics(x, relative_bias)?.output)
+    }
+
+    fn forward_diagnostics(
+        &self,
+        x: &MlxArray,
+        relative_bias: &MlxArray,
+    ) -> Result<Phi4MMAudioBlockDiagnostics, String> {
+        let after_ff_in = mlxcel_core::add(
             x,
             &mlxcel_core::multiply_scalar(&self.ff_in.forward(x), 0.5),
         );
-        let attention_input = self.attention_norm.forward(&h);
-        let attn = self.attention.forward(&attention_input, relative_bias);
-        let h = mlxcel_core::add(&h, &attn);
-        let conv = self.conv.forward(&h)?;
-        let h = mlxcel_core::add(&h, &conv);
-        let ff = mlxcel_core::multiply_scalar(&self.ff_out.forward(&h), 0.5);
-        let h = mlxcel_core::add(&h, &ff);
-        let h = self.output_norm.forward(&h);
-        Ok(h)
+        let attention_input = self.attention_norm.forward(&after_ff_in);
+        let attention = self.attention.forward(&attention_input, relative_bias);
+        let after_attention = mlxcel_core::add(&after_ff_in, &attention);
+        let convolution = self.conv.forward(&after_attention)?;
+        let after_convolution = mlxcel_core::add(&after_attention, &convolution);
+        let ff_out = mlxcel_core::multiply_scalar(&self.ff_out.forward(&after_convolution), 0.5);
+        let pre_norm = mlxcel_core::add(&after_convolution, &ff_out);
+        let output = self.output_norm.forward(&pre_norm);
+        Ok(Phi4MMAudioBlockDiagnostics {
+            after_ff_in,
+            attention,
+            after_attention,
+            convolution,
+            after_convolution,
+            ff_out,
+            output,
+        })
     }
 }
 
+#[doc(hidden)]
+pub struct Phi4MMAudioEncoderDiagnostics {
+    pub subsample: Phi4MMAudioSubsampleDiagnostics,
+    pub block0: Phi4MMAudioBlockDiagnostics,
+    pub selected_blocks: Vec<(usize, UniquePtr<MlxArray>)>,
+    pub encoded: UniquePtr<MlxArray>,
+}
+
 pub struct Phi4MMAudioEncoder {
     mean: UniquePtr<MlxArray>,
     invstd: UniquePtr<MlxArray>,
@@ -539,6 +599,44 @@ impl Phi4MMAudioEncoder {
         Ok(crate::vision::encoders::qwen2_vl::concat_many(&chunks, 1))
     }
 
+    /// Real-checkpoint parity checkpoints used by the independent IREE audio
+    /// diagnostic. This is deliberately not part of the production inference
+    /// path and only accepts the single <=500-row Conformer chunk exercised by
+    /// the pinned #874 fixture.
+    #[doc(hidden)]
+    pub fn forward_diagnostics(
+        &self,
+        features: &MlxArray,
+    ) -> Result<Phi4MMAudioEncoderDiagnostics, String> {
+        let features = mlxcel_core::astype(features, mlxcel_core::array_dtype(&self.mean));
+        let normalized =
+            mlxcel_core::multiply(&mlxcel_core::subtract(&features, &self.mean), &self.invstd);
+        let subsample = self.subsampler.forward_diagnostics(&normalized)?;
+        let shape = mlxcel_core::array_shape(&subsample.projected);
+        if shape[1] > 500 {
+            return Err(format!(
+                "Phi4MM audio diagnostic requires one <=500-row chunk, got {}",
+                shape[1]
+            ));
+        }
+        let bias = self.relative_bias(shape[1]);
+        let block0 = self.blocks[0].forward_diagnostics(&subsample.projected, &bias)?;
+        let mut h = mlxcel_core::copy(&block0.output);
+        let mut selected_blocks = Vec::new();
+        for (index, block) in self.blocks.iter().enumerate().skip(1) {
+            h = block.forward(&h, &bias)?;
+            if matches!(index, 1 | 5 | 11 | 17 | 23) {
+                selected_blocks.push((index, mlxcel_core::copy(&h)));
+            }
+        }
+        Ok(Phi4MMAudioEncoderDiagnostics {
+            subsample,
+            block0,
+            selected_blocks,
+            encoded: h,
+        })
+    }
+
     fn relative_bias(&self, length: i32) -> UniquePtr<MlxArray> {
         let mut indices = Vec::with_capacity((length * length) as usize);
         for query in 0..length {
@@ -563,6 +661,12 @@ pub struct Phi4MMAudioProjection {
     vision_2: UnifiedLinear,
 }
 
+#[doc(hidden)]
+pub struct Phi4MMAudioProjectionDiagnostics {
+    pub first: UniquePtr<MlxArray>,
+    pub output: UniquePtr<MlxArray>,
+}
+
 impl Phi4MMAudioProjection {
     pub fn from_weights(weights: &WeightMap, prefix: &str) -> Result<Self, String> {
         Ok(Self {
@@ -574,11 +678,22 @@ impl Phi4MMAudioProjection {
     }
 
     pub fn forward(&self, encoded: &MlxArray, vision_mode: bool) -> UniquePtr<MlxArray> {
+        self.forward_diagnostics(encoded, vision_mode).output
+    }
+
+    #[doc(hidden)]
+    pub fn forward_diagnostics(
+        &self,
+        encoded: &MlxArray,
+        vision_mode: bool,
+    ) -> Phi4MMAudioProjectionDiagnostics {
         let (first, second) = if vision_mode {
             (&self.vision_1, &self.vision_2)
         } else {
             (&self.speech_1, &self.speech_2)
         };
-        second.forward(&mlxcel_core::gelu(&first.forward(encoded)))
+        let first = first.forward(encoded);
+        let output = second.forward(&mlxcel_core::gelu(&first));
+        Phi4MMAudioProjectionDiagnostics { first, output }
     }
 }
diff --git a/src/commands/generate.rs b/src/commands/generate.rs
index 225ec01d4..e4fe4dd1e 100644
--- a/src/commands/generate.rs
+++ b/src/commands/generate.rs
@@ -1054,6 +1054,9 @@ fn validate_xla_cli_image_cardinality(declared: usize, decoded: usize) -> Result
 
 #[cfg(feature = "xla-backend")]
 fn decode_xla_cli_images(paths: &[std::path::PathBuf]) -> Result<Vec<image::DynamicImage>> {
+    if paths.is_empty() {
+        return Ok(Vec::new());
+    }
     let limits = mlxcel::current_image_input_limits();
     ensure!(
         paths.len() <= limits.max_images_per_request,
@@ -1098,18 +1101,16 @@ fn decode_xla_cli_images(paths: &[std::path::PathBuf]) -> Result<Vec<image::Dyna
 fn generate_xla(
     model_path: &Path,
     num_layers: usize,
+    prompt: &str,
     prompt_tokens: &[i32],
+    tokenizer: &mlxcel::tokenizer::MlxcelTokenizer,
     image_paths: &[std::path::PathBuf],
-    has_audio: bool,
+    audio_path: Option<&Path>,
     has_video: bool,
     max_tokens: usize,
     kv_cache_mode: KVCacheMode,
     token_bias: TokenBiasMap,
 ) -> Result<(Vec<i32>, GenerationStats)> {
-    ensure!(
-        !has_audio,
-        "the OpenXLA backend does not support audio input yet"
-    );
     ensure!(
         !has_video,
         "the OpenXLA backend does not support video input yet"
@@ -1120,7 +1121,82 @@ fn generate_xla(
     let tokens = match &mut session {
         mlxcel::Session::Xla(s) => {
             let eos = s.eos_token_ids().to_vec();
-            if image_paths.is_empty() {
+            let phi4mm_media = mlxcel::models::get_model_type(model_path)?
+                == mlxcel::models::ModelType::Phi4MMVLM
+                && (audio_path.is_some() || !image_paths.is_empty());
+            if phi4mm_media {
+                #[cfg(not(feature = "xla-iree"))]
+                {
+                    let _ = (audio_path, prompt, tokenizer, image_paths);
+                    anyhow::bail!(
+                        "OpenXLA Phi4MM media input requires a build with the xla-iree feature"
+                    );
+                }
+                #[cfg(feature = "xla-iree")]
+                {
+                    let mut tokenization_error = None;
+                    let normalized = mlxcel::phi4mm_prompt::prepare_phi4mm_prompt_tokens(
+                        prompt,
+                        image_paths.len(),
+                        usize::from(audio_path.is_some()),
+                        |text, add_special| match tokenizer.encode(text, add_special) {
+                            Ok(tokens) => tokens.into_iter().map(|token| token as i32).collect(),
+                            Err(error) => {
+                                tokenization_error = Some(error.to_string());
+                                Vec::new()
+                            }
+                        },
+                    )
+                    .map_err(|error| anyhow!("Phi4MM prompt normalization failed: {error}"))?;
+                    if let Some(error) = tokenization_error {
+                        return Err(anyhow!("Phi4MM tokenization failed: {error}"));
+                    }
+                    let cancelled = std::sync::atomic::AtomicBool::new(false);
+                    let images = decode_xla_cli_images(image_paths)?;
+                    let device = std::env::var("MLXCEL_XLA_DEVICE")
+                        .unwrap_or_else(|_| mlxcel_xla::default_device().to_string());
+                    let mut producer = if images.is_empty() {
+                        mlxcel::multimodal::phi4mm_xla_audio::Phi4MmXlaAudioProducer::load(
+                            model_path,
+                            &device,
+                            s.context_capacity(),
+                        )
+                    } else {
+                        mlxcel::multimodal::phi4mm_xla_audio::Phi4MmXlaAudioProducer::load_multimodal(
+                            model_path,
+                            &device,
+                            s.context_capacity(),
+                        )
+                    }
+                    .map_err(anyhow::Error::msg)?;
+                    let prepared = if let Some(audio_path) = audio_path {
+                        let policy =
+                            mlxcel::multimodal::phi4mm_xla_audio::load_phi4mm_audio_policy(
+                                model_path,
+                            )
+                            .map_err(anyhow::Error::msg)?;
+                        let waveforms =
+                            mlxcel::audio::preprocess_wav_file(audio_path, policy, &cancelled)
+                                .map_err(|error| {
+                                    anyhow!("OpenXLA audio waveform preprocessing failed: {error}")
+                                })?;
+                        producer.prepare_media(waveforms, normalized.tokens, &images, &cancelled)
+                    } else {
+                        producer.prepare_images(normalized.tokens, &images, &cancelled)
+                    }
+                    .map_err(|error| {
+                        anyhow!("OpenXLA Phi4MM media preprocessing failed: {error}")
+                    })?;
+                    s.generate_prepared_greedy(&prepared, max_tokens, &eos)
+                        .map_err(|error| {
+                            anyhow!("OpenXLA Phi4MM media generation failed: {error}")
+                        })?
+                }
+            } else if audio_path.is_some() {
+                anyhow::bail!(
+                    "the loaded OpenXLA model/runtime bundle does not support audio input"
+                );
+            } else if image_paths.is_empty() {
                 s.generate_greedy(prompt_tokens, max_tokens, &eos)
                     .map_err(|e| anyhow!("OpenXLA generation failed: {e}"))?
             } else {
@@ -1902,9 +1978,11 @@ fn run_generate_once(mut args: GenerateArgs) -> Result<()> {
         let (generated_tokens, stats) = generate_xla(
             &args.model.model,
             num_layers,
+            &prompt,
             &prompt_tokens,
+            &tokenizer,
             &args.generation.image,
-            args.generation.audio.is_some(),
+            args.generation.audio.as_deref(),
             !args.generation.video.is_empty(),
             args.generation.max_tokens,
             kv_cache_mode,
diff --git a/src/commands/generate_tests.rs b/src/commands/generate_tests.rs
index b20e4ef08..bfd7b4a6b 100644
--- a/src/commands/generate_tests.rs
+++ b/src/commands/generate_tests.rs
@@ -12,6 +12,8 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
+#[cfg(feature = "xla-backend")]
+use super::decode_xla_cli_images;
 use super::{
     apply_user_chat_template, apply_vlm_chat_template, cli_pipeline_requested,
     estimate_delta_label_and_bytes, generated_suffix, generation_stats_from_duration,
@@ -56,6 +58,12 @@ fn xla_cli_rejects_partial_multi_image_decode() {
     assert!(validate_xla_cli_image_cardinality(1, images.len()).is_ok());
 }
 
+#[cfg(feature = "xla-backend")]
+#[test]
+fn xla_cli_accepts_empty_image_list_for_audio_only_requests() {
+    assert!(decode_xla_cli_images(&[]).unwrap().is_empty());
+}
+
 // issue #166: the offline MTP loop is constructed only when the operator
 // explicitly passes `--draft-kind mtp`. An auto-detected MTP shape (no explicit
 // flag) keeps the classic SpeculativeGenerator path, and DFlash / InternalMtp
diff --git a/src/lib.rs b/src/lib.rs
index 2c9ec15ad..d9520a12f 100644
--- a/src/lib.rs
+++ b/src/lib.rs
@@ -98,9 +98,9 @@ pub use multimodal::{
 // fold to the single MLX variant with no runtime dispatch.
 pub use backend::{Backend, ComputeBackend, MlxBackend, Session, select_backend};
 pub use mlxcel_core::session::{
-    InferenceSession, MlxInferenceSession, OwnedTensor, PreparedAttentionBias, PreparedModality,
-    PreparedPositions, PreparedPrefill, PreparedPrefillError, PreparedTensorDType,
-    SessionCapabilities,
+    InferenceSession, MlxInferenceSession, OwnedTensor, PreparedAdapterMode, PreparedAttentionBias,
+    PreparedModality, PreparedPositions, PreparedPrefill, PreparedPrefillError,
+    PreparedTensorDType, SessionCapabilities,
 };
 pub use server::ImageInputLimits;
 
diff --git a/src/lib/mlxcel-core/src/lib.rs b/src/lib/mlxcel-core/src/lib.rs
index 8e83baa7b..49e6c98f0 100644
--- a/src/lib/mlxcel-core/src/lib.rs
+++ b/src/lib/mlxcel-core/src/lib.rs
@@ -2596,8 +2596,8 @@ pub use sampling::TokenBiasMap;
 // Re-export the inference-session contract (issue #448, ADR 0004) so the root
 // crate's compute-backend seam can construct and dispatch MLX sessions.
 pub use multimodal::{
-    OwnedTensor, PreparedAttentionBias, PreparedModality, PreparedPositions, PreparedPrefill,
-    PreparedPrefillError, PreparedTensorDType,
+    OwnedTensor, PreparedAdapterMode, PreparedAttentionBias, PreparedModality, PreparedPositions,
+    PreparedPrefill, PreparedPrefillError, PreparedTensorDType,
 };
 pub use session::{InferenceSession, MlxInferenceSession, SessionCapabilities};
 // Re-export N-gram loop-detection so the server control plane and CLI decode
diff --git a/src/lib/mlxcel-core/src/multimodal.rs b/src/lib/mlxcel-core/src/multimodal.rs
index 103477b75..65049625c 100644
--- a/src/lib/mlxcel-core/src/multimodal.rs
+++ b/src/lib/mlxcel-core/src/multimodal.rs
@@ -147,6 +147,27 @@ pub struct PreparedModality {
     pub token_count: usize,
 }
 
+/// Request-scoped decoder adapter selected by a prepared multimodal prefill.
+///
+/// The value is carried into the engine slot and remains fixed for every
+/// decode step. Plain text and families without request-time adapters use
+/// [`Language`](Self::Language).
+#[derive(Debug, Clone, Copy, Default, PartialEq, Eq, Serialize, Deserialize)]
+#[repr(i32)]
+pub enum PreparedAdapterMode {
+    #[default]
+    Language = 0,
+    Speech = 1,
+    Vision = 2,
+}
+
+impl PreparedAdapterMode {
+    #[must_use]
+    pub const fn code(self) -> i32 {
+        self as i32
+    }
+}
+
 /// Owned prefill payload crossing from host preprocessing to an engine worker.
 #[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
 pub struct PreparedPrefill {
@@ -158,6 +179,9 @@ pub struct PreparedPrefill {
     pub attention_bias: PreparedAttentionBias,
     pub sequence_len: usize,
     pub modalities: Vec<PreparedModality>,
+    /// Immutable request mode used by adapter-aware language graphs.
+    #[serde(default)]
+    pub adapter_mode: PreparedAdapterMode,
 }
 
 impl PreparedPrefill {
@@ -248,8 +272,16 @@ impl PreparedPrefill {
             attention_bias,
             sequence_len,
             modalities,
+            adapter_mode: PreparedAdapterMode::Language,
         })
     }
+
+    /// Attach the immutable language-adapter mode selected by the producer.
+    #[must_use]
+    pub fn with_adapter_mode(mut self, mode: PreparedAdapterMode) -> Self {
+        self.adapter_mode = mode;
+        self
+    }
 }
 
 /// Structural errors at the owned prefill boundary.
diff --git a/src/lib/mlxcel-core/src/session.rs b/src/lib/mlxcel-core/src/session.rs
index d33cca7a1..426c219fd 100644
--- a/src/lib/mlxcel-core/src/session.rs
+++ b/src/lib/mlxcel-core/src/session.rs
@@ -70,8 +70,8 @@ use crate::sampling::TokenBiasMap;
 // Keep the prepared-prefill DTO discoverable from the session boundary while
 // its implementation remains in a focused module below the 500-line file cap.
 pub use crate::multimodal::{
-    OwnedTensor, PreparedAttentionBias, PreparedModality, PreparedPositions, PreparedPrefill,
-    PreparedPrefillError, PreparedTensorDType,
+    OwnedTensor, PreparedAdapterMode, PreparedAttentionBias, PreparedModality, PreparedPositions,
+    PreparedPrefill, PreparedPrefillError, PreparedTensorDType,
 };
 
 /// Capabilities a backend inference session advertises so the control plane can
diff --git a/src/lib/mlxcel-xla/csrc/xla_iree.c b/src/lib/mlxcel-xla/csrc/xla_iree.c
index 1417c7725..32578ed45 100644
--- a/src/lib/mlxcel-xla/csrc/xla_iree.c
+++ b/src/lib/mlxcel-xla/csrc/xla_iree.c
@@ -114,6 +114,7 @@ typedef struct xla_ctx {
   int32_t deepstack_visual_positions;
   int32_t* deepstack_target_layers;
   int32_t has_deepstack;
+  int32_t has_adapter_modes;
   int32_t has_prefill_diagnostics;
   iree_hal_buffer_view_t** weights;  // resident weights, uploaded once
   iree_hal_buffer_view_t* kcache;    // threaded KV (set by prefill, advanced by decode)
@@ -307,6 +308,37 @@ static iree_status_t xla_alloc_bv(xla_ctx* c, iree_host_size_t rank,
       iree_make_const_byte_span(data, nbytes), out);
 }
 
+// Push the request-scoped Phi4MM adapter mode immediately after resident
+// weights. Adapter-aware single-sequence graphs consume a scalar; ragged graphs
+// consume one mode per row. Non-adapter bundles reject non-language modes.
+static iree_status_t xla_push_adapter_modes(
+    xla_ctx* c, iree_runtime_call_t* call, const int32_t* modes,
+    int32_t count, bool vector, iree_hal_buffer_view_t** out) {
+  *out = NULL;
+  if (!modes || count <= 0) {
+    return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
+                            "adapter mode payload is empty");
+  }
+  for (int32_t i = 0; i < count; ++i) {
+    if (modes[i] < 0 || modes[i] > 2) {
+      return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
+                              "adapter mode %d is outside 0..=2", modes[i]);
+    }
+    if (!c->has_adapter_modes && modes[i] != 0) {
+      return iree_make_status(
+          IREE_STATUS_FAILED_PRECONDITION,
+          "non-language adapter mode supplied to an incompatible bundle");
+    }
+  }
+  if (!c->has_adapter_modes) return iree_ok_status();
+  iree_hal_dim_t shape[1] = {(iree_hal_dim_t)count};
+  IREE_RETURN_IF_ERROR(xla_alloc_bv(
+      c, vector ? 1 : 0, vector ? shape : NULL,
+      IREE_HAL_ELEMENT_TYPE_INT_32, modes,
+      (iree_host_size_t)count * sizeof(int32_t), out));
+  return iree_runtime_call_inputs_push_back_buffer_view(call, *out);
+}
+
 static int xla_validate_tensor_desc(const char* name,
                                     const xla_tensor_desc* desc,
                                     int32_t expected_dtype,
@@ -608,9 +640,11 @@ xla_ctx* xla_llama_create(const char* device_uri, const char* prefill_vmfb,
                           int32_t dense_ple_hidden, int32_t model_layers,
                           int32_t deepstack_layers,
                           int32_t deepstack_visual_positions,
-                          const int32_t* deepstack_target_layers) {
+                          const int32_t* deepstack_target_layers,
+                          int32_t has_adapter_modes) {
   xla_ctx* c = (xla_ctx*)calloc(1, sizeof(xla_ctx));
   if (!c) return NULL;
+  c->has_adapter_modes = has_adapter_modes != 0;
   iree_status_t s =
       xla_llama_create_impl(c, device_uri, prefill_vmfb,
                             prefill_embeddings_vmfb,
@@ -640,7 +674,8 @@ xla_ctx* xla_llama_create(const char* device_uri, const char* prefill_vmfb,
 // prefill: tokens[lp], positions[lp] or positions[3, lp], real_len -> first token
 // id; the returned
 // KV cache becomes the resident cache decode threads forward.
-int xla_llama_prefill(xla_ctx* c, const int32_t* tokens, int32_t lp,
+int xla_llama_prefill(xla_ctx* c, int32_t adapter_mode,
+                      const int32_t* tokens, int32_t lp,
                       const int32_t* positions, int32_t real_len,
                       int32_t* out_token) {
   if (lp != c->context_capacity || real_len < 0 || real_len > lp) {
@@ -656,6 +691,9 @@ int xla_llama_prefill(xla_ctx* c, const int32_t* tokens, int32_t lp,
     XLA_CHECK(
         iree_runtime_call_inputs_push_back_buffer_view(&call, c->weights[i]));
   }
+  iree_hal_buffer_view_t* mode_bv = NULL;
+  XLA_CHECK(
+      xla_push_adapter_modes(c, &call, &adapter_mode, 1, false, &mode_bv));
   iree_hal_dim_t seq_shape[1] = {(iree_hal_dim_t)lp};
   iree_hal_dim_t position_shape[2] = {3, (iree_hal_dim_t)lp};
   iree_host_size_t seq_bytes = (iree_host_size_t)lp * sizeof(int32_t);
@@ -691,6 +729,7 @@ int xla_llama_prefill(xla_ctx* c, const int32_t* tokens, int32_t lp,
     iree_hal_buffer_view_release(tok_bv);
     iree_hal_buffer_view_release(pos_bv);
     iree_hal_buffer_view_release(len_bv);
+    if (mode_bv) iree_hal_buffer_view_release(mode_bv);
     iree_runtime_call_deinitialize(&call);
     return 1;
   }
@@ -705,13 +744,15 @@ int xla_llama_prefill(xla_ctx* c, const int32_t* tokens, int32_t lp,
   iree_hal_buffer_view_release(tok_bv);
   iree_hal_buffer_view_release(pos_bv);
   iree_hal_buffer_view_release(len_bv);
+  if (mode_bv) iree_hal_buffer_view_release(mode_bv);
   iree_runtime_call_deinitialize(&call);
   return 0;
 }
 
 // Shared decode implementation. `position_mode` is an explicit ABI contract,
 // never inferred from the position buffer rank.
-static int xla_llama_decode_impl(xla_ctx* c, int32_t token,
+static int xla_llama_decode_impl(xla_ctx* c, int32_t adapter_mode,
+                                 int32_t token,
                                  const int32_t* positions,
                                  int32_t position_mode, int32_t cache_len,
                                  int32_t* out_token) {
@@ -737,6 +778,9 @@ static int xla_llama_decode_impl(xla_ctx* c, int32_t token,
     XLA_CHECK(
         iree_runtime_call_inputs_push_back_buffer_view(&call, c->weights[i]));
   }
+  iree_hal_buffer_view_t* mode_bv = NULL;
+  XLA_CHECK(
+      xla_push_adapter_modes(c, &call, &adapter_mode, 1, false, &mode_bv));
   iree_hal_buffer_view_t* tok_bv = NULL;
   iree_hal_buffer_view_t* pos_bv = NULL;
   iree_hal_buffer_view_t* len_bv = NULL;
@@ -771,6 +815,7 @@ static int xla_llama_decode_impl(xla_ctx* c, int32_t token,
     iree_hal_buffer_view_release(tok_bv);
     iree_hal_buffer_view_release(pos_bv);
     iree_hal_buffer_view_release(len_bv);
+    if (mode_bv) iree_hal_buffer_view_release(mode_bv);
     iree_runtime_call_deinitialize(&call);
     return 1;
   }
@@ -785,6 +830,7 @@ static int xla_llama_decode_impl(xla_ctx* c, int32_t token,
   iree_hal_buffer_view_release(tok_bv);
   iree_hal_buffer_view_release(pos_bv);
   iree_hal_buffer_view_release(len_bv);
+  if (mode_bv) iree_hal_buffer_view_release(mode_bv);
   iree_runtime_call_deinitialize(&call);
   // The old KV was an input to this call; the call held a ref during invoke and
   // dropped it at deinitialize, so release our own ref now.
@@ -793,15 +839,17 @@ static int xla_llama_decode_impl(xla_ctx* c, int32_t token,
   return 0;
 }
 
-int xla_llama_decode(xla_ctx* c, int32_t token, int32_t pos,
-                     int32_t cache_len, int32_t* out_token) {
-  return xla_llama_decode_impl(c, token, &pos, 0, cache_len, out_token);
+int xla_llama_decode(xla_ctx* c, int32_t adapter_mode, int32_t token,
+                     int32_t pos, int32_t cache_len, int32_t* out_token) {
+  return xla_llama_decode_impl(c, adapter_mode, token, &pos, 0, cache_len,
+                               out_token);
 }
 
-int xla_llama_decode_mrope(xla_ctx* c, int32_t token,
+int xla_llama_decode_mrope(xla_ctx* c, int32_t adapter_mode, int32_t token,
                            const int32_t positions[3], int32_t cache_len,
                            int32_t* out_token) {
-  return xla_llama_decode_impl(c, token, positions, 1, cache_len, out_token);
+  return xla_llama_decode_impl(c, adapter_mode, token, positions, 1, cache_len,
+                               out_token);
 }
 
 // ===========================================================================
@@ -933,10 +981,10 @@ int xla_llama_ragged_reset(xla_ctx* c, int32_t bsz) {
 // DEVICE-SIDE into the slot's region of the rank-5 cache: only this slot's bytes
 // move (offset slot*per), so live slots are not disturbed and there is no host
 // round-trip. The caller (engine) samples the first token from `out_logits`.
-int xla_llama_prefill_slot_logits(xla_ctx* c, int32_t slot, const int32_t* tokens,
-                                  int32_t lp, const int32_t* positions,
-                                  int32_t real_len, int32_t vocab,
-                                  float* out_logits) {
+int xla_llama_prefill_slot_logits(
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, const int32_t* tokens,
+    int32_t lp, const int32_t* positions, int32_t real_len, int32_t vocab,
+    float* out_logits) {
   if (slot < 0 || slot >= c->rg_bsz) {
     fprintf(stderr, "xla_llama_prefill_slot_logits: slot %d out of range [0,%d)\n",
             slot, c->rg_bsz);
@@ -955,6 +1003,9 @@ int xla_llama_prefill_slot_logits(xla_ctx* c, int32_t slot, const int32_t* token
     XLA_CHECK(
         iree_runtime_call_inputs_push_back_buffer_view(&call, c->weights[i]));
   }
+  iree_hal_buffer_view_t* mode_bv = NULL;
+  XLA_CHECK(
+      xla_push_adapter_modes(c, &call, &adapter_mode, 1, false, &mode_bv));
   iree_hal_dim_t seq_shape[1] = {(iree_hal_dim_t)lp};
   iree_hal_dim_t position_shape[2] = {3, (iree_hal_dim_t)lp};
   iree_host_size_t seq_bytes = (iree_host_size_t)lp * sizeof(int32_t);
@@ -990,6 +1041,7 @@ int xla_llama_prefill_slot_logits(xla_ctx* c, int32_t slot, const int32_t* token
     iree_hal_buffer_view_release(tok_bv);
     iree_hal_buffer_view_release(pos_bv);
     iree_hal_buffer_view_release(len_bv);
+    if (mode_bv) iree_hal_buffer_view_release(mode_bv);
     iree_runtime_call_deinitialize(&call);
     return 1;
   }
@@ -999,6 +1051,7 @@ int xla_llama_prefill_slot_logits(xla_ctx* c, int32_t slot, const int32_t* token
   iree_hal_buffer_view_release(tok_bv);
   iree_hal_buffer_view_release(pos_bv);
   iree_hal_buffer_view_release(len_bv);
+  if (mode_bv) iree_hal_buffer_view_release(mode_bv);
   iree_runtime_call_deinitialize(&call);
   if (!iree_status_is_ok(rs)) {
     int code = (int)iree_status_code(rs);
@@ -1018,9 +1071,9 @@ int xla_llama_prefill_slot_logits(xla_ctx* c, int32_t slot, const int32_t* token
 // out f32 vector; K/V are handled exactly like ordinary slot prefill so the
 // same loaded context can continue through decode and greedy validation.
 int xla_llama_prefill_diagnostics_slot(
-    xla_ctx* c, int32_t slot, const int32_t* tokens, int32_t lp,
-    const int32_t* positions, int32_t real_len, int32_t diagnostic_len,
-    float* out_diagnostics) {
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, const int32_t* tokens,
+    int32_t lp, const int32_t* positions, int32_t real_len,
+    int32_t diagnostic_len, float* out_diagnostics) {
   if (!c || !c->has_prefill_diagnostics || slot < 0 || slot >= c->rg_bsz ||
       lp != c->context_capacity || real_len <= 0 || real_len > lp ||
       diagnostic_len <= 0 || !out_diagnostics) {
@@ -1031,6 +1084,7 @@ int xla_llama_prefill_diagnostics_slot(
   int rc = 0;
   bool call_initialized = false;
   iree_runtime_call_t call;
+  iree_hal_buffer_view_t* mode_bv = NULL;
   iree_hal_buffer_view_t* tok_bv = NULL;
   iree_hal_buffer_view_t* pos_bv = NULL;
   iree_hal_buffer_view_t* len_bv = NULL;
@@ -1047,6 +1101,9 @@ int xla_llama_prefill_diagnostics_slot(
         iree_runtime_call_inputs_push_back_buffer_view(&call, c->weights[i]),
         cleanup, rc);
   }
+  XLA_CHECK_GOTO(
+      xla_push_adapter_modes(c, &call, &adapter_mode, 1, false, &mode_bv),
+      cleanup, rc);
   iree_hal_dim_t seq_shape[1] = {(iree_hal_dim_t)lp};
   iree_hal_dim_t position_shape[2] = {3, (iree_hal_dim_t)lp};
   iree_host_size_t seq_bytes = (iree_host_size_t)lp * sizeof(int32_t);
@@ -1091,6 +1148,7 @@ int xla_llama_prefill_diagnostics_slot(
   rc = xla_store_prefill_kv_slot(c, slot, kc, vc);
 
 cleanup:
+  if (mode_bv) iree_hal_buffer_view_release(mode_bv);
   if (output) iree_hal_buffer_view_release(output);
   if (kc) iree_hal_buffer_view_release(kc);
   if (vc) iree_hal_buffer_view_release(vc);
@@ -1102,7 +1160,7 @@ int xla_llama_prefill_diagnostics_slot(
 }
 
 static int xla_llama_prefill_embeddings_impl(
-    xla_ctx* c, int32_t slot, int32_t position_mode,
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, int32_t position_mode,
     const xla_tensor_desc* embeddings,
     const xla_tensor_desc* dense_ple,
     const xla_tensor_desc* positions, const xla_tensor_desc* attention_bias,
@@ -1151,6 +1209,7 @@ static int xla_llama_prefill_embeddings_impl(
   int rc = 0;
   bool call_initialized = false;
   iree_runtime_call_t call;
+  iree_hal_buffer_view_t* mode_bv = NULL;
   iree_hal_buffer_view_t* embeddings_bv = NULL;
   iree_hal_buffer_view_t* dense_ple_bv = NULL;
   iree_hal_buffer_view_t* positions_bv = NULL;
@@ -1174,6 +1233,9 @@ static int xla_llama_prefill_embeddings_impl(
         iree_runtime_call_inputs_push_back_buffer_view(&call, c->weights[i]),
         cleanup, rc);
   }
+  XLA_CHECK_GOTO(
+      xla_push_adapter_modes(c, &call, &adapter_mode, 1, false, &mode_bv),
+      cleanup, rc);
   XLA_CHECK_GOTO(xla_alloc_desc_bv(c, embeddings, &embeddings_bv), cleanup, rc);
   if (dense_ple) {
     XLA_CHECK_GOTO(xla_alloc_desc_bv(c, dense_ple, &dense_ple_bv), cleanup, rc);
@@ -1231,6 +1293,7 @@ static int xla_llama_prefill_embeddings_impl(
   }
 
 cleanup:
+  if (mode_bv) iree_hal_buffer_view_release(mode_bv);
   if (output) iree_hal_buffer_view_release(output);
   if (kc) iree_hal_buffer_view_release(kc);
   if (vc) iree_hal_buffer_view_release(vc);
@@ -1244,58 +1307,60 @@ static int xla_llama_prefill_embeddings_impl(
 }
 
 int xla_llama_prefill_embeddings(
-    xla_ctx* c, int32_t position_mode, const xla_tensor_desc* embeddings,
+    xla_ctx* c, int32_t adapter_mode, int32_t position_mode,
+    const xla_tensor_desc* embeddings,
     const xla_tensor_desc* positions, const xla_tensor_desc* attention_bias,
     int32_t real_len, int32_t* out_token) {
   return xla_llama_prefill_embeddings_impl(
-      c, -1, position_mode, embeddings, NULL, positions, attention_bias,
-      real_len, 0, out_token, NULL, 0, NULL);
+      c, -1, adapter_mode, position_mode, embeddings, NULL, positions,
+      attention_bias, real_len, 0, out_token, NULL, 0, NULL);
 }
 
 int xla_llama_prefill_embeddings_slot_logits(
-    xla_ctx* c, int32_t slot, int32_t position_mode,
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, int32_t position_mode,
     const xla_tensor_desc* embeddings,
     const xla_tensor_desc* positions, const xla_tensor_desc* attention_bias,
     int32_t real_len, int32_t vocab, float* out_logits) {
   return xla_llama_prefill_embeddings_impl(
-      c, slot, position_mode, embeddings, NULL, positions, attention_bias,
-      real_len, vocab, NULL, out_logits, 0, NULL);
+      c, slot, adapter_mode, position_mode, embeddings, NULL, positions,
+      attention_bias, real_len, vocab, NULL, out_logits, 0, NULL);
 }
 
 int xla_llama_prefill_embeddings_slot_diagnostics(
-    xla_ctx* c, int32_t slot, int32_t position_mode,
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, int32_t position_mode,
     const xla_tensor_desc* embeddings,
     const xla_tensor_desc* positions, const xla_tensor_desc* attention_bias,
     int32_t real_len, int32_t vocab, int32_t kv_width, float* out_logits,
     float* out_kv) {
   return xla_llama_prefill_embeddings_impl(
-      c, slot, position_mode, embeddings, NULL, positions, attention_bias,
-      real_len, vocab, NULL, out_logits, kv_width, out_kv);
+      c, slot, adapter_mode, position_mode, embeddings, NULL, positions,
+      attention_bias, real_len, vocab, NULL, out_logits, kv_width, out_kv);
 }
 
 int xla_llama_prefill_embeddings_ple(
-    xla_ctx* c, int32_t position_mode, const xla_tensor_desc* embeddings,
+    xla_ctx* c, int32_t adapter_mode, int32_t position_mode,
+    const xla_tensor_desc* embeddings,
     const xla_tensor_desc* dense_ple, const xla_tensor_desc* positions,
     const xla_tensor_desc* attention_bias, int32_t real_len,
     int32_t* out_token) {
   return xla_llama_prefill_embeddings_impl(
-      c, -1, position_mode, embeddings, dense_ple, positions, attention_bias,
-      real_len, 0, out_token, NULL, 0, NULL);
+      c, -1, adapter_mode, position_mode, embeddings, dense_ple, positions,
+      attention_bias, real_len, 0, out_token, NULL, 0, NULL);
 }
 
 int xla_llama_prefill_embeddings_ple_slot_logits(
-    xla_ctx* c, int32_t slot, int32_t position_mode,
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, int32_t position_mode,
     const xla_tensor_desc* embeddings,
     const xla_tensor_desc* dense_ple, const xla_tensor_desc* positions,
     const xla_tensor_desc* attention_bias, int32_t real_len, int32_t vocab,
     float* out_logits) {
   return xla_llama_prefill_embeddings_impl(
-      c, slot, position_mode, embeddings, dense_ple, positions, attention_bias,
-      real_len, vocab, NULL, out_logits, 0, NULL);
+      c, slot, adapter_mode, position_mode, embeddings, dense_ple, positions,
+      attention_bias, real_len, vocab, NULL, out_logits, 0, NULL);
 }
 
 static int xla_llama_prefill_embeddings_deepstack_impl(
-    xla_ctx* c, int32_t slot, int32_t position_mode,
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, int32_t position_mode,
     const xla_tensor_desc* embeddings,
     const xla_tensor_desc* positions, const xla_tensor_desc* attention_bias,
     const xla_tensor_desc* visual_positions,
@@ -1402,6 +1467,7 @@ static int xla_llama_prefill_embeddings_deepstack_impl(
   int rc = 0;
   bool call_initialized = false;
   iree_runtime_call_t call;
+  iree_hal_buffer_view_t* mode_bv = NULL;
   iree_hal_buffer_view_t* embeddings_bv = NULL;
   iree_hal_buffer_view_t* positions_bv = NULL;
   iree_hal_buffer_view_t* len_bv = NULL;
@@ -1427,6 +1493,9 @@ static int xla_llama_prefill_embeddings_deepstack_impl(
         iree_runtime_call_inputs_push_back_buffer_view(&call, c->weights[i]),
         cleanup, rc);
   }
+  XLA_CHECK_GOTO(
+      xla_push_adapter_modes(c, &call, &adapter_mode, 1, false, &mode_bv),
+      cleanup, rc);
   XLA_CHECK_GOTO(xla_alloc_desc_bv(c, embeddings, &embeddings_bv), cleanup, rc);
   XLA_CHECK_GOTO(xla_alloc_desc_bv(c, positions, &positions_bv), cleanup, rc);
   XLA_CHECK_GOTO(xla_alloc_bv(c, 0, NULL, IREE_HAL_ELEMENT_TYPE_INT_32,
@@ -1484,6 +1553,7 @@ static int xla_llama_prefill_embeddings_deepstack_impl(
   }
 
 cleanup:
+  if (mode_bv) iree_hal_buffer_view_release(mode_bv);
   if (output) iree_hal_buffer_view_release(output);
   if (kc) iree_hal_buffer_view_release(kc);
   if (vc) iree_hal_buffer_view_release(vc);
@@ -1501,20 +1571,21 @@ static int xla_llama_prefill_embeddings_deepstack_impl(
 }
 
 int xla_llama_prefill_embeddings_deepstack(
-    xla_ctx* c, int32_t position_mode, const xla_tensor_desc* embeddings,
+    xla_ctx* c, int32_t adapter_mode, int32_t position_mode,
+    const xla_tensor_desc* embeddings,
     const xla_tensor_desc* positions, const xla_tensor_desc* attention_bias,
     const xla_tensor_desc* visual_positions,
     const xla_tensor_desc* layer_features,
     const xla_tensor_desc* layer_indices, int32_t actual_layer_count,
     int32_t actual_visual_count, int32_t real_len, int32_t* out_token) {
   return xla_llama_prefill_embeddings_deepstack_impl(
-      c, -1, position_mode, embeddings, positions, attention_bias, visual_positions,
-      layer_features, layer_indices, actual_layer_count, actual_visual_count,
-      real_len, 0, out_token, NULL);
+      c, -1, adapter_mode, position_mode, embeddings, positions,
+      attention_bias, visual_positions, layer_features, layer_indices,
+      actual_layer_count, actual_visual_count, real_len, 0, out_token, NULL);
 }
 
 int xla_llama_prefill_embeddings_deepstack_slot_logits(
-    xla_ctx* c, int32_t slot, int32_t position_mode,
+    xla_ctx* c, int32_t slot, int32_t adapter_mode, int32_t position_mode,
     const xla_tensor_desc* embeddings,
     const xla_tensor_desc* positions, const xla_tensor_desc* attention_bias,
     const xla_tensor_desc* visual_positions,
@@ -1523,9 +1594,10 @@ int xla_llama_prefill_embeddings_deepstack_slot_logits(
     int32_t actual_visual_count, int32_t real_len, int32_t vocab,
     float* out_logits) {
   return xla_llama_prefill_embeddings_deepstack_impl(
-      c, slot, position_mode, embeddings, positions, attention_bias, visual_positions,
-      layer_features, layer_indices, actual_layer_count, actual_visual_count,
-      real_len, vocab, NULL, out_logits);
+      c, slot, adapter_mode, position_mode, embeddings, positions,
+      attention_bias, visual_positions, layer_features, layer_indices,
+      actual_layer_count, actual_visual_count, real_len, vocab, NULL,
+      out_logits);
 }
 
 // Ragged decode_step: token[B], pos[B], cache_len[B] (per row), rank-5 KV ->
@@ -1533,7 +1605,8 @@ int xla_llama_prefill_embeddings_deepstack_slot_logits(
 // KV in place. The caller (engine) samples a token per row. Inactive rows
 // (token/pos/cache_len 0) are masked no-ops whose logits the caller discards.
 static int xla_llama_decode_ragged_impl(
-    xla_ctx* c, int32_t bsz, const int32_t* tokens,
+    xla_ctx* c, int32_t bsz, const int32_t* adapter_modes,
+    const int32_t* tokens,
     const int32_t* positions, int32_t position_mode,
     const int32_t* cache_len, int32_t vocab, float* out_logits) {
   if (bsz != c->rg_bsz || !c->kcache_b || !c->vcache_b ||
@@ -1569,6 +1642,9 @@ static int xla_llama_decode_ragged_impl(
     XLA_CHECK(
         iree_runtime_call_inputs_push_back_buffer_view(&call, c->weights[i]));
   }
+  iree_hal_buffer_view_t* mode_bv = NULL;
+  XLA_CHECK(
+      xla_push_adapter_modes(c, &call, adapter_modes, bsz, true, &mode_bv));
   iree_hal_dim_t vshape[1] = {(iree_hal_dim_t)bsz};
   iree_hal_dim_t position_shape[2] = {(iree_hal_dim_t)bsz, 3};
   iree_host_size_t vbytes = (iree_host_size_t)bsz * sizeof(int32_t);
@@ -1608,6 +1684,7 @@ static int xla_llama_decode_ragged_impl(
     iree_hal_buffer_view_release(tok_bv);
     iree_hal_buffer_view_release(pos_bv);
     iree_hal_buffer_view_release(len_bv);
+    if (mode_bv) iree_hal_buffer_view_release(mode_bv);
     iree_runtime_call_deinitialize(&call);
     return 1;
   }
@@ -1624,6 +1701,7 @@ static int xla_llama_decode_ragged_impl(
   iree_hal_buffer_view_release(tok_bv);
   iree_hal_buffer_view_release(pos_bv);
   iree_hal_buffer_view_release(len_bv);
+  if (mode_bv) iree_hal_buffer_view_release(mode_bv);
   iree_runtime_call_deinitialize(&call);
   iree_hal_buffer_view_release(old_k);
   iree_hal_buffer_view_release(old_v);
@@ -1631,20 +1709,22 @@ static int xla_llama_decode_ragged_impl(
 }
 
 int xla_llama_decode_ragged_logits(xla_ctx* c, int32_t bsz,
+                                   const int32_t* adapter_modes,
                                    const int32_t* tokens,
                                    const int32_t* positions,
                                    const int32_t* cache_len, int32_t vocab,
                                    float* out_logits) {
-  return xla_llama_decode_ragged_impl(c, bsz, tokens, positions, 0, cache_len,
-                                      vocab, out_logits);
+  return xla_llama_decode_ragged_impl(c, bsz, adapter_modes, tokens, positions,
+                                      0, cache_len, vocab, out_logits);
 }
 
 int xla_llama_decode_ragged_mrope_logits(
-    xla_ctx* c, int32_t bsz, const int32_t* tokens,
+    xla_ctx* c, int32_t bsz, const int32_t* adapter_modes,
+    const int32_t* tokens,
     const int32_t* positions, const int32_t* cache_len, int32_t vocab,
     float* out_logits) {
-  return xla_llama_decode_ragged_impl(c, bsz, tokens, positions, 1, cache_len,
-                                      vocab, out_logits);
+  return xla_llama_decode_ragged_impl(c, bsz, adapter_modes, tokens, positions,
+                                      1, cache_len, vocab, out_logits);
 }
 
 void xla_llama_free(xla_ctx* c) {
diff --git a/src/lib/mlxcel-xla/src/batch.rs b/src/lib/mlxcel-xla/src/batch.rs
index cb31b1df2..b08cb4b2f 100644
--- a/src/lib/mlxcel-xla/src/batch.rs
+++ b/src/lib/mlxcel-xla/src/batch.rs
@@ -131,6 +131,8 @@ pub enum EngineEvent {
 /// An active slot's running state.
 struct Slot {
     req_id: u64,
+    /// Request-scoped Phi4MM language/speech/vision LoRA selector.
+    adapter_mode: i32,
     /// Last emitted token; the next decode input for this row.
     cur: i32,
     /// Tokens currently in this slot's KV (== the next write position).
@@ -195,6 +197,13 @@ impl PendingInput {
             }
         }
     }
+
+    fn adapter_mode(&self) -> i32 {
+        match self {
+            Self::Prepared(prepared) => prepared.adapter_mode,
+            Self::Tokens(_) | Self::Gemma3nPrepared { .. } | Self::DeepStackPrepared { .. } => 0,
+        }
+    }
 }
 
 /// A queued, not-yet-admitted request.
@@ -680,6 +689,7 @@ impl XlaBatchEngine {
             });
             let slot = Slot {
                 req_id: p.req_id,
+                adapter_mode: p.input.adapter_mode(),
                 cur: first,
                 cache_len: p.input.effective_len() as i32,
                 rope_delta: p.input.rope_delta(),
@@ -708,22 +718,31 @@ impl XlaBatchEngine {
         // zeros (masked no-ops) and their logits are discarded.
         let b = self.sched.b_max;
         let mut tok = vec![0i32; b];
+        let mut adapter_modes = vec![0i32; b];
         let mut clen = vec![0i32; b];
         let mut pos = vec![0i32; b];
         for (s, slot) in self.sched.slots.iter().enumerate() {
             if let Some(st) = slot {
+                adapter_modes[s] = st.adapter_mode;
                 tok[s] = st.cur;
                 clen[s] = st.cache_len;
                 pos[s] = st.cache_len;
             }
         }
         let logits = match self.engine.position_mode() {
-            PreparedPositionMode::OneD => self.engine.decode_ragged_logits(&tok, &pos, &clen)?,
+            PreparedPositionMode::OneD => {
+                self.engine
+                    .decode_ragged_logits_with_modes(&adapter_modes, &tok, &pos, &clen)?
+            }
             PreparedPositionMode::Mrope3D => {
                 let mrope_positions =
                     mrope_slot_coordinates(&self.sched.slots).map_err(|error| error.to_string())?;
-                self.engine
-                    .decode_ragged_mrope_logits(&tok, &mrope_positions, &clen)?
+                self.engine.decode_ragged_mrope_logits_with_modes(
+                    &adapter_modes,
+                    &tok,
+                    &mrope_positions,
+                    &clen,
+                )?
             }
         };
         let vocab = self.engine.vocab();
@@ -1253,8 +1272,8 @@ impl XlaReferenceEngine {
 mod tests {
     use super::*;
     use mlxcel_core::session::{
-        OwnedTensor, PreparedAttentionBias, PreparedModality, PreparedPositions,
-        PreparedTensorDType,
+        OwnedTensor, PreparedAdapterMode, PreparedAttentionBias, PreparedModality,
+        PreparedPositions, PreparedTensorDType,
     };
 
     const EOS: [i32; 1] = [42];
@@ -1273,6 +1292,15 @@ mod tests {
     }
 
     fn prepared_input(tokens: Vec<i32>, hidden: usize, capacity: usize) -> PreparedIreePrefill {
+        prepared_input_with_mode(tokens, hidden, capacity, PreparedAdapterMode::Language)
+    }
+
+    fn prepared_input_with_mode(
+        tokens: Vec<i32>,
+        hidden: usize,
+        capacity: usize,
+        adapter_mode: PreparedAdapterMode,
+    ) -> PreparedIreePrefill {
         let sequence = tokens.len();
         let embedding_bytes = vec![0; sequence * hidden * std::mem::size_of::<f32>()];
         let bias_bytes = vec![0; sequence * std::mem::size_of::<f32>()];
@@ -1303,7 +1331,8 @@ mod tests {
                 token_count: 1,
             }],
         )
-        .unwrap();
+        .unwrap()
+        .with_adapter_mode(adapter_mode);
         PreparedIreePrefill::prepare(&value, hidden, capacity).unwrap()
     }
 
@@ -1381,6 +1410,7 @@ mod tests {
         let req_id = pending.req_id;
         scheduler.slots[slot_index] = Some(Slot {
             req_id,
+            adapter_mode: pending.input.adapter_mode(),
             cur: 7,
             cache_len: pending.input.effective_len() as i32,
             rope_delta: pending.input.rope_delta(),
@@ -1448,6 +1478,7 @@ mod tests {
         let mut s = Scheduler::new(2, EOS.to_vec());
         s.slots[0] = Some(Slot {
             req_id: 41,
+            adapter_mode: 0,
             cur: 7,
             cache_len: 300,
             rope_delta: 0,
@@ -1601,6 +1632,7 @@ mod tests {
 
         s.slots[0] = Some(Slot {
             req_id: queued_prepared,
+            adapter_mode: 0,
             cur: 7,
             cache_len: effective_len as i32,
             rope_delta: 0,
@@ -1634,6 +1666,44 @@ mod tests {
         assert_eq!(s.free_slots(), vec![0]);
     }
 
+    #[test]
+    fn adapter_mode_is_request_scoped_and_cleared_on_slot_reuse() {
+        let mut scheduler = Scheduler::new(2, EOS.to_vec());
+        let speech = prepared_input_with_mode(vec![4, 5, 6], 2, 8, PreparedAdapterMode::Speech);
+        let speech_id =
+            scheduler.submit_input(PendingInput::Prepared(speech), 8, SampleParams::greedy());
+        assert_eq!(admit_next(&mut scheduler, 0), speech_id);
+        assert_eq!(
+            scheduler.slots[0].as_ref().unwrap().adapter_mode,
+            PreparedAdapterMode::Speech.code()
+        );
+
+        assert!(scheduler.cancel(speech_id));
+        let language_id = scheduler.submit(vec![9, 10], 4, g());
+        assert_eq!(admit_next(&mut scheduler, 0), language_id);
+        assert_eq!(
+            scheduler.slots[0].as_ref().unwrap().adapter_mode,
+            PreparedAdapterMode::Language.code(),
+            "slot reuse must not leak the cancelled request's adapter"
+        );
+
+        let vision = prepared_input_with_mode(vec![7, 8], 2, 8, PreparedAdapterMode::Vision);
+        let vision_id =
+            scheduler.submit_input(PendingInput::Prepared(vision), 8, SampleParams::greedy());
+        assert_eq!(admit_next(&mut scheduler, 1), vision_id);
+        assert_eq!(
+            scheduler
+                .slots
+                .iter()
+                .map(|slot| slot.as_ref().map_or(0, |slot| slot.adapter_mode))
+                .collect::<Vec<_>>(),
+            vec![
+                PreparedAdapterMode::Language.code(),
+                PreparedAdapterMode::Vision.code()
+            ]
+        );
+    }
+
     #[test]
     fn deepstack_side_tensors_are_dropped_on_cancel_and_slot_can_be_reused() {
         let mut s = Scheduler::new(1, EOS.to_vec());
@@ -1669,6 +1739,7 @@ mod tests {
         let p = s.pop_next_pending().unwrap();
         s.slots[0] = Some(Slot {
             req_id: p.req_id,
+            adapter_mode: 0,
             cur: 5,
             cache_len: 1,
             rope_delta: 0,
diff --git a/src/lib/mlxcel-xla/src/emitter/builder.rs b/src/lib/mlxcel-xla/src/emitter/builder.rs
index 5f37bdc87..ae182c154 100644
--- a/src/lib/mlxcel-xla/src/emitter/builder.rs
+++ b/src/lib/mlxcel-xla/src/emitter/builder.rs
@@ -728,6 +728,31 @@ impl Builder {
         Val { name: r, ty }
     }
 
+    /// Dense i32 constant tensor from row-major values.
+    pub fn const_tensor_i32(&mut self, data: &[i32], shape: Vec<usize>) -> Val {
+        assert_eq!(data.len(), shape.iter().product::<usize>());
+        let flat_ty = Ty::new(vec![data.len()], "i32");
+        let values = data
+            .iter()
+            .map(i32::to_string)
+            .collect::<Vec<_>>()
+            .join(", ");
+        let r = self.fresh();
+        self.line(format!(
+            "{r} = stablehlo.constant dense<[{values}]> : {}",
+            flat_ty.render()
+        ));
+        let flat = Val {
+            name: r,
+            ty: flat_ty,
+        };
+        if shape == flat.ty.shape {
+            flat
+        } else {
+            self.reshape(&flat, shape)
+        }
+    }
+
     // --- structural --------------------------------------------------------
 
     pub fn iota(&mut self, n: usize) -> Val {
@@ -1040,6 +1065,179 @@ impl Builder {
         Val { name: r, ty }
     }
 
+    /// Gather rows from `[N, M]` with an arbitrary-rank index tensor whose final
+    /// dimension is the one-element index vector. The output is the leading
+    /// index shape followed by `M`.
+    pub fn gather_rows_nd(&mut self, operand: &Val, indices: &Val) -> Val {
+        assert_eq!(operand.ty.shape.len(), 2);
+        assert_eq!(indices.ty.elt, "i32");
+        assert_eq!(indices.ty.shape.last(), Some(&1));
+        let mut shape = indices.ty.shape[..indices.ty.shape.len() - 1].to_vec();
+        shape.push(operand.ty.shape[1]);
+        let ty = Ty::new(shape.clone(), operand.ty.elt);
+        let offset_dim = shape.len() - 1;
+        let index_vector_dim = indices.ty.shape.len() - 1;
+        let r = self.fresh();
+        self.line(format!(
+            "{r} = \"stablehlo.gather\"({}, {}) <{{dimension_numbers = #stablehlo.gather<offset_dims = [{offset_dim}], collapsed_slice_dims = [0], start_index_map = [0], index_vector_dim = {index_vector_dim}>, slice_sizes = array<i64: 1, {}>}}> : ({}, {}) -> {}",
+            operand.name,
+            indices.name,
+            operand.ty.shape[1],
+            operand.ty.render(),
+            indices.ty.render(),
+            ty.render()
+        ));
+        Val { name: r, ty }
+    }
+
+    /// Positive-edge padding with a scalar value. Interior padding is zero.
+    pub fn pad(&mut self, operand: &Val, value: &Val, low: &[usize], high: &[usize]) -> Val {
+        assert_eq!(low.len(), operand.ty.shape.len());
+        assert_eq!(high.len(), operand.ty.shape.len());
+        assert!(value.ty.shape.is_empty());
+        assert_eq!(value.ty.elt, operand.ty.elt);
+        let shape = operand
+            .ty
+            .shape
+            .iter()
+            .zip(low)
+            .zip(high)
+            .map(|((&dim, &low), &high)| dim + low + high)
+            .collect::<Vec<_>>();
+        let ty = Ty::new(shape, operand.ty.elt);
+        let render = |values: &[usize]| {
+            values
+                .iter()
+                .map(usize::to_string)
+                .collect::<Vec<_>>()
+                .join(", ")
+        };
+        let interior = vec![0; low.len()];
+        let r = self.fresh();
+        self.line(format!(
+            "{r} = \"stablehlo.pad\"({}, {}) {{edge_padding_low = array<i64: {}>, edge_padding_high = array<i64: {}>, interior_padding = array<i64: {}>}} : ({}, {}) -> {}",
+            operand.name,
+            value.name,
+            render(low),
+            render(high),
+            render(&interior),
+            operand.ty.render(),
+            value.ty.render(),
+            ty.render()
+        ));
+        Val { name: r, ty }
+    }
+
+    /// Channels-last 1-D/2-D convolution. `kernel` is
+    /// `[spatial..., input_per_group, output]`.
+    pub fn convolution(
+        &mut self,
+        input: &Val,
+        kernel: &Val,
+        strides: &[usize],
+        padding: &[(usize, usize)],
+        feature_groups: usize,
+    ) -> Val {
+        assert!(feature_groups > 0);
+        assert_eq!(
+            input.ty.elt, kernel.ty.elt,
+            "convolution input/kernel dtypes must match"
+        );
+        let spatial_rank = input.ty.shape.len() - 2;
+        assert!(matches!(spatial_rank, 1 | 2));
+        assert_eq!(kernel.ty.shape.len(), input.ty.shape.len());
+        assert_eq!(strides.len(), spatial_rank);
+        assert!(strides.iter().all(|stride| *stride > 0));
+        assert_eq!(padding.len(), spatial_rank);
+        assert_eq!(input.ty.shape[spatial_rank + 1] % feature_groups, 0);
+        assert_eq!(
+            kernel.ty.shape[spatial_rank],
+            input.ty.shape[spatial_rank + 1] / feature_groups
+        );
+        assert!(kernel.ty.shape[spatial_rank + 1] > 0);
+        let mut shape = Vec::with_capacity(input.ty.shape.len());
+        shape.push(input.ty.shape[0]);
+        for axis in 0..spatial_rank {
+            let padded = input.ty.shape[axis + 1] + padding[axis].0 + padding[axis].1;
+            let kernel_size = kernel.ty.shape[axis];
+            assert!(
+                padded >= kernel_size,
+                "convolution kernel exceeds padded input on spatial axis {axis}"
+            );
+            shape.push((padded - kernel_size) / strides[axis] + 1);
+        }
+        shape.push(kernel.ty.shape[spatial_rank + 1]);
+        let output_ty = Ty::new(shape, input.ty.elt);
+        let input_demoted;
+        let kernel_demoted;
+        let (input, kernel) = match self.precision.dot_elt() {
+            Some(elt) => {
+                input_demoted = if input.ty.elt == "f32" {
+                    self.convert(input, elt)
+                } else {
+                    input.clone()
+                };
+                kernel_demoted = if kernel.ty.elt == "f32" {
+                    self.convert(kernel, elt)
+                } else {
+                    kernel.clone()
+                };
+                (&input_demoted, &kernel_demoted)
+            }
+            None => (input, kernel),
+        };
+        let convolution_ty = Ty::new(output_ty.shape.clone(), input.ty.elt);
+        let array = |values: &[usize]| {
+            values
+                .iter()
+                .map(usize::to_string)
+                .collect::<Vec<_>>()
+                .join(", ")
+        };
+        let padding = padding
+            .iter()
+            .map(|&(low, high)| format!("[{low}, {high}]"))
+            .collect::<Vec<_>>()
+            .join(", ");
+        let spatial = (0..spatial_rank)
+            .map(|axis| axis.to_string())
+            .collect::<Vec<_>>()
+            .join(", ");
+        let input_dims = format!("[b, {spatial}, f]");
+        let kernel_dims = format!("[{spatial}, i, o]");
+        let output_dims = input_dims.clone();
+        let dilation = vec![1; spatial_rank];
+        let reversal = vec!["false"; spatial_rank].join(", ");
+        let r = self.fresh();
+        self.line(format!(
+            "{r} = \"stablehlo.convolution\"({}, {}) {{window_strides = array<i64: {}>, padding = dense<[{}]> : tensor<{}x2xi64>, lhs_dilation = array<i64: {}>, rhs_dilation = array<i64: {}>, window_reversal = array<i1: {}>, dimension_numbers = #stablehlo.conv<{}x{}->{}>, batch_group_count = 1 : i64, feature_group_count = {} : i64, precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]}} : ({}, {}) -> {}",
+            input.name,
+            kernel.name,
+            array(strides),
+            padding,
+            spatial_rank,
+            array(&dilation),
+            array(&dilation),
+            reversal,
+            input_dims,
+            kernel_dims,
+            output_dims,
+            feature_groups,
+            input.ty.render(),
+            kernel.ty.render(),
+            convolution_ty.render()
+        ));
+        let convolved = Val {
+            name: r,
+            ty: convolution_ty,
+        };
+        if convolved.ty.elt == output_ty.elt {
+            convolved
+        } else {
+            self.convert(&convolved, output_ty.elt)
+        }
+    }
+
     // --- elementwise -------------------------------------------------------
 
     fn binary(&mut self, op: &str, a: &Val, b: &Val) -> Val {
@@ -1065,6 +1263,9 @@ impl Builder {
     pub fn multiply(&mut self, a: &Val, b: &Val) -> Val {
         self.binary("multiply", a, b)
     }
+    pub fn maximum(&mut self, a: &Val, b: &Val) -> Val {
+        self.binary("maximum", a, b)
+    }
     pub fn divide(&mut self, a: &Val, b: &Val) -> Val {
         self.binary("divide", a, b)
     }
diff --git a/src/lib/mlxcel-xla/src/emitter/config.rs b/src/lib/mlxcel-xla/src/emitter/config.rs
index 918dd4055..8a4739bb7 100644
--- a/src/lib/mlxcel-xla/src/emitter/config.rs
+++ b/src/lib/mlxcel-xla/src/emitter/config.rs
@@ -42,6 +42,22 @@ pub enum RopeScaling {
         high_freq_factor: f64,
         orig_ctx: usize,
     },
+    /// Phi-3/Phi-4 SuScaled LongRoPE. The released Phi4MM runtime always
+    /// materializes the checkpoint's long-factor table and applies the
+    /// amplitude correction derived from max/original context.
+    LongRope { factors: Vec<f64>, amplitude: f64 },
+}
+
+/// Official Phi4MM request-time LoRA contract.
+///
+/// The ranks and scales affect graph shapes and numerics, so they are part of
+/// the emitted artifact identity rather than mutable serving configuration.
+#[derive(Clone, Copy, Debug, PartialEq)]
+pub struct Phi4LoraConfig {
+    pub speech_rank: usize,
+    pub speech_scale: f32,
+    pub vision_rank: usize,
+    pub vision_scale: f32,
 }
 
 /// Per-layer norm placement. The three dense patterns differ in where the
@@ -384,6 +400,9 @@ pub struct Config {
     /// loader splits it (gate first, up second) into the emitter's `gate`/`up` args.
     /// Consumed by the weight loader (`iree.rs`); the emitter graph is unaffected.
     pub fused_gate_up: bool,
+    /// Per-request speech/vision adapters carried by the Phi4MM checkpoint.
+    /// `None` preserves the existing dense graph signatures byte-for-byte.
+    pub phi4_lora: Option<Phi4LoraConfig>,
 }
 
 impl Config {
@@ -441,6 +460,7 @@ impl Config {
             logit_div: None,
             fused_qkv: false,
             fused_gate_up: false,
+            phi4_lora: None,
         }
     }
 
@@ -460,10 +480,9 @@ impl Config {
     pub fn from_json_str(s: &str) -> Result<Self, String> {
         let root: serde_json::Value =
             serde_json::from_str(s).map_err(|e| format!("parse config.json: {e}"))?;
-        let v = if matches!(
-            root.get("model_type").and_then(serde_json::Value::as_str),
-            Some("llava" | "llava_next")
-        ) {
+        let wrapper_model_type = root.get("model_type").and_then(serde_json::Value::as_str);
+        let is_phi4mm = wrapper_model_type == Some("phi4mm");
+        let v = if matches!(wrapper_model_type, Some("llava" | "llava_next")) {
             let mut text = root
                 .get("text_config")
                 .and_then(serde_json::Value::as_object)
@@ -841,7 +860,7 @@ impl Config {
                 sliding_pattern = ou("sliding_window_pattern").unwrap_or(4).max(1);
                 logit_mul = Some(of("logit_scale").unwrap_or(0.0625) as f32);
             }
-            Some("phi3") => {
+            Some("phi3" | "phi4mm") => {
                 // Fused qkv_proj / gate_up_proj (split at load); RMSNorm; untied.
                 fused_qkv = true;
                 fused_gate_up = true;
@@ -981,7 +1000,7 @@ impl Config {
                 return Err(format!(
                     "the OpenXLA emitter supports the dense architectures Llama, Qwen2, \
                      Qwen3, Gemma1/2/3, SmolLM3, OLMo2/3, Seed-OSS, MiMo, InternLM3, ExaOne, \
-                     Cohere, Cohere2, Phi3, StableLM, StarCoder2, Granite, and MiniCPM, plus \
+                     Cohere, Cohere2, Phi3, Phi4MM, StableLM, StarCoder2, Granite, and MiniCPM, plus \
                      the Mixtral, Qwen2-MoE, Qwen3-MoE, and OLMoE mixture-of-experts \
                      architectures; config.json model_type = {other:?} (other MoE / MLA / \
                      novel-activation variants are follow-ups)"
@@ -1076,6 +1095,50 @@ impl Config {
                             orig_ctx,
                         }
                     }
+                    Some("longrope") if is_phi4mm => {
+                        let values = scaling
+                            .get("long_factor")
+                            .and_then(serde_json::Value::as_array)
+                            .ok_or_else(|| {
+                                "Phi4MM LongRoPE requires rope_scaling.long_factor".to_string()
+                            })?;
+                        let factors = values
+                            .iter()
+                            .enumerate()
+                            .map(|(index, value)| {
+                                value.as_f64().ok_or_else(|| {
+                                    format!(
+                                        "Phi4MM rope_scaling.long_factor[{index}] must be finite"
+                                    )
+                                })
+                            })
+                            .collect::<Result<Vec<_>, _>>()?;
+                        if factors.len() != rotary_dim.unwrap_or(head_dim) / 2
+                            || factors
+                                .iter()
+                                .any(|factor| !factor.is_finite() || *factor <= 0.0)
+                        {
+                            return Err(format!(
+                                "Phi4MM LongRoPE requires {} positive finite factors, got {}",
+                                rotary_dim.unwrap_or(head_dim) / 2,
+                                factors.len()
+                            ));
+                        }
+                        let original = ou("original_max_position_embeddings").ok_or_else(|| {
+                            "Phi4MM LongRoPE requires original_max_position_embeddings".to_string()
+                        })?;
+                        let maximum = ou("max_position_embeddings").ok_or_else(|| {
+                            "Phi4MM LongRoPE requires max_position_embeddings".to_string()
+                        })?;
+                        if original <= 1 || maximum < original {
+                            return Err(format!(
+                                "invalid Phi4MM LongRoPE context bounds: original={original}, maximum={maximum}"
+                            ));
+                        }
+                        let ratio = maximum as f64 / original as f64;
+                        let amplitude = (1.0 + ratio.ln() / (original as f64).ln()).sqrt();
+                        RopeScaling::LongRope { factors, amplitude }
+                    }
                     other => {
                         return Err(format!(
                             "the OpenXLA emitter supports plain / default / (in-context) dynamic \
@@ -1243,6 +1306,38 @@ impl Config {
         // HF `PretrainedConfig` defaults this to `true`, but some arches default
         // untied (`tie_default`); an explicit config field always wins.
         let tie_word_embeddings = ob("tie_word_embeddings").unwrap_or(tie_default);
+        let phi4_lora = if is_phi4mm {
+            let parse = |name: &str| -> Result<(usize, f32), String> {
+                let lora = v
+                    .get(name)
+                    .and_then(serde_json::Value::as_object)
+                    .ok_or_else(|| format!("Phi4MM config missing object `{name}`"))?;
+                let rank = lora
+                    .get("r")
+                    .and_then(serde_json::Value::as_u64)
+                    .and_then(|value| usize::try_from(value).ok())
+                    .filter(|rank| *rank > 0)
+                    .ok_or_else(|| format!("Phi4MM `{name}.r` must be positive"))?;
+                let alpha = lora
+                    .get("lora_alpha")
+                    .and_then(serde_json::Value::as_f64)
+                    .filter(|alpha| alpha.is_finite() && *alpha > 0.0)
+                    .ok_or_else(|| {
+                        format!("Phi4MM `{name}.lora_alpha` must be positive and finite")
+                    })?;
+                Ok((rank, (alpha / rank as f64) as f32))
+            };
+            let (speech_rank, speech_scale) = parse("speech_lora")?;
+            let (vision_rank, vision_scale) = parse("vision_lora")?;
+            Some(Phi4LoraConfig {
+                speech_rank,
+                speech_scale,
+                vision_rank,
+                vision_scale,
+            })
+        } else {
+            None
+        };
 
         Ok(Config {
             context_capacity: crate::DEFAULT_CONTEXT_CAPACITY,
@@ -1291,6 +1386,7 @@ impl Config {
             logit_div,
             fused_qkv,
             fused_gate_up,
+            phi4_lora,
         })
     }
 
@@ -1683,4 +1779,75 @@ mod tests {
             })
         );
     }
+
+    #[test]
+    fn phi4mm_pins_longrope_and_both_adapter_contracts() {
+        let config = Config::from_json_str(
+            r#"{
+                "model_type":"phi4mm","hidden_size":8,"intermediate_size":16,
+                "num_hidden_layers":2,"num_attention_heads":2,"num_key_value_heads":1,
+                "head_dim":4,"partial_rotary_factor":1.0,"rms_norm_eps":1e-5,
+                "rope_theta":10000.0,"vocab_size":32,"tie_word_embeddings":true,
+                "max_position_embeddings":16,"original_max_position_embeddings":4,
+                "hidden_act":"silu",
+                "rope_scaling":{"type":"longrope","long_factor":[1.0,2.0]},
+                "speech_lora":{"r":3,"lora_alpha":6},
+                "vision_lora":{"r":2,"lora_alpha":5}
+            }"#,
+        )
+        .expect("Phi4MM config");
+        assert!(config.fused_qkv);
+        assert!(config.fused_gate_up);
+        assert_eq!(
+            config.phi4_lora,
+            Some(Phi4LoraConfig {
+                speech_rank: 3,
+                speech_scale: 2.0,
+                vision_rank: 2,
+                vision_scale: 2.5,
+            })
+        );
+        let RopeScaling::LongRope { factors, amplitude } = config.rope else {
+            panic!("Phi4MM must preserve LongRoPE");
+        };
+        assert_eq!(factors, vec![1.0, 2.0]);
+        let expected = (1.0 + (4.0f64).ln() / (4.0f64).ln()).sqrt();
+        assert!((amplitude - expected).abs() < f64::EPSILON);
+    }
+
+    #[test]
+    fn phi4mm_rejects_incomplete_lora_and_longrope_shapes() {
+        let base = r#"{
+            "model_type":"phi4mm","hidden_size":8,"intermediate_size":16,
+            "num_hidden_layers":2,"num_attention_heads":2,"num_key_value_heads":1,
+            "head_dim":4,"partial_rotary_factor":1.0,"rms_norm_eps":1e-5,
+            "rope_theta":10000.0,"vocab_size":32,
+            "max_position_embeddings":16,"original_max_position_embeddings":4,
+            "hidden_act":"silu",
+            "rope_scaling":{"type":"longrope","long_factor":[1.0,2.0]},
+            "speech_lora":{"r":3,"lora_alpha":6},
+            "vision_lora":{"r":2,"lora_alpha":5}
+        }"#;
+        let missing = base.replace(
+            r#""vision_lora":{"r":2,"lora_alpha":5}"#,
+            r#""removed_vision_lora":null"#,
+        );
+        assert!(
+            Config::from_json_str(&missing)
+                .unwrap_err()
+                .contains("vision_lora")
+        );
+        let wrong_factors = base.replace("[1.0,2.0]", "[1.0]");
+        assert!(
+            Config::from_json_str(&wrong_factors)
+                .unwrap_err()
+                .contains("requires 2")
+        );
+        let invalid_scale = base.replace(r#""lora_alpha":6"#, r#""lora_alpha":0"#);
+        assert!(
+            Config::from_json_str(&invalid_scale)
+                .unwrap_err()
+                .contains("positive and finite")
+        );
+    }
 }
diff --git a/src/lib/mlxcel-xla/src/emitter/mod.rs b/src/lib/mlxcel-xla/src/emitter/mod.rs
index 0bfc9c4f9..2f998a4fc 100644
--- a/src/lib/mlxcel-xla/src/emitter/mod.rs
+++ b/src/lib/mlxcel-xla/src/emitter/mod.rs
@@ -57,6 +57,7 @@ mod gemma3n_schema;
 mod gemma3n_weights;
 mod model;
 mod moe;
+mod phi4_audio;
 mod rope;
 mod vision;
 mod vision_config;
@@ -66,7 +67,8 @@ mod context_tests;
 
 #[allow(unused_imports)]
 pub(crate) use config::{
-    Config, DeepStackConfig, MropeConfig, MropeLayout, NormStyle, QkNorm, QuantConfig, WeightScheme,
+    Config, DeepStackConfig, MropeConfig, MropeLayout, NormStyle, Phi4LoraConfig, QkNorm,
+    QuantConfig, WeightScheme,
 };
 #[allow(unused_imports)]
 pub(crate) use gemma3n::{
@@ -111,6 +113,12 @@ pub use gemma3n_qmv::{
 };
 #[allow(unused_imports)]
 pub(crate) use gemma3n_weights::{Gemma3nWeightSpec, gemma3n_weight_specs};
+#[allow(unused_imports)]
+pub(crate) use phi4_audio::{
+    Phi4AudioConfig, Phi4AudioDiagnosticSpec, Phi4AudioWeightSpec, emit_phi4_audio,
+    emit_phi4_audio_diagnostic_with, emit_phi4_audio_with, phi4_audio_diagnostic_specs,
+    phi4_audio_weight_specs, validate_phi4_audio_weight_shapes,
+};
 // MoE FFN config types (issue #500), read by the weight loader (`iree.rs`) and the
 // validation harness. `SharedExpertConfig` is only named in some build cfgs.
 #[allow(unused_imports)]
@@ -143,7 +151,7 @@ pub(crate) use vision_config::{LlavaVisionConfig, VisionActivation, VisionWeight
 
 #[cfg(test)]
 mod tests {
-    use super::config::{NormStyle, QkNorm, RopeScaling};
+    use super::config::{NormStyle, Phi4LoraConfig, QkNorm, RopeScaling};
     use super::*;
 
     const CONFIG_JSON: &str = include_str!("../../assets/llama-3.2-1b/config.json");
@@ -204,6 +212,69 @@ mod tests {
         }
     }
 
+    fn phi4_like() -> Config {
+        Config {
+            context_capacity: 8,
+            hidden: 8,
+            inter: 16,
+            n_layers: 2,
+            n_q: 2,
+            n_kv: 1,
+            head_dim: 4,
+            rotary_dim: Some(4),
+            vocab: 32,
+            tie_word_embeddings: true,
+            fused_qkv: true,
+            fused_gate_up: true,
+            phi4_lora: Some(Phi4LoraConfig {
+                speech_rank: 3,
+                speech_scale: 2.0,
+                vision_rank: 2,
+                vision_scale: 2.5,
+            }),
+            ..Config::llama_3_2_1b()
+        }
+    }
+
+    #[test]
+    fn phi4mm_graphs_carry_explicit_scalar_and_per_row_adapter_modes() {
+        let config = phi4_like();
+        let prefill = emit_prefill(&config, false);
+        let decode = emit_decode(&config, false);
+        let ragged = emit_decode_ragged(&config, 4, false);
+        assert!(prefill.contains("tensor<i32> loc(\"adapter_mode\")"));
+        assert!(decode.contains("tensor<i32> loc(\"adapter_mode\")"));
+        assert!(ragged.contains("tensor<4xi32> loc(\"adapter_mode\")"));
+        for graph in [&prefill, &decode, &ragged] {
+            for adapter in ["speech", "vision"] {
+                for projection in ["qkv", "o", "gate_up", "down"] {
+                    assert!(
+                        graph.contains(&format!("{projection}.lora_A.{adapter}")),
+                        "missing {projection} A/{adapter}"
+                    );
+                    assert!(
+                        graph.contains(&format!("{projection}.lora_B.{adapter}")),
+                        "missing {projection} B/{adapter}"
+                    );
+                }
+            }
+            assert!(
+                graph.contains("stablehlo.compare"),
+                "mode gating must compare the request input"
+            );
+            assert!(
+                graph.contains("stablehlo.select"),
+                "mode gating must select adapter deltas without mutating base weights"
+            );
+        }
+    }
+
+    #[test]
+    #[should_panic(expected = "Phi4MM requires per-row ragged decode")]
+    fn phi4mm_rejects_uniform_batched_decode() {
+        let _ = emit_decode_batched(&phi4_like(), 4, false);
+    }
+
     fn mrope_qwen(layout: MropeLayout) -> Config {
         Config {
             context_capacity: 8,
diff --git a/src/lib/mlxcel-xla/src/emitter/model.rs b/src/lib/mlxcel-xla/src/emitter/model.rs
index d76e3fa23..16a748cc6 100644
--- a/src/lib/mlxcel-xla/src/emitter/model.rs
+++ b/src/lib/mlxcel-xla/src/emitter/model.rs
@@ -252,6 +252,29 @@ struct LayerW {
     /// MoE FFN weight handles (issue #500), `Some` on a MoE layer (router, stacked
     /// experts, optional shared expert); `None` on a dense layer.
     moe: Option<MoeLayerW>,
+    /// Phi4MM's immutable speech/vision adapter pairs. The request's numeric
+    /// mode gates their deltas per row; the base weights are never mutated.
+    phi4_lora: Option<Phi4LayerLoraW>,
+}
+
+#[derive(Clone)]
+struct LoraPairW {
+    a: Val,
+    b: Val,
+    scale: f32,
+}
+
+#[derive(Clone)]
+struct ModalLoraW {
+    speech: LoraPairW,
+    vision: LoraPairW,
+}
+
+struct Phi4LayerLoraW {
+    qkv: ModalLoraW,
+    o: ModalLoraW,
+    gate_up: ModalLoraW,
+    down: ModalLoraW,
 }
 
 impl LayerW {
@@ -272,6 +295,7 @@ struct Args {
     /// Untied LM head (`None` when tied; the tail then reuses `embed`).
     lm_head: Option<Val>,
     layers: Vec<LayerW>,
+    adapter_mode: Option<Val>,
     token: Val,
     pos: Val,
     cache_len: Val,
@@ -501,6 +525,63 @@ fn take_layer_weights(
     // appended after the attention weights / biases / FF norms on a MoE layer, in the
     // same order `weight_specs` (`weights.rs`) lists them so the args line up.
     let moe = moe_layer.then(|| take_moe_weights(decls, idx, c, li));
+    let phi4_lora = c.phi4_lora.map(|lora| {
+        let pair = |decls: &mut Vec<ArgDecl>,
+                    idx: &mut usize,
+                    stem: &str,
+                    input: usize,
+                    output: usize,
+                    rank: usize,
+                    scale: f32,
+                    adapter: &str| {
+            LoraPairW {
+                a: take_arg(
+                    decls,
+                    idx,
+                    Ty::f32(vec![rank, input]),
+                    format!("params['layers'][{li}]['{stem}.lora_A.{adapter}']"),
+                ),
+                b: take_arg(
+                    decls,
+                    idx,
+                    Ty::f32(vec![output, rank]),
+                    format!("params['layers'][{li}]['{stem}.lora_B.{adapter}']"),
+                ),
+                scale,
+            }
+        };
+        let modal =
+            |decls: &mut Vec<ArgDecl>, idx: &mut usize, stem: &str, input: usize, output: usize| {
+                ModalLoraW {
+                    speech: pair(
+                        decls,
+                        idx,
+                        stem,
+                        input,
+                        output,
+                        lora.speech_rank,
+                        lora.speech_scale,
+                        "speech",
+                    ),
+                    vision: pair(
+                        decls,
+                        idx,
+                        stem,
+                        input,
+                        output,
+                        lora.vision_rank,
+                        lora.vision_scale,
+                        "vision",
+                    ),
+                }
+            };
+        Phi4LayerLoraW {
+            qkv: modal(decls, idx, "qkv", h, qd + 2 * kv),
+            o: modal(decls, idx, "o", qd, h),
+            gate_up: modal(decls, idx, "gate_up", h, 2 * inter),
+            down: modal(decls, idx, "down", inter, h),
+        }
+    });
     LayerW {
         down,
         gate,
@@ -525,6 +606,7 @@ fn take_layer_weights(
         gate_bias,
         up_bias,
         moe,
+        phi4_lora,
     }
 }
 
@@ -626,6 +708,14 @@ fn build_arg_schema(b: &mut Builder, c: &Config) -> (Vec<ArgDecl>, Args) {
         layers.push(take_layer_weights(b, &mut decls, &mut idx, c, li));
     }
 
+    let adapter_mode = c.phi4_lora.map(|_| {
+        take_arg(
+            &mut decls,
+            &mut idx,
+            Ty::scalar("i32"),
+            "adapter_mode".into(),
+        )
+    });
     let token = take_arg(&mut decls, &mut idx, Ty::scalar("i32"), "token".into());
     let pos = take_arg(
         &mut decls,
@@ -655,6 +745,7 @@ fn build_arg_schema(b: &mut Builder, c: &Config) -> (Vec<ArgDecl>, Args) {
             final_norm_bias,
             lm_head,
             layers,
+            adapter_mode,
             token,
             pos,
             cache_len,
@@ -1065,8 +1156,36 @@ fn emit_mlp_body(
         return layout.add_bias(b, down, lw.down_bias.as_ref());
     }
     let gate = layout.linear(b, hn, lw.gate.as_ref().expect("gated MLP has a gate"));
+    let gate_up_lora = lw
+        .phi4_lora
+        .as_ref()
+        .map(|lora| layout.modal_lora_delta(b, hn, &lora.gate_up));
+    let gate = match &gate_up_lora {
+        Some(delta) if delta.ty.shape.len() == 1 => {
+            let delta = b.slice(delta, &[(0, c.inter)]);
+            b.add(&gate, &delta)
+        }
+        Some(delta) => {
+            let rows = delta.ty.shape[0];
+            let delta = b.slice(delta, &[(0, rows), (0, c.inter)]);
+            b.add(&gate, &delta)
+        }
+        None => gate,
+    };
     let gate = layout.add_bias(b, gate, lw.gate_bias.as_ref());
     let up = layout.linear(b, hn, LayerW::dense(&lw.up));
+    let up = match &gate_up_lora {
+        Some(delta) if delta.ty.shape.len() == 1 => {
+            let delta = b.slice(delta, &[(c.inter, 2 * c.inter)]);
+            b.add(&up, &delta)
+        }
+        Some(delta) => {
+            let rows = delta.ty.shape[0];
+            let delta = b.slice(delta, &[(0, rows), (c.inter, 2 * c.inter)]);
+            b.add(&up, &delta)
+        }
+        None => up,
+    };
     let up = layout.add_bias(b, up, lw.up_bias.as_ref());
     let act = if c.mlp_geglu {
         gelu_tanh(b, &gate)
@@ -1075,6 +1194,13 @@ fn emit_mlp_body(
     };
     let act = b.multiply(&act, &up);
     let down = layout.linear(b, &act, LayerW::dense(&lw.down));
+    let down = match &lw.phi4_lora {
+        Some(lora) => {
+            let delta = layout.modal_lora_delta(b, &act, &lora.down);
+            b.add(&down, &delta)
+        }
+        None => down,
+    };
     let down = layout.add_bias(b, down, lw.down_bias.as_ref());
     if c.has_post_ff_norm() {
         let w = norm_w(
@@ -1249,6 +1375,7 @@ enum AttnLayout {
     /// `[d]` RoPE vector, an `[S]` key mask, and a shared-offset KV write at
     /// `cache_len`.
     Single {
+        adapter_mode: Option<Val>,
         cos: Val,
         sin: Val,
         /// Gemma3 / OLMo3 local-base RoPE vectors for the sliding layers (`Some`
@@ -1264,6 +1391,7 @@ enum AttnLayout {
     /// write at each row's own `pos[b]`.
     Ragged {
         bsz: usize,
+        adapter_mode: Option<Val>,
         cos: Val,
         sin: Val,
         cos_local: Option<Val>,
@@ -1278,6 +1406,7 @@ enum AttnLayout {
     /// Scores read the freshly projected K/V directly (no cache read-back).
     Prefill {
         lp: usize,
+        adapter_mode: Option<Val>,
         cos: Val,
         sin: Val,
         cos_local: Option<Val>,
@@ -1288,6 +1417,60 @@ enum AttnLayout {
 }
 
 impl AttnLayout {
+    fn adapter_mode(&self) -> Option<&Val> {
+        match self {
+            Self::Single { adapter_mode, .. }
+            | Self::Ragged { adapter_mode, .. }
+            | Self::Prefill { adapter_mode, .. } => adapter_mode.as_ref(),
+        }
+    }
+
+    fn adapter_predicate(&self, b: &mut Builder, code: i32, output_shape: &[usize]) -> Option<Val> {
+        let mode = self.adapter_mode()?;
+        let code = b.const_i32(code);
+        let code = if mode.ty.shape.is_empty() {
+            code
+        } else {
+            b.broadcast(&code, &[], mode.ty.shape.clone())
+        };
+        let predicate = b.compare("EQ", mode, &code, "SIGNED");
+        Some(if predicate.ty.shape == output_shape {
+            predicate
+        } else if predicate.ty.shape.is_empty() {
+            b.broadcast(&predicate, &[], output_shape.to_vec())
+        } else {
+            b.broadcast(&predicate, &[0], output_shape.to_vec())
+        })
+    }
+
+    fn lora_branch(&self, b: &mut Builder, x: &Val, pair: &LoraPairW) -> Val {
+        let low_rank = self.linear(b, x, &pair.a);
+        let delta = self.linear(b, &low_rank, &pair.b);
+        let scale = b.const_f32(pair.scale);
+        let scale = b.broadcast(&scale, &[], delta.ty.shape.clone());
+        b.multiply(&delta, &scale)
+    }
+
+    /// Compute a Phi4MM adapter delta without mutating resident base weights.
+    /// Mode 0 is language/base-only, 1 speech, and 2 vision (including the
+    /// official mixed vision-speech policy).
+    fn modal_lora_delta(&self, b: &mut Builder, x: &Val, lora: &ModalLoraW) -> Val {
+        let speech = self.lora_branch(b, x, &lora.speech);
+        let vision = self.lora_branch(b, x, &lora.vision);
+        let shape = speech.ty.shape.clone();
+        let zero = b.const_f32(0.0);
+        let zeros = b.broadcast(&zero, &[], shape.clone());
+        let speech_predicate = self
+            .adapter_predicate(b, 1, &shape)
+            .expect("Phi4MM LoRA requires an adapter-mode graph input");
+        let vision_predicate = self
+            .adapter_predicate(b, 2, &shape)
+            .expect("Phi4MM LoRA requires an adapter-mode graph input");
+        let speech = b.select(&speech_predicate, &speech, &zeros);
+        let vision = b.select(&vision_predicate, &vision, &zeros);
+        b.add(&speech, &vision)
+    }
+
     /// Normalization at this kind's activation rank (rank-reduced for single
     /// decode, per-row otherwise): RMSNorm for the Llama family, mean-subtract
     /// LayerNorm with the optional `bias` for a LayerNorm arch (issue #498).
@@ -1343,24 +1526,50 @@ impl AttnLayout {
         let (nq, nkv) = (c.n_q, c.n_kv);
         match self {
             AttnLayout::Single { .. } => {
+                let lora = lw
+                    .phi4_lora
+                    .as_ref()
+                    .map(|lora| self.modal_lora_delta(b, hn, &lora.qkv));
                 let q = b.linear(hn, &lw.wq);
+                let q = match &lora {
+                    Some(delta) => {
+                        let delta = b.slice(delta, &[(0, nq * d)]);
+                        b.add(&q, &delta)
+                    }
+                    None => q,
+                };
                 let q = add_proj_bias(b, q, &lw.bq);
                 let q = b.reshape(&q, vec![nq, d]);
                 let kk = b.linear(hn, &lw.wk);
+                let kk = match &lora {
+                    Some(delta) => {
+                        let delta = b.slice(delta, &[(nq * d, nq * d + nkv * d)]);
+                        b.add(&kk, &delta)
+                    }
+                    None => kk,
+                };
                 let kk = add_proj_bias(b, kk, &lw.bk);
                 let kk = b.reshape(&kk, vec![nkv, d]);
                 let vv = b.linear(hn, &lw.wv);
+                let vv = match &lora {
+                    Some(delta) => {
+                        let delta = b.slice(delta, &[(nq * d + nkv * d, nq * d + 2 * nkv * d)]);
+                        b.add(&vv, &delta)
+                    }
+                    None => vv,
+                };
                 let vv = add_proj_bias(b, vv, &lw.bv);
                 let vv = b.reshape(&vv, vec![nkv, d]);
                 (q, kk, vv)
             }
-            AttnLayout::Ragged { bsz, .. } => Self::project_qkv_seq(b, c, hn, lw, *bsz),
-            AttnLayout::Prefill { lp, .. } => Self::project_qkv_seq(b, c, hn, lw, *lp),
+            AttnLayout::Ragged { bsz, .. } => self.project_qkv_seq(b, c, hn, lw, *bsz),
+            AttnLayout::Prefill { lp, .. } => self.project_qkv_seq(b, c, hn, lw, *lp),
         }
     }
 
     /// The `[N, ...]` (seq) q/k/v projection shared by ragged decode and prefill.
     fn project_qkv_seq(
+        &self,
         b: &mut Builder,
         c: &Config,
         hn: &Val,
@@ -1369,13 +1578,38 @@ impl AttnLayout {
     ) -> (Val, Val, Val) {
         let d = c.head_dim;
         let (nq, nkv) = (c.n_q, c.n_kv);
+        let lora = lw
+            .phi4_lora
+            .as_ref()
+            .map(|lora| self.modal_lora_delta(b, hn, &lora.qkv));
         let q = b.linear_seq(hn, &lw.wq);
+        let q = match &lora {
+            Some(delta) => {
+                let delta = b.slice(delta, &[(0, n), (0, nq * d)]);
+                b.add(&q, &delta)
+            }
+            None => q,
+        };
         let q = add_proj_bias_seq(b, q, &lw.bq, n, nq * d);
         let q = b.reshape(&q, vec![n, nq, d]);
         let kk = b.linear_seq(hn, &lw.wk);
+        let kk = match &lora {
+            Some(delta) => {
+                let delta = b.slice(delta, &[(0, n), (nq * d, nq * d + nkv * d)]);
+                b.add(&kk, &delta)
+            }
+            None => kk,
+        };
         let kk = add_proj_bias_seq(b, kk, &lw.bk, n, nkv * d);
         let kk = b.reshape(&kk, vec![n, nkv, d]);
         let vv = b.linear_seq(hn, &lw.wv);
+        let vv = match &lora {
+            Some(delta) => {
+                let delta = b.slice(delta, &[(0, n), (nq * d + nkv * d, nq * d + 2 * nkv * d)]);
+                b.add(&vv, &delta)
+            }
+            None => vv,
+        };
         let vv = add_proj_bias_seq(b, vv, &lw.bv, n, nkv * d);
         let vv = b.reshape(&vv, vec![n, nkv, d]);
         (q, kk, vv)
@@ -1645,6 +1879,13 @@ impl AttnLayout {
             AttnLayout::Single { .. } => b.linear(o, &lw.wo),
             _ => b.linear_seq(o, &lw.wo),
         };
+        let out = match &lw.phi4_lora {
+            Some(lora) => {
+                let delta = self.modal_lora_delta(b, o, &lora.o);
+                b.add(&out, &delta)
+            }
+            None => out,
+        };
         self.add_bias(b, out, lw.wo_bias.as_ref())
     }
 }
@@ -1903,6 +2144,7 @@ pub fn emit_decode_with(c: &Config, sample: bool, precision: Precision) -> Strin
     });
 
     let layout = AttnLayout::Single {
+        adapter_mode: a.adapter_mode.clone(),
         cos: cos_vec,
         sin: sin_vec,
         cos_local,
@@ -1970,6 +2212,7 @@ struct BatchedArgs {
     final_norm_bias: Option<Val>,
     lm_head: Option<Val>,
     layers: Vec<LayerW>,
+    adapter_mode: Option<Val>,
     token: Val,     // [B] i32
     pos: Val,       // scalar i32 (shared across the batch)
     cache_len: Val, // scalar i32 (shared across the batch)
@@ -2008,6 +2251,14 @@ fn build_batched_arg_schema(
         layers.push(take_layer_weights(b, &mut decls, &mut idx, c, li));
     }
 
+    let adapter_mode = c.phi4_lora.map(|_| {
+        take_arg(
+            &mut decls,
+            &mut idx,
+            Ty::new(vec![bsz], "i32"),
+            "adapter_mode".into(),
+        )
+    });
     let token = take_arg(
         &mut decls,
         &mut idx,
@@ -2054,6 +2305,7 @@ fn build_batched_arg_schema(
             final_norm_bias,
             lm_head,
             layers,
+            adapter_mode,
             token,
             pos,
             cache_len,
@@ -2099,6 +2351,10 @@ pub fn emit_decode_batched_with(
     sample: bool,
     precision: Precision,
 ) -> String {
+    assert!(
+        c.phi4_lora.is_none(),
+        "Phi4MM requires per-row ragged decode; uniform batched decode cannot carry request-scoped adapter modes"
+    );
     let mut b = Builder::new().with_precision(precision);
     let (decls, a) = build_batched_arg_schema(&mut b, c, bsz);
     let k = emit_consts(&mut b, c);
@@ -2318,6 +2574,7 @@ struct RaggedArgs {
     final_norm_bias: Option<Val>,
     lm_head: Option<Val>,
     layers: Vec<LayerW>,
+    adapter_mode: Option<Val>,
     token: Val,     // [B] i32
     pos: Val,       // [B] i32 (per row)
     cache_len: Val, // [B] i32 (per row)
@@ -2352,6 +2609,14 @@ fn build_ragged_arg_schema(b: &mut Builder, c: &Config, bsz: usize) -> (Vec<ArgD
         layers.push(take_layer_weights(b, &mut decls, &mut idx, c, li));
     }
 
+    let adapter_mode = c.phi4_lora.map(|_| {
+        take_arg(
+            &mut decls,
+            &mut idx,
+            Ty::new(vec![bsz], "i32"),
+            "adapter_mode".into(),
+        )
+    });
     let token = take_arg(
         &mut decls,
         &mut idx,
@@ -2410,6 +2675,7 @@ fn build_ragged_arg_schema(b: &mut Builder, c: &Config, bsz: usize) -> (Vec<ArgD
             final_norm_bias,
             lm_head,
             layers,
+            adapter_mode,
             token,
             pos,
             cache_len,
@@ -2485,6 +2751,7 @@ pub fn emit_decode_ragged_with(
 
     let layout = AttnLayout::Ragged {
         bsz,
+        adapter_mode: a.adapter_mode.clone(),
         cos,
         sin,
         cos_local,
@@ -2597,6 +2864,7 @@ struct PrefillArgs {
     final_norm_bias: Option<Val>,
     lm_head: Option<Val>,
     layers: Vec<LayerW>,
+    adapter_mode: Option<Val>,
     input: PrefillInput,
     positions: Val,
     real_len: Val,
@@ -2634,6 +2902,14 @@ fn build_prefill_arg_schema(
         layers.push(take_layer_weights(b, &mut decls, &mut idx, c, li));
     }
 
+    let adapter_mode = c.phi4_lora.map(|_| {
+        take_arg(
+            &mut decls,
+            &mut idx,
+            Ty::scalar("i32"),
+            "adapter_mode".into(),
+        )
+    });
     let (tokens, embeddings) = match input_kind {
         PrefillInputKind::Tokens => (
             Some(take_arg(
@@ -2737,6 +3013,7 @@ fn build_prefill_arg_schema(
             final_norm_bias,
             lm_head,
             layers,
+            adapter_mode,
             input,
             positions,
             real_len,
@@ -3144,6 +3421,7 @@ fn emit_prefill_module(
 
     let layout = AttnLayout::Prefill {
         lp,
+        adapter_mode: a.adapter_mode.clone(),
         cos,
         sin,
         cos_local,
diff --git a/src/lib/mlxcel-xla/src/emitter/phi4_audio.rs b/src/lib/mlxcel-xla/src/emitter/phi4_audio.rs
new file mode 100644
index 000000000..38b2587fc
--- /dev/null
+++ b/src/lib/mlxcel-xla/src/emitter/phi4_audio.rs
@@ -0,0 +1,1469 @@
+// Copyright 2025-2026 Lablup Inc. and Jeongkyu Shin
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+//! Phi4MM's pinned SpeechLib/Conformer audio artifact contract.
+//!
+//! The audio encoder is compiled as a module separate from the decoder bundle:
+//! it owns only audio encoder/projection weights and returns projected decoder
+//! embeddings plus their exact valid length. Keeping this schema separate stops
+//! an audio-capable artifact from silently loading a decoder-only or partial
+//! checkpoint.
+
+use std::collections::{HashMap, HashSet};
+
+use serde_json::Value;
+
+use super::builder::{Builder, Precision, Ty, Val};
+
+const RAW_PREFIX: &str = "model.embed_tokens_extend.audio_embed";
+
+#[derive(Clone, Debug, Eq, PartialEq)]
+pub(crate) struct Phi4AudioConfig {
+    pub input_size: usize,
+    pub attention_dim: usize,
+    pub attention_heads: usize,
+    pub num_blocks: usize,
+    pub linear_units: usize,
+    pub time_reduction: usize,
+    pub conv_channels: usize,
+    pub kernel_size: usize,
+    pub relative_bias_max_distance: usize,
+    pub projection_hidden: usize,
+}
+
+impl Phi4AudioConfig {
+    pub(crate) fn from_json_str(text: &str) -> Result<Self, String> {
+        let root: Value =
+            serde_json::from_str(text).map_err(|error| format!("parse config.json: {error}"))?;
+        if root.get("model_type").and_then(Value::as_str) != Some("phi4mm") {
+            return Err("Phi4MM audio requires config.json model_type = \"phi4mm\"".to_string());
+        }
+        let config = root
+            .pointer("/audio_processor/config")
+            .ok_or("Phi4MM config is missing audio_processor.config")?;
+        let integer = |name: &str| {
+            config
+                .get(name)
+                .and_then(Value::as_u64)
+                .and_then(|value| usize::try_from(value).ok())
+                .filter(|value| *value > 0)
+                .ok_or_else(|| format!("Phi4MM audio config requires positive integer `{name}`"))
+        };
+        let nested_integer = |path: &str, label: &str| {
+            config
+                .pointer(path)
+                .and_then(Value::as_u64)
+                .and_then(|value| usize::try_from(value).ok())
+                .filter(|value| *value > 0)
+                .ok_or_else(|| format!("Phi4MM audio config requires positive integer `{label}`"))
+        };
+        let parsed = Self {
+            input_size: integer("input_size")?,
+            attention_dim: integer("attention_dim")?,
+            attention_heads: integer("attention_heads")?,
+            num_blocks: integer("num_blocks")?,
+            linear_units: integer("linear_units")?,
+            time_reduction: integer("time_reduction")?,
+            conv_channels: nested_integer(
+                "/nemo_conv_settings/conv_channels",
+                "nemo_conv_settings.conv_channels",
+            )?,
+            kernel_size: integer("kernel_size")?,
+            relative_bias_max_distance: nested_integer(
+                "/relative_attention_bias_args/t5_bias_max_distance",
+                "relative_attention_bias_args.t5_bias_max_distance",
+            )?,
+            projection_hidden: root
+                .get("hidden_size")
+                .and_then(Value::as_u64)
+                .and_then(|value| usize::try_from(value).ok())
+                .filter(|value| *value > 0)
+                .ok_or("Phi4MM config requires positive integer hidden_size")?,
+        };
+        parsed.validate_published(config, &root)?;
+        Ok(parsed)
+    }
+
+    fn validate_published(&self, config: &Value, root: &Value) -> Result<(), String> {
+        let string_is =
+            |name: &str, expected: &str| config.get(name).and_then(Value::as_str) == Some(expected);
+        let bool_is = |name: &str, expected: bool| {
+            config.get(name).and_then(Value::as_bool) == Some(expected)
+        };
+        let published = self.input_size == 80
+            && self.attention_dim == 1024
+            && self.attention_heads == 16
+            && self.num_blocks == 24
+            && self.linear_units == 1536
+            && self.time_reduction == 8
+            && self.conv_channels == 1024
+            && self.kernel_size == 3
+            && self.relative_bias_max_distance == 500
+            && self.projection_hidden == 3072
+            && string_is("input_layer", "nemo_conv")
+            && string_is("activation", "swish")
+            && string_is("conv_activation", "swish")
+            && string_is("conv_glu_type", "swish")
+            && bool_is("causal", true)
+            && bool_is("batch_norm", false)
+            && bool_is("bias_in_glu", true)
+            && config.get("depthwise_multiplier").and_then(Value::as_u64) == Some(1)
+            && config
+                .get("depthwise_seperable_out_channel")
+                .and_then(Value::as_u64)
+                == Some(1024)
+            && config
+                .pointer("/encoder_embedding_config/input_size")
+                .and_then(Value::as_u64)
+                == Some(80)
+            && config
+                .pointer("/relative_attention_bias_args/type")
+                .and_then(Value::as_str)
+                == Some("t5")
+            && !config
+                .pointer("/relative_attention_bias_args/t5_bias_symmetric")
+                .and_then(Value::as_bool)
+                .unwrap_or(false)
+            && !config
+                .pointer("/nemo_conv_settings/is_causal")
+                .and_then(Value::as_bool)
+                .unwrap_or(false);
+        let projection = root.pointer("/embd_layer/audio_embd_layer");
+        let published_projection = projection.is_some_and(|value| {
+            value.get("compression_rate").and_then(Value::as_u64) == Some(8)
+                && value.get("downsample_rate").and_then(Value::as_u64) == Some(1)
+                && value.get("projection_cls").and_then(Value::as_str) == Some("mlp")
+                && value.get("use_conv_downsample").and_then(Value::as_bool) == Some(false)
+                && value.get("use_qformer").and_then(Value::as_bool) == Some(false)
+        });
+        if !published || !published_projection {
+            return Err(
+                "unsupported Phi4MM audio architecture: expected the pinned published Cascades/MLP contract"
+                    .to_string(),
+            );
+        }
+        Ok(())
+    }
+
+    pub(crate) fn encoded_bucket_len(&self, frame_bucket: usize) -> Result<usize, String> {
+        if frame_bucket == 0 {
+            return Err("Phi4MM audio frame bucket must be positive".to_string());
+        }
+        if self.time_reduction != 8 {
+            return Err(format!(
+                "Phi4MM audio graph implements the pinned 3x stride-2 reduction, not time_reduction={}",
+                self.time_reduction
+            ));
+        }
+        // The pinned NeMo subsampler is three symmetric-pad Conv2D stages
+        // (kernel=3, pad=1, stride=2). Derive the emitted row count from that
+        // exact operator chain instead of assuming the configured aggregate
+        // reduction remains equivalent for every boundary length.
+        Ok((0..3).fold(frame_bucket, |length, _| (length - 1) / 2 + 1))
+    }
+
+    pub(crate) fn encoded_valid_len(&self, frame_len: usize) -> Result<usize, String> {
+        self.encoded_bucket_len(frame_len)
+    }
+}
+
+#[derive(Clone, Debug, Eq, PartialEq)]
+pub(crate) struct Phi4AudioWeightSpec {
+    pub name: String,
+    pub shape: Vec<usize>,
+}
+
+fn push(specs: &mut Vec<Phi4AudioWeightSpec>, name: String, shape: &[usize]) {
+    specs.push(Phi4AudioWeightSpec {
+        name,
+        shape: shape.to_vec(),
+    });
+}
+
+fn push_linear(specs: &mut Vec<Phi4AudioWeightSpec>, stem: &str, output: usize, input: usize) {
+    push(specs, format!("{stem}.weight"), &[output, input]);
+    push(specs, format!("{stem}.bias"), &[output]);
+}
+
+fn push_layer_norm(specs: &mut Vec<Phi4AudioWeightSpec>, stem: &str, dim: usize) {
+    push(specs, format!("{stem}.weight"), &[dim]);
+    push(specs, format!("{stem}.bias"), &[dim]);
+}
+
+fn push_ffn(specs: &mut Vec<Phi4AudioWeightSpec>, stem: &str, dim: usize, inner: usize) {
+    push_layer_norm(specs, &format!("{stem}.layer_norm"), dim);
+    push_linear(specs, &format!("{stem}.net.0.linear"), inner * 2, dim);
+    push_linear(specs, &format!("{stem}.net.2"), dim, inner);
+}
+
+pub(crate) fn phi4_audio_weight_specs(config: &Phi4AudioConfig) -> Vec<Phi4AudioWeightSpec> {
+    let dim = config.attention_dim;
+    let channels = config.conv_channels;
+    let mut specs = Vec::with_capacity(887);
+    let encoder = format!("{RAW_PREFIX}.encoder");
+
+    push(
+        &mut specs,
+        format!("{encoder}.encoder_embedding.global_mean"),
+        &[config.input_size],
+    );
+    push(
+        &mut specs,
+        format!("{encoder}.encoder_embedding.global_invstd"),
+        &[config.input_size],
+    );
+    let embed = format!("{encoder}.embed");
+    push(
+        &mut specs,
+        format!("{embed}.conv.0.weight"),
+        &[channels, 1, 3, 3],
+    );
+    push(&mut specs, format!("{embed}.conv.0.bias"), &[channels]);
+    for (depthwise, pointwise) in [(2, 3), (5, 6)] {
+        push(
+            &mut specs,
+            format!("{embed}.conv.{depthwise}.weight"),
+            &[channels, 1, 3, 3],
+        );
+        push(
+            &mut specs,
+            format!("{embed}.conv.{depthwise}.bias"),
+            &[channels],
+        );
+        push(
+            &mut specs,
+            format!("{embed}.conv.{pointwise}.weight"),
+            &[channels, channels, 1, 1],
+        );
+        push(
+            &mut specs,
+            format!("{embed}.conv.{pointwise}.bias"),
+            &[channels],
+        );
+    }
+    let reduced_frequency = config.input_size.div_ceil(8);
+    push_linear(
+        &mut specs,
+        &format!("{embed}.out"),
+        dim,
+        channels * reduced_frequency,
+    );
+
+    for layer in 0..config.num_blocks {
+        let block = format!("{encoder}.encoders.{layer}");
+        push_ffn(
+            &mut specs,
+            &format!("{block}.feed_forward_in"),
+            dim,
+            config.linear_units,
+        );
+        push_layer_norm(&mut specs, &format!("{block}.layer_norm_att"), dim);
+        for projection in ["linear_q", "linear_k", "linear_v", "linear_out"] {
+            push_linear(
+                &mut specs,
+                &format!("{block}.self_attn.{projection}"),
+                dim,
+                dim,
+            );
+        }
+
+        let conv = format!("{block}.conv");
+        push_layer_norm(&mut specs, &format!("{conv}.layer_norm"), dim);
+        push(&mut specs, format!("{conv}.glu.b1"), &[1, dim, 1]);
+        push(&mut specs, format!("{conv}.glu.b2"), &[1, dim, 1]);
+        push(
+            &mut specs,
+            format!("{conv}.glu.ext_pw_conv_1d.weight"),
+            &[dim * 2, dim, 1],
+        );
+        push(
+            &mut specs,
+            format!("{conv}.glu.ext_pw_conv_1d.bias"),
+            &[dim * 2],
+        );
+        push(
+            &mut specs,
+            format!("{conv}.dw_sep_conv_1d.dw_conv.weight"),
+            &[dim, 1, config.kernel_size],
+        );
+        push(
+            &mut specs,
+            format!("{conv}.dw_sep_conv_1d.dw_conv.bias"),
+            &[dim],
+        );
+        for projection in ["dw_sep_conv_1d.pw_conv", "ext_pw_conv_1d"] {
+            push(
+                &mut specs,
+                format!("{conv}.{projection}.weight"),
+                &[dim, dim, 1],
+            );
+            push(&mut specs, format!("{conv}.{projection}.bias"), &[dim]);
+        }
+        push_ffn(
+            &mut specs,
+            &format!("{block}.feed_forward_out"),
+            dim,
+            config.linear_units,
+        );
+        push_layer_norm(&mut specs, &format!("{block}.layer_norm"), dim);
+    }
+    push(
+        &mut specs,
+        format!("{encoder}.relative_attention_bias_layer.bias_values.weight"),
+        &[
+            config.relative_bias_max_distance * 2,
+            config.attention_heads,
+        ],
+    );
+
+    for mode in ["speech", "vision"] {
+        let projection = format!("{RAW_PREFIX}.audio_projection.{mode}");
+        push_linear(
+            &mut specs,
+            &format!("{projection}.0"),
+            config.projection_hidden,
+            dim,
+        );
+        push_linear(
+            &mut specs,
+            &format!("{projection}.2"),
+            config.projection_hidden,
+            config.projection_hidden,
+        );
+    }
+    specs
+}
+
+pub(crate) fn validate_phi4_audio_weight_shapes<F>(
+    config: &Phi4AudioConfig,
+    mut actual_shape: F,
+) -> Result<(), String>
+where
+    F: FnMut(&str) -> Option<Vec<usize>>,
+{
+    let specs = phi4_audio_weight_specs(config);
+    let unique = specs
+        .iter()
+        .map(|spec| spec.name.as_str())
+        .collect::<HashSet<_>>();
+    if unique.len() != specs.len() {
+        return Err("Phi4MM audio weight schema contains duplicate tensor names".to_string());
+    }
+    for spec in specs {
+        let actual = actual_shape(&spec.name)
+            .ok_or_else(|| format!("Phi4MM audio checkpoint is missing `{}`", spec.name))?;
+        if actual != spec.shape {
+            return Err(format!(
+                "Phi4MM audio weight `{}` has shape {actual:?}, expected {:?}",
+                spec.name, spec.shape
+            ));
+        }
+    }
+    Ok(())
+}
+
+fn lookup<'a>(weights: &'a HashMap<String, Val>, name: &str) -> &'a Val {
+    weights
+        .get(name)
+        .unwrap_or_else(|| panic!("Phi4MM audio emitter schema is missing {name}"))
+}
+
+fn add_last_bias(b: &mut Builder, value: &Val, bias: &Val) -> Val {
+    let axis = value.ty.shape.len() - 1;
+    let broadcast = b.broadcast(bias, &[axis], value.ty.shape.clone());
+    b.add(value, &broadcast)
+}
+
+fn linear_bias(b: &mut Builder, weights: &HashMap<String, Val>, input: &Val, stem: &str) -> Val {
+    let projected = b.linear_seq(input, lookup(weights, &format!("{stem}.weight")));
+    add_last_bias(b, &projected, lookup(weights, &format!("{stem}.bias")))
+}
+
+fn layer_norm(b: &mut Builder, weights: &HashMap<String, Val>, input: &Val, stem: &str) -> Val {
+    let axis = input.ty.shape.len() - 1;
+    let dim = input.ty.shape[axis];
+    let zero = b.const_f32(0.0);
+    let sum = b.reduce_add(input, axis, &zero);
+    let count = b.const_f32(dim as f32);
+    let count = b.broadcast(&count, &[], sum.ty.shape.clone());
+    let mean = b.divide(&sum, &count);
+    let leading = (0..axis).collect::<Vec<_>>();
+    let mean = b.broadcast(&mean, &leading, input.ty.shape.clone());
+    let centered = b.subtract(input, &mean);
+    let squared = b.multiply(&centered, &centered);
+    let variance = b.reduce_add(&squared, axis, &zero);
+    let variance = b.divide(&variance, &count);
+    let epsilon = b.const_f32(1e-5);
+    let epsilon = b.broadcast(&epsilon, &[], variance.ty.shape.clone());
+    let variance = b.add(&variance, &epsilon);
+    let inverse_std = b.rsqrt(&variance);
+    let inverse_std = b.broadcast(&inverse_std, &leading, input.ty.shape.clone());
+    let normalized = b.multiply(&centered, &inverse_std);
+    let weight = b.broadcast(
+        lookup(weights, &format!("{stem}.weight")),
+        &[axis],
+        input.ty.shape.clone(),
+    );
+    let normalized = b.multiply(&normalized, &weight);
+    let bias = b.broadcast(
+        lookup(weights, &format!("{stem}.bias")),
+        &[axis],
+        input.ty.shape.clone(),
+    );
+    b.add(&normalized, &bias)
+}
+
+fn silu(b: &mut Builder, input: &Val) -> Val {
+    let one = b.const_f32(1.0);
+    let one = b.broadcast(&one, &[], input.ty.shape.clone());
+    let negative = b.negate(input);
+    let exponential = b.exponential(&negative);
+    let denominator = b.add(&one, &exponential);
+    b.divide(input, &denominator)
+}
+
+fn gelu(b: &mut Builder, input: &Val) -> Val {
+    let shape = input.ty.shape.clone();
+    let inverse_sqrt_two = b.const_f32(std::f32::consts::FRAC_1_SQRT_2);
+    let inverse_sqrt_two = b.broadcast(&inverse_sqrt_two, &[], shape.clone());
+    let scaled = b.multiply(input, &inverse_sqrt_two);
+    let erf = b.erf(&scaled);
+    let one = b.const_f32(1.0);
+    let one = b.broadcast(&one, &[], shape.clone());
+    let half = b.const_f32(0.5);
+    let half = b.broadcast(&half, &[], shape);
+    let cdf = b.add(&one, &erf);
+    let cdf = b.multiply(&half, &cdf);
+    b.multiply(input, &cdf)
+}
+
+fn softmax_last(b: &mut Builder, input: &Val) -> Val {
+    let axis = input.ty.shape.len() - 1;
+    let negative_infinity = b.const_f32(f32::NEG_INFINITY);
+    let maximum = b.reduce_max(input, axis, &negative_infinity);
+    let leading = (0..axis).collect::<Vec<_>>();
+    let maximum = b.broadcast(&maximum, &leading, input.ty.shape.clone());
+    let shifted = b.subtract(input, &maximum);
+    let exponent = b.exponential(&shifted);
+    let zero = b.const_f32(0.0);
+    let sum = b.reduce_add(&exponent, axis, &zero);
+    let sum = b.broadcast(&sum, &leading, input.ty.shape.clone());
+    b.divide(&exponent, &sum)
+}
+
+fn mask_rows(b: &mut Builder, input: &Val, active: &Val) -> Val {
+    let active = b.broadcast(active, &[0], input.ty.shape.clone());
+    let zero = b.const_f32(0.0);
+    let zero = b.broadcast(&zero, &[], input.ty.shape.clone());
+    b.select(&active, input, &zero)
+}
+
+fn round_bf16(b: &mut Builder, input: &Val) -> Val {
+    let narrow = b.convert(input, "bf16");
+    b.convert(&narrow, "f32")
+}
+
+fn conv2d(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    input: &Val,
+    stem: &str,
+    stride: usize,
+    padding: usize,
+    groups: usize,
+) -> Val {
+    let kernel = b.transpose(lookup(weights, &format!("{stem}.weight")), &[2, 3, 1, 0]);
+    let convolved = b.convolution(
+        input,
+        &kernel,
+        &[stride, stride],
+        &[(padding, padding), (padding, padding)],
+        groups,
+    );
+    add_last_bias(b, &convolved, lookup(weights, &format!("{stem}.bias")))
+}
+
+fn conv1d(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    input: &Val,
+    stem: &str,
+    causal_padding: usize,
+    groups: usize,
+) -> Val {
+    let zero = b.const_f32(0.0);
+    let input = if causal_padding == 0 {
+        input.clone()
+    } else {
+        b.pad(input, &zero, &[0, causal_padding, 0], &[0, 0, 0])
+    };
+    let kernel = b.transpose(lookup(weights, &format!("{stem}.weight")), &[2, 1, 0]);
+    let convolved = b.convolution(&input, &kernel, &[1], &[(0, 0)], groups);
+    add_last_bias(b, &convolved, lookup(weights, &format!("{stem}.bias")))
+}
+
+fn feed_forward(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    input: &Val,
+    stem: &str,
+    inner: usize,
+) -> Val {
+    let normalized = layer_norm(b, weights, input, &format!("{stem}.layer_norm"));
+    let gated = linear_bias(b, weights, &normalized, &format!("{stem}.net.0.linear"));
+    let rows = gated.ty.shape[0];
+    let left = b.slice(&gated, &[(0, rows), (0, inner)]);
+    let gate = b.slice(&gated, &[(0, rows), (inner, inner * 2)]);
+    let activated = silu(b, &gate);
+    let combined = b.multiply(&left, &activated);
+    linear_bias(b, weights, &combined, &format!("{stem}.net.2"))
+}
+
+fn relative_attention_bias(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    config: &Phi4AudioConfig,
+    length: usize,
+) -> Val {
+    let distance = config.relative_bias_max_distance as isize;
+    let mut indices = Vec::with_capacity(length * length);
+    for query in 0..length {
+        for key in 0..length {
+            indices.push(
+                ((key as isize - query as isize).clamp(-distance, distance - 1) + distance) as i32,
+            );
+        }
+    }
+    let indices = b.const_tensor_i32(&indices, vec![length, length, 1]);
+    let gathered = b.gather_rows_nd(
+        lookup(
+            weights,
+            &format!("{RAW_PREFIX}.encoder.relative_attention_bias_layer.bias_values.weight"),
+        ),
+        &indices,
+    );
+    b.transpose(&gathered, &[2, 0, 1])
+}
+
+fn self_attention(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    config: &Phi4AudioConfig,
+    input: &Val,
+    relative_bias: &Val,
+    active_rows: &Val,
+    stem: &str,
+) -> Val {
+    let length = input.ty.shape[0];
+    let heads = config.attention_heads;
+    let head_dim = config.attention_dim / heads;
+    let project = |b: &mut Builder, suffix: &str| {
+        let value = linear_bias(b, weights, input, &format!("{stem}.self_attn.{suffix}"));
+        let value = b.reshape(&value, vec![length, heads, head_dim]);
+        b.transpose(&value, &[1, 0, 2])
+    };
+    let query = project(b, "linear_q");
+    let scale = b.const_f32((head_dim as f32).powf(-0.5));
+    let scale = b.broadcast(&scale, &[], query.ty.shape.clone());
+    let query = b.multiply(&query, &scale);
+    let key = project(b, "linear_k");
+    let value = project(b, "linear_v");
+    let key = b.transpose(&key, &[0, 2, 1]);
+    let scores = b.dot_general(
+        &query,
+        &key,
+        &[0],
+        &[0],
+        &[2],
+        &[1],
+        vec![heads, length, length],
+    );
+    let scores = b.add(&scores, relative_bias);
+    let active_keys = b.broadcast(active_rows, &[2], vec![heads, length, length]);
+    let zero = b.const_f32(0.0);
+    let zero = b.broadcast(&zero, &[], vec![heads, length, length]);
+    let negative = b.const_f32(-1e30);
+    let negative = b.broadcast(&negative, &[], vec![heads, length, length]);
+    let key_mask = b.select(&active_keys, &zero, &negative);
+    let scores = b.add(&scores, &key_mask);
+    let probabilities = softmax_last(b, &scores);
+    let context = b.dot_general(
+        &probabilities,
+        &value,
+        &[0],
+        &[0],
+        &[2],
+        &[1],
+        vec![heads, length, head_dim],
+    );
+    let context = b.transpose(&context, &[1, 0, 2]);
+    let context = b.reshape(&context, vec![length, config.attention_dim]);
+    linear_bias(
+        b,
+        weights,
+        &context,
+        &format!("{stem}.self_attn.linear_out"),
+    )
+}
+
+fn convolution_module(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    config: &Phi4AudioConfig,
+    input: &Val,
+    stem: &str,
+) -> Val {
+    let length = input.ty.shape[0];
+    let dim = config.attention_dim;
+    let normalized = layer_norm(b, weights, input, &format!("{stem}.conv.layer_norm"));
+    let normalized = b.reshape(&normalized, vec![1, length, dim]);
+    let glu_stem = format!("{stem}.conv.glu.ext_pw_conv_1d");
+    let gated = conv1d(b, weights, &normalized, &glu_stem, 0, 1);
+    let left = b.slice(&gated, &[(0, 1), (0, length), (0, dim)]);
+    let gate = b.slice(&gated, &[(0, 1), (0, length), (dim, dim * 2)]);
+    let bias = |b: &mut Builder, suffix: &str| {
+        let value = b.transpose(
+            lookup(weights, &format!("{stem}.conv.glu.{suffix}")),
+            &[0, 2, 1],
+        );
+        b.broadcast(&value, &[0, 1, 2], vec![1, length, dim])
+    };
+    let left_bias = bias(b, "b1");
+    let gate_bias = bias(b, "b2");
+    let left = b.add(&left, &left_bias);
+    let gate = b.add(&gate, &gate_bias);
+    let activated_gate = silu(b, &gate);
+    let gated = b.multiply(&left, &activated_gate);
+    let depthwise = conv1d(
+        b,
+        weights,
+        &gated,
+        &format!("{stem}.conv.dw_sep_conv_1d.dw_conv"),
+        config.kernel_size - 1,
+        dim,
+    );
+    let pointwise = conv1d(
+        b,
+        weights,
+        &depthwise,
+        &format!("{stem}.conv.dw_sep_conv_1d.pw_conv"),
+        0,
+        1,
+    );
+    let activated = silu(b, &pointwise);
+    let output = conv1d(
+        b,
+        weights,
+        &activated,
+        &format!("{stem}.conv.ext_pw_conv_1d"),
+        0,
+        1,
+    );
+    b.reshape(&output, vec![length, dim])
+}
+
+fn conformer_block(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    config: &Phi4AudioConfig,
+    input: &Val,
+    relative_bias: &Val,
+    active_rows: &Val,
+    layer: usize,
+) -> ConformerBlockValues {
+    let stem = format!("{RAW_PREFIX}.encoder.encoders.{layer}");
+    let scale = b.const_f32(0.5);
+    let scale = b.broadcast(&scale, &[], input.ty.shape.clone());
+    let ff = feed_forward(
+        b,
+        weights,
+        input,
+        &format!("{stem}.feed_forward_in"),
+        config.linear_units,
+    );
+    let ff = b.multiply(&ff, &scale);
+    let after_ff_in = b.add(input, &ff);
+    let attention_input = layer_norm(b, weights, &after_ff_in, &format!("{stem}.layer_norm_att"));
+    let attention = self_attention(
+        b,
+        weights,
+        config,
+        &attention_input,
+        relative_bias,
+        active_rows,
+        &stem,
+    );
+    let after_attention = b.add(&after_ff_in, &attention);
+    let convolution = convolution_module(b, weights, config, &after_attention, &stem);
+    let after_convolution = b.add(&after_attention, &convolution);
+    let ff = feed_forward(
+        b,
+        weights,
+        &after_convolution,
+        &format!("{stem}.feed_forward_out"),
+        config.linear_units,
+    );
+    let ff = b.multiply(&ff, &scale);
+    let hidden = b.add(&after_convolution, &ff);
+    let output = layer_norm(b, weights, &hidden, &format!("{stem}.layer_norm"));
+    let output = mask_rows(b, &output, active_rows);
+    ConformerBlockValues {
+        after_ff_in,
+        attention,
+        after_attention,
+        convolution,
+        after_convolution,
+        ff_out: ff,
+        output,
+    }
+}
+
+struct ConformerBlockValues {
+    after_ff_in: Val,
+    attention: Val,
+    after_attention: Val,
+    convolution: Val,
+    after_convolution: Val,
+    ff_out: Val,
+    output: Val,
+}
+
+struct SubsampleValues {
+    conv0: Val,
+    conv1_depthwise: Val,
+    conv1_pointwise: Val,
+    conv1: Val,
+    conv2_depthwise: Val,
+    conv2_pointwise: Val,
+    conv2: Val,
+    projected: Val,
+}
+
+fn subsample(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    config: &Phi4AudioConfig,
+    input: &Val,
+) -> SubsampleValues {
+    let bucket = input.ty.shape[0];
+    let mut hidden = b.reshape(input, vec![1, bucket, config.input_size, 1]);
+    let embed = format!("{RAW_PREFIX}.encoder.embed");
+    hidden = conv2d(b, weights, &hidden, &format!("{embed}.conv.0"), 2, 1, 1);
+    // MLX evaluates the first convolution with BF16 features/weights/bias and
+    // materializes its BF16 result. `relu()`'s F32 zero then promotes the
+    // remaining subsampler and all later captured stages to F32.
+    hidden = round_bf16(b, &hidden);
+    let relu = |b: &mut Builder, value: &Val| {
+        let zero = b.const_f32(0.0);
+        let zero = b.broadcast(&zero, &[], value.ty.shape.clone());
+        b.maximum(value, &zero)
+    };
+    let conv0 = relu(b, &hidden);
+    let conv1_depthwise = conv2d(
+        b,
+        weights,
+        &conv0,
+        &format!("{embed}.conv.2"),
+        2,
+        1,
+        config.conv_channels,
+    );
+    let conv1_pointwise = conv2d(
+        b,
+        weights,
+        &conv1_depthwise,
+        &format!("{embed}.conv.3"),
+        1,
+        0,
+        1,
+    );
+    let conv1 = relu(b, &conv1_pointwise);
+    let conv2_depthwise = conv2d(
+        b,
+        weights,
+        &conv1,
+        &format!("{embed}.conv.5"),
+        2,
+        1,
+        config.conv_channels,
+    );
+    let conv2_pointwise = conv2d(
+        b,
+        weights,
+        &conv2_depthwise,
+        &format!("{embed}.conv.6"),
+        1,
+        0,
+        1,
+    );
+    let conv2 = relu(b, &conv2_pointwise);
+    let hidden = b.transpose(&conv2, &[0, 1, 3, 2]);
+    let shape = hidden.ty.shape.clone();
+    let hidden = b.reshape(&hidden, vec![shape[1], shape[2] * shape[3]]);
+    let projected = linear_bias(b, weights, &hidden, &format!("{embed}.out"));
+    SubsampleValues {
+        conv0,
+        conv1_depthwise,
+        conv1_pointwise,
+        conv1,
+        conv2_depthwise,
+        conv2_pointwise,
+        conv2,
+        projected,
+    }
+}
+
+struct ProjectionValues {
+    speech_first: Val,
+    speech: Val,
+    vision_first: Val,
+    vision: Val,
+    selected: Val,
+}
+
+fn project_audio(
+    b: &mut Builder,
+    weights: &HashMap<String, Val>,
+    config: &Phi4AudioConfig,
+    encoded: &Val,
+    projection_mode: &Val,
+    active_rows: &Val,
+) -> ProjectionValues {
+    let branch = |b: &mut Builder, mode: &str| {
+        let first = linear_bias(
+            b,
+            weights,
+            encoded,
+            &format!("{RAW_PREFIX}.audio_projection.{mode}.0"),
+        );
+        let activated = gelu(b, &first);
+        let projected = linear_bias(
+            b,
+            weights,
+            &activated,
+            &format!("{RAW_PREFIX}.audio_projection.{mode}.2"),
+        );
+        (first, projected)
+    };
+    let (speech_first, speech) = branch(b, "speech");
+    let (vision_first, vision) = branch(b, "vision");
+    let shape = vec![encoded.ty.shape[0], config.projection_hidden];
+    let zero = b.const_f32(0.0);
+    let zero = b.broadcast(&zero, &[], shape.clone());
+    let code = |b: &mut Builder, value: i32| {
+        let value = b.const_i32(value);
+        b.compare("EQ", projection_mode, &value, "SIGNED")
+    };
+    let speech_active = code(b, 1);
+    let vision_active = code(b, 2);
+    let speech_active = b.broadcast(&speech_active, &[], shape.clone());
+    let vision_active = b.broadcast(&vision_active, &[], shape);
+    let selected = b.select(&speech_active, &speech, &zero);
+    let selected = b.select(&vision_active, &vision, &selected);
+    let selected = mask_rows(b, &selected, active_rows);
+    ProjectionValues {
+        speech_first,
+        speech,
+        vision_first,
+        vision,
+        selected,
+    }
+}
+
+/// Emit `audio.main` for a static frame bucket.
+///
+/// Inputs after the 887 resident weights are the exact #874 host-oracle shapes
+/// `features[1,frame_bucket,80]`, `frame_mask[1,frame_bucket]`, the unpadded
+/// frame length, and projection mode (1=speech, 2=vision/mixed). Outputs are
+/// projected decoder embeddings `[1,subsampled_bucket,3072]` and the exact
+/// valid subsampled length.
+pub(crate) fn emit_phi4_audio_with(
+    config: &Phi4AudioConfig,
+    frame_bucket: usize,
+    precision: Precision,
+) -> Result<String, String> {
+    emit_phi4_audio_entry(config, frame_bucket, precision, false)
+}
+
+#[derive(Clone, Debug, Eq, PartialEq)]
+pub(crate) struct Phi4AudioDiagnosticSpec {
+    pub name: &'static str,
+    pub shape: Vec<usize>,
+}
+
+pub(crate) fn phi4_audio_diagnostic_specs(
+    config: &Phi4AudioConfig,
+    frame_bucket: usize,
+) -> Result<Vec<Phi4AudioDiagnosticSpec>, String> {
+    let reduce = |length: usize| (length - 1) / 2 + 1;
+    let time0 = reduce(frame_bucket);
+    let time1 = reduce(time0);
+    let time2 = reduce(time1);
+    let frequency0 = reduce(config.input_size);
+    let frequency1 = reduce(frequency0);
+    let frequency2 = reduce(frequency1);
+    let encoder = vec![1, time2, config.attention_dim];
+    let projection = vec![1, time2, config.projection_hidden];
+    let mut specs = vec![
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.conv0",
+            shape: vec![1, time0, frequency0, config.conv_channels],
+        },
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.conv1.depthwise",
+            shape: vec![1, time1, frequency1, config.conv_channels],
+        },
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.conv1.pointwise",
+            shape: vec![1, time1, frequency1, config.conv_channels],
+        },
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.conv1",
+            shape: vec![1, time1, frequency1, config.conv_channels],
+        },
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.conv2.depthwise",
+            shape: vec![1, time2, frequency2, config.conv_channels],
+        },
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.conv2.pointwise",
+            shape: vec![1, time2, frequency2, config.conv_channels],
+        },
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.conv2",
+            shape: vec![1, time2, frequency2, config.conv_channels],
+        },
+        Phi4AudioDiagnosticSpec {
+            name: "subsample.projected",
+            shape: encoder.clone(),
+        },
+    ];
+    for name in [
+        "block0.after_ff_in",
+        "block0.attention",
+        "block0.after_attention",
+        "block0.convolution",
+        "block0.after_convolution",
+        "block0.ff_out",
+        "block0.output",
+        "block1.output",
+        "block5.output",
+        "block11.output",
+        "block17.output",
+        "block23.output",
+        "encoder.output",
+    ] {
+        specs.push(Phi4AudioDiagnosticSpec {
+            name,
+            shape: encoder.clone(),
+        });
+    }
+    for name in [
+        "projection.speech.first",
+        "projection.speech.output",
+        "projection.vision.first",
+        "projection.vision.output",
+    ] {
+        specs.push(Phi4AudioDiagnosticSpec {
+            name,
+            shape: projection.clone(),
+        });
+    }
+    if time2 != config.encoded_bucket_len(frame_bucket)? {
+        return Err("Phi4MM diagnostic subsampling shape drifted from audio contract".to_string());
+    }
+    Ok(specs)
+}
+
+pub(crate) fn emit_phi4_audio_diagnostic_with(
+    config: &Phi4AudioConfig,
+    frame_bucket: usize,
+    precision: Precision,
+) -> Result<String, String> {
+    emit_phi4_audio_entry(config, frame_bucket, precision, true)
+}
+
+fn emit_phi4_audio_entry(
+    config: &Phi4AudioConfig,
+    frame_bucket: usize,
+    precision: Precision,
+    diagnostic: bool,
+) -> Result<String, String> {
+    if config.attention_dim % config.attention_heads != 0 {
+        return Err("Phi4MM audio attention_dim must be divisible by heads".to_string());
+    }
+    let encoded_len = config.encoded_bucket_len(frame_bucket)?;
+    if encoded_len > config.relative_bias_max_distance {
+        return Err(format!(
+            "Phi4MM audio bucket encodes to {encoded_len} rows, exceeding the pinned 500-row Conformer chunk"
+        ));
+    }
+    let specs = phi4_audio_weight_specs(config);
+    let mut declarations = Vec::with_capacity(specs.len() + 4);
+    let mut weights = HashMap::with_capacity(specs.len());
+    for (index, spec) in specs.iter().enumerate() {
+        let ty = Ty::f32(spec.shape.clone());
+        declarations.push((ty.clone(), spec.name.clone()));
+        weights.insert(spec.name.clone(), Builder::arg(index, ty));
+    }
+    let features_index = declarations.len();
+    declarations.push((
+        Ty::f32(vec![1, frame_bucket, config.input_size]),
+        "features".to_string(),
+    ));
+    let mask_index = declarations.len();
+    declarations.push((
+        Ty::new(vec![1, frame_bucket], "i32"),
+        "frame_mask".to_string(),
+    ));
+    let length_index = declarations.len();
+    declarations.push((Ty::scalar("i32"), "frame_length".to_string()));
+    let mode_index = declarations.len();
+    declarations.push((Ty::scalar("i32"), "projection_mode".to_string()));
+
+    let mut b = Builder::new().with_precision(precision);
+    let features = Builder::arg(
+        features_index,
+        Ty::f32(vec![1, frame_bucket, config.input_size]),
+    );
+    let frame_mask = Builder::arg(mask_index, Ty::new(vec![1, frame_bucket], "i32"));
+    let frame_length = Builder::arg(length_index, Ty::scalar("i32"));
+    let projection_mode = Builder::arg(mode_index, Ty::scalar("i32"));
+
+    let features = b.reshape(&features, vec![frame_bucket, config.input_size]);
+    let frame_mask = b.reshape(&frame_mask, vec![frame_bucket]);
+    let zero_i32 = b.const_i32(0);
+    let zero_i32 = b.broadcast(&zero_i32, &[], vec![frame_bucket]);
+    let frame_active = b.compare("GT", &frame_mask, &zero_i32, "SIGNED");
+    let mean = b.broadcast(
+        lookup(
+            &weights,
+            &format!("{RAW_PREFIX}.encoder.encoder_embedding.global_mean"),
+        ),
+        &[1],
+        vec![frame_bucket, config.input_size],
+    );
+    let inverse_std = b.broadcast(
+        lookup(
+            &weights,
+            &format!("{RAW_PREFIX}.encoder.encoder_embedding.global_invstd"),
+        ),
+        &[1],
+        vec![frame_bucket, config.input_size],
+    );
+    // The MLX oracle casts features to the BF16 mean tensor. Its subtract and
+    // multiply therefore each materialize a BF16 boundary; the first ReLU then
+    // promotes the Conv2D stream and every captured Conformer stage to F32.
+    // Mirror those actual boundaries instead of demoting every contraction.
+    let centered = b.subtract(&features, &mean);
+    let centered = round_bf16(&mut b, &centered);
+    let normalized = b.multiply(&centered, &inverse_std);
+    let normalized = round_bf16(&mut b, &normalized);
+    // Host bucket padding is zero in transport space. Re-apply the frame mask
+    // after normalization so padded rows become the encoder's mathematical
+    // zero padding rather than `(0 - global_mean) * global_invstd`.
+    let normalized = mask_rows(&mut b, &normalized, &frame_active);
+    let subsampled = subsample(&mut b, &weights, config, &normalized);
+
+    let reduction_minus_one = b.const_i32((config.time_reduction - 1) as i32);
+    let reduction = b.const_i32(config.time_reduction as i32);
+    let valid_length = b.add(&frame_length, &reduction_minus_one);
+    let valid_length = b.divide(&valid_length, &reduction);
+    let positions = b.iota(encoded_len);
+    let valid_length_vector = b.broadcast(&valid_length, &[], vec![encoded_len]);
+    let active_rows = b.compare("LT", &positions, &valid_length_vector, "SIGNED");
+    let mut encoded = mask_rows(&mut b, &subsampled.projected, &active_rows);
+    let relative_bias = relative_attention_bias(&mut b, &weights, config, encoded_len);
+    let mut block0 = None;
+    let mut selected_blocks = Vec::new();
+    for layer in 0..config.num_blocks {
+        let values = conformer_block(
+            &mut b,
+            &weights,
+            config,
+            &encoded,
+            &relative_bias,
+            &active_rows,
+            layer,
+        );
+        encoded = values.output.clone();
+        if layer == 0 {
+            block0 = Some(values);
+        } else if matches!(layer, 1 | 5 | 11 | 17 | 23) {
+            selected_blocks.push((layer, encoded.clone()));
+        }
+    }
+    let projection = project_audio(
+        &mut b,
+        &weights,
+        config,
+        &encoded,
+        &projection_mode,
+        &active_rows,
+    );
+    let signature = declarations
+        .iter()
+        .enumerate()
+        .map(|(index, (ty, location))| {
+            format!("%arg{index}: {} loc(\"{}\")", ty.render(), location)
+        })
+        .collect::<Vec<_>>()
+        .join(", ");
+    if diagnostic {
+        let block0 = block0.ok_or("Phi4MM diagnostic requires Conformer block 0")?;
+        let selected = |layer: usize| {
+            selected_blocks
+                .iter()
+                .find_map(|(candidate, value)| (*candidate == layer).then_some(value.clone()))
+                .ok_or_else(|| format!("Phi4MM diagnostic requires Conformer block {layer}"))
+        };
+        let reshape_encoder = |b: &mut Builder, value: &Val| {
+            b.reshape(value, vec![1, encoded_len, config.attention_dim])
+        };
+        let reshape_projection = |b: &mut Builder, value: &Val| {
+            b.reshape(value, vec![1, encoded_len, config.projection_hidden])
+        };
+        let mut checkpoints = vec![
+            subsampled.conv0,
+            subsampled.conv1_depthwise,
+            subsampled.conv1_pointwise,
+            subsampled.conv1,
+            subsampled.conv2_depthwise,
+            subsampled.conv2_pointwise,
+            subsampled.conv2,
+            reshape_encoder(&mut b, &subsampled.projected),
+            reshape_encoder(&mut b, &block0.after_ff_in),
+            reshape_encoder(&mut b, &block0.attention),
+            reshape_encoder(&mut b, &block0.after_attention),
+            reshape_encoder(&mut b, &block0.convolution),
+            reshape_encoder(&mut b, &block0.after_convolution),
+            reshape_encoder(&mut b, &block0.ff_out),
+            reshape_encoder(&mut b, &block0.output),
+        ];
+        for layer in [1, 5, 11, 17, 23] {
+            checkpoints.push(reshape_encoder(&mut b, &selected(layer)?));
+        }
+        checkpoints.extend([
+            reshape_encoder(&mut b, &encoded),
+            reshape_projection(&mut b, &projection.speech_first),
+            reshape_projection(&mut b, &projection.speech),
+            reshape_projection(&mut b, &projection.vision_first),
+            reshape_projection(&mut b, &projection.vision),
+        ]);
+        let specs = phi4_audio_diagnostic_specs(config, frame_bucket)?;
+        if checkpoints.len() != specs.len() {
+            return Err("Phi4MM diagnostic checkpoint schema length drifted".to_string());
+        }
+        for (checkpoint, spec) in checkpoints.iter().zip(&specs) {
+            if checkpoint.ty.shape != spec.shape {
+                return Err(format!(
+                    "Phi4MM diagnostic checkpoint {} has shape {:?}, expected {:?}",
+                    spec.name, checkpoint.ty.shape, spec.shape
+                ));
+            }
+        }
+        let return_values = checkpoints
+            .iter()
+            .map(|value| value.name.as_str())
+            .chain(std::iter::once(valid_length.name.as_str()))
+            .collect::<Vec<_>>()
+            .join(", ");
+        let return_types = checkpoints
+            .iter()
+            .map(|value| value.ty.render())
+            .chain(std::iter::once("tensor<i32>".to_string()))
+            .collect::<Vec<_>>()
+            .join(", ");
+        Ok(format!(
+            "module @audio {{\n  func.func public @diagnostic({signature}) -> ({return_types}) {{\n{body}    return {return_values} : {return_types} loc(\"audio diagnostic checkpoints\")\n  }}\n}}\n",
+            body = b.body(),
+        ))
+    } else {
+        let projected = b.reshape(
+            &projection.selected,
+            vec![1, encoded_len, config.projection_hidden],
+        );
+        Ok(format!(
+            "module @audio {{\n  // #874 oracle stages: features tensor<1x{frame_bucket}x{input_size}xf32>, encoder tensor<1x{encoded_len}x{attention_dim}xf32>, projection {projected_ty}\n  func.func public @main({signature}) -> ({projected_ty}, tensor<i32>) {{\n{body}    return {projected}, {length} : {projected_ty}, tensor<i32> loc(\"projection\")\n  }}\n}}\n",
+            projected_ty = projected.ty.render(),
+            body = b.body(),
+            projected = projected.name,
+            length = valid_length.name,
+            input_size = config.input_size,
+            attention_dim = config.attention_dim,
+        ))
+    }
+}
+
+pub(crate) fn emit_phi4_audio(
+    config: &Phi4AudioConfig,
+    frame_bucket: usize,
+) -> Result<String, String> {
+    emit_phi4_audio_with(config, frame_bucket, Precision::F32)
+}
+
+#[cfg(test)]
+mod tests {
+    use std::collections::HashMap;
+    use std::fs::File;
+    use std::io::Read;
+    use std::path::PathBuf;
+    use std::process::Command;
+
+    use super::*;
+
+    fn pinned_config() -> Phi4AudioConfig {
+        Phi4AudioConfig {
+            input_size: 80,
+            attention_dim: 1024,
+            attention_heads: 16,
+            num_blocks: 24,
+            linear_units: 1536,
+            time_reduction: 8,
+            conv_channels: 1024,
+            kernel_size: 3,
+            relative_bias_max_distance: 500,
+            projection_hidden: 3072,
+        }
+    }
+
+    fn minimum_compile_config() -> Phi4AudioConfig {
+        Phi4AudioConfig {
+            // Keep the released 80-bin input so all grouped Conv2D stages
+            // exercise production-shaped frequency reductions.
+            input_size: 80,
+            attention_dim: 4,
+            attention_heads: 2,
+            num_blocks: 1,
+            linear_units: 3,
+            time_reduction: 8,
+            conv_channels: 2,
+            kernel_size: 3,
+            relative_bias_max_distance: 4,
+            projection_hidden: 6,
+        }
+    }
+
+    fn safetensors_header(path: &std::path::Path) -> Value {
+        let mut file = File::open(path).expect("open safetensors shard");
+        let mut length = [0_u8; 8];
+        file.read_exact(&mut length).expect("read header length");
+        let length = usize::try_from(u64::from_le_bytes(length)).expect("header length fits usize");
+        let mut bytes = vec![0_u8; length];
+        file.read_exact(&mut bytes)
+            .expect("read safetensors header");
+        serde_json::from_slice(&bytes).expect("parse safetensors header")
+    }
+
+    fn compile_audio_graph_cpu(
+        config: &Phi4AudioConfig,
+        frame_bucket: usize,
+        precision: Precision,
+        label: &str,
+    ) {
+        let compiler =
+            PathBuf::from(std::env::var("MLXCEL_XLA_IREE_COMPILE").expect("set compiler path"));
+        let graph =
+            emit_phi4_audio_with(config, frame_bucket, precision).expect("emit audio graph");
+        let stem = format!("mlxcel-phi4mm-audio-{label}-{}", std::process::id());
+        let input = std::env::temp_dir().join(format!("{stem}.mlir"));
+        let output = std::env::temp_dir().join(format!("{stem}.vmfb"));
+        std::fs::write(&input, graph).expect("write temporary StableHLO");
+        let result = Command::new(compiler)
+            .arg("--iree-input-type=stablehlo")
+            .arg("--iree-hal-target-device=local")
+            .arg("--iree-hal-local-target-device-backends=llvm-cpu")
+            .arg(&input)
+            .arg("-o")
+            .arg(&output)
+            .output()
+            .expect("run iree-compile");
+        let _ = std::fs::remove_file(&input);
+        let _ = std::fs::remove_file(&output);
+        assert!(
+            result.status.success(),
+            "iree-compile failed:\nstdout:\n{}\nstderr:\n{}",
+            String::from_utf8_lossy(&result.stdout),
+            String::from_utf8_lossy(&result.stderr)
+        );
+    }
+
+    #[test]
+    fn pinned_schema_is_complete_and_propagates_exact_lengths() {
+        let config = pinned_config();
+        let specs = phi4_audio_weight_specs(&config);
+        assert_eq!(specs.len(), 887);
+        // #874's official host oracle uses the released non-causal NeMo
+        // subsampler: three kernel-3, symmetric-pad-1, stride-2 stages.
+        for (frames, projected_rows) in [
+            (1, 1),
+            (7, 1),
+            (8, 1),
+            (9, 2),
+            (148, 19),
+            (351, 44),
+            (500, 63),
+        ] {
+            assert_eq!(
+                config.encoded_bucket_len(frames).unwrap(),
+                projected_rows,
+                "{frames} host feature frames"
+            );
+            assert_eq!(
+                config.encoded_valid_len(frames).unwrap(),
+                projected_rows,
+                "{frames} valid host feature frames"
+            );
+        }
+        let oracle: Value = serde_json::from_str(include_str!(
+            "../../../../../tests/fixtures/phi4mm_audio_parity.json"
+        ))
+        .expect("parse #874 pinned official oracle");
+        assert_eq!(oracle["feature_shape"], serde_json::json!([1, 351, 80]));
+        assert_eq!(oracle["audio_embed_size"], 44);
+        assert_eq!(
+            config
+                .encoded_valid_len(oracle["feature_shape"][1].as_u64().unwrap() as usize)
+                .unwrap(),
+            oracle["audio_embed_size"].as_u64().unwrap() as usize
+        );
+    }
+
+    #[test]
+    fn shape_validation_rejects_missing_and_drifted_weights() {
+        let config = pinned_config();
+        let specs = phi4_audio_weight_specs(&config);
+        let mut shapes = specs
+            .iter()
+            .map(|spec| (spec.name.clone(), spec.shape.clone()))
+            .collect::<HashMap<_, _>>();
+        validate_phi4_audio_weight_shapes(&config, |name| shapes.get(name).cloned()).unwrap();
+        let first = specs[0].name.clone();
+        shapes.remove(&first);
+        assert!(
+            validate_phi4_audio_weight_shapes(&config, |name| shapes.get(name).cloned())
+                .unwrap_err()
+                .contains("missing")
+        );
+        shapes.insert(first.clone(), vec![79]);
+        let error = validate_phi4_audio_weight_shapes(&config, |name| shapes.get(name).cloned())
+            .unwrap_err();
+        assert!(error.contains(&first));
+        assert!(error.contains("expected [80]"));
+    }
+
+    #[test]
+    fn emitted_graph_tracks_oracle_stages_projection_mode_and_length() {
+        let graph = emit_phi4_audio(&minimum_compile_config(), 32).expect("emit minimum graph");
+        assert!(graph.contains("module @audio"));
+        assert!(graph.contains("func.func public @main"));
+        assert!(graph.contains("loc(\"features\")"));
+        assert!(graph.contains("loc(\"frame_mask\")"));
+        assert!(graph.contains("loc(\"frame_length\")"));
+        assert!(graph.contains("loc(\"projection_mode\")"));
+        assert!(graph.contains(
+            "#874 oracle stages: features tensor<1x32x80xf32>, encoder tensor<1x4x4xf32>, projection tensor<1x4x6xf32>"
+        ));
+        assert!(graph.contains("stablehlo.convolution"));
+        assert!(graph.contains("stablehlo.gather"));
+        assert!(graph.contains("stablehlo.dot_general"));
+        assert!(graph.contains("chlo.erf"));
+        assert!(graph.contains("audio_projection.speech.0.weight"));
+        assert!(graph.contains("audio_projection.vision.0.weight"));
+        assert!(graph.contains("stablehlo.divide"));
+        assert!(graph.contains("-> (tensor<1x4x6xf32>, tensor<i32>)"));
+        assert!(!graph.contains("model.layers.0.self_attn"));
+    }
+
+    #[test]
+    fn emitted_graph_rejects_invalid_attention_and_oversized_chunk() {
+        let mut config = minimum_compile_config();
+        config.attention_heads = 3;
+        assert!(
+            emit_phi4_audio(&config, 8)
+                .unwrap_err()
+                .contains("divisible")
+        );
+
+        let config = minimum_compile_config();
+        assert!(
+            emit_phi4_audio(&config, 33)
+                .unwrap_err()
+                .contains("exceeding the pinned")
+        );
+    }
+
+    #[test]
+    #[ignore = "requires MLXCEL_XLA_IREE_COMPILE pointing at iree-compile"]
+    fn minimum_audio_graph_compiles_with_iree_cpu() {
+        compile_audio_graph_cpu(&minimum_compile_config(), 32, Precision::F32, "minimum");
+    }
+
+    #[test]
+    #[ignore = "requires MLXCEL_XLA_IREE_COMPILE pointing at iree-compile"]
+    fn pinned_full_audio_graph_compiles_with_iree_cpu() {
+        // The 512-frame production bucket contains #874's real 351-frame
+        // oracle sample and emits all 24 Conformer blocks and 887 weights.
+        compile_audio_graph_cpu(&pinned_config(), 512, Precision::F32, "pinned-full");
+    }
+
+    #[test]
+    #[ignore = "requires MLXCEL_XLA_IREE_COMPILE pointing at iree-compile"]
+    fn pinned_full_mixed_bf16_audio_graph_compiles_with_iree_cpu() {
+        compile_audio_graph_cpu(
+            &pinned_config(),
+            512,
+            Precision::Bf16,
+            "pinned-full-mixed-bf16",
+        );
+    }
+
+    #[test]
+    #[ignore = "requires PHI4MM_MODEL_DIR pointing at the pinned official checkpoint"]
+    fn real_checkpoint_has_exact_audio_schema_and_shapes() {
+        let model_dir =
+            PathBuf::from(std::env::var("PHI4MM_MODEL_DIR").expect("set PHI4MM_MODEL_DIR"));
+        let config_text =
+            std::fs::read_to_string(model_dir.join("config.json")).expect("read config");
+        let config = Phi4AudioConfig::from_json_str(&config_text).expect("pinned audio config");
+        let index: Value = serde_json::from_str(
+            &std::fs::read_to_string(model_dir.join("model.safetensors.index.json"))
+                .expect("read weight index"),
+        )
+        .expect("parse weight index");
+        let weight_map = index["weight_map"].as_object().expect("weight_map object");
+        let mut headers = HashMap::new();
+        for file_name in weight_map.values().filter_map(Value::as_str) {
+            headers
+                .entry(file_name.to_string())
+                .or_insert_with(|| safetensors_header(&model_dir.join(file_name)));
+        }
+        validate_phi4_audio_weight_shapes(&config, |name| {
+            let file_name = weight_map.get(name)?.as_str()?;
+            let metadata = headers.get(file_name)?.get(name)?;
+            assert_eq!(
+                metadata.get("dtype").and_then(Value::as_str),
+                Some("BF16"),
+                "{name} dtype"
+            );
+            metadata
+                .get("shape")?
+                .as_array()?
+                .iter()
+                .map(|dim| dim.as_u64().and_then(|value| usize::try_from(value).ok()))
+                .collect()
+        })
+        .expect("complete exact Phi4MM audio schema");
+    }
+}
diff --git a/src/lib/mlxcel-xla/src/emitter/rope.rs b/src/lib/mlxcel-xla/src/emitter/rope.rs
index 88a4e0dbc..ea01dd30c 100644
--- a/src/lib/mlxcel-xla/src/emitter/rope.rs
+++ b/src/lib/mlxcel-xla/src/emitter/rope.rs
@@ -28,6 +28,16 @@ pub fn inv_freq(c: &Config) -> Vec<f64> {
             high_freq_factor,
             orig_ctx,
         } => llama3_inv_freq(c, *factor, *low_freq_factor, *high_freq_factor, *orig_ctx),
+        RopeScaling::LongRope { factors, .. } => {
+            let d = c.rotary_width();
+            factors
+                .iter()
+                .enumerate()
+                .map(|(index, factor)| {
+                    1.0 / (factor * c.rope_theta.powf((2 * index) as f64 / d as f64))
+                })
+                .collect()
+        }
     }
 }
 
@@ -130,7 +140,14 @@ pub fn rope_tables_from_inv(
 /// Build cos and sin tables of shape [max_seq, rotary_width] for the config's
 /// global RoPE scheme (half-split or interleaved).
 pub fn rope_tables(c: &Config, max_seq: usize) -> (Vec<f32>, Vec<f32>) {
-    rope_tables_from_inv(&inv_freq(c), c.rotary_width(), max_seq, c.rope_interleaved)
+    let (mut cos, mut sin) =
+        rope_tables_from_inv(&inv_freq(c), c.rotary_width(), max_seq, c.rope_interleaved);
+    if let RopeScaling::LongRope { amplitude, .. } = &c.rope {
+        let amplitude = *amplitude as f32;
+        cos.iter_mut().for_each(|value| *value *= amplitude);
+        sin.iter_mut().for_each(|value| *value *= amplitude);
+    }
+    (cos, sin)
 }
 
 /// Build the local cos/sin tables for a config with a distinct local RoPE base
diff --git a/src/lib/mlxcel-xla/src/iree.rs b/src/lib/mlxcel-xla/src/iree.rs
index afe160776..e31100190 100644
--- a/src/lib/mlxcel-xla/src/iree.rs
+++ b/src/lib/mlxcel-xla/src/iree.rs
@@ -226,9 +226,11 @@ unsafe extern "C" {
         deepstack_layers: c_int,
         deepstack_visual_positions: c_int,
         deepstack_target_layers: *const c_int,
+        has_adapter_modes: c_int,
     ) -> *mut XlaCtx;
     fn xla_llama_prefill(
         c: *mut XlaCtx,
+        adapter_mode: c_int,
         tokens: *const c_int,
         lp: c_int,
         positions: *const c_int,
@@ -237,6 +239,7 @@ unsafe extern "C" {
     ) -> c_int;
     fn xla_llama_decode(
         c: *mut XlaCtx,
+        adapter_mode: c_int,
         token: c_int,
         pos: c_int,
         cache_len: c_int,
@@ -244,6 +247,7 @@ unsafe extern "C" {
     ) -> c_int;
     fn xla_llama_decode_mrope(
         c: *mut XlaCtx,
+        adapter_mode: c_int,
         token: c_int,
         positions: *const c_int,
         cache_len: c_int,
@@ -257,6 +261,7 @@ unsafe extern "C" {
     fn xla_llama_prefill_slot_logits(
         c: *mut XlaCtx,
         slot: c_int,
+        adapter_mode: c_int,
         tokens: *const c_int,
         lp: c_int,
         positions: *const c_int,
@@ -268,6 +273,7 @@ unsafe extern "C" {
     fn xla_llama_prefill_diagnostics_slot(
         c: *mut XlaCtx,
         slot: c_int,
+        adapter_mode: c_int,
         tokens: *const c_int,
         lp: c_int,
         positions: *const c_int,
@@ -277,6 +283,7 @@ unsafe extern "C" {
     ) -> c_int;
     fn xla_llama_prefill_embeddings(
         c: *mut XlaCtx,
+        adapter_mode: c_int,
         position_mode: c_int,
         embeddings: *const XlaTensorDesc,
         positions: *const XlaTensorDesc,
@@ -287,6 +294,7 @@ unsafe extern "C" {
     fn xla_llama_prefill_embeddings_slot_logits(
         c: *mut XlaCtx,
         slot: c_int,
+        adapter_mode: c_int,
         position_mode: c_int,
         embeddings: *const XlaTensorDesc,
         positions: *const XlaTensorDesc,
@@ -299,6 +307,7 @@ unsafe extern "C" {
     fn xla_llama_prefill_embeddings_slot_diagnostics(
         c: *mut XlaCtx,
         slot: c_int,
+        adapter_mode: c_int,
         position_mode: c_int,
         embeddings: *const XlaTensorDesc,
         positions: *const XlaTensorDesc,
@@ -311,6 +320,7 @@ unsafe extern "C" {
     ) -> c_int;
     fn xla_llama_prefill_embeddings_ple(
         c: *mut XlaCtx,
+        adapter_mode: c_int,
         position_mode: c_int,
         embeddings: *const XlaTensorDesc,
         dense_ple: *const XlaTensorDesc,
@@ -322,6 +332,7 @@ unsafe extern "C" {
     fn xla_llama_prefill_embeddings_ple_slot_logits(
         c: *mut XlaCtx,
         slot: c_int,
+        adapter_mode: c_int,
         position_mode: c_int,
         embeddings: *const XlaTensorDesc,
         dense_ple: *const XlaTensorDesc,
@@ -333,6 +344,7 @@ unsafe extern "C" {
     ) -> c_int;
     fn xla_llama_prefill_embeddings_deepstack(
         c: *mut XlaCtx,
+        adapter_mode: c_int,
         position_mode: c_int,
         embeddings: *const XlaTensorDesc,
         positions: *const XlaTensorDesc,
@@ -348,6 +360,7 @@ unsafe extern "C" {
     fn xla_llama_prefill_embeddings_deepstack_slot_logits(
         c: *mut XlaCtx,
         slot: c_int,
+        adapter_mode: c_int,
         position_mode: c_int,
         embeddings: *const XlaTensorDesc,
         positions: *const XlaTensorDesc,
@@ -364,6 +377,7 @@ unsafe extern "C" {
     fn xla_llama_decode_ragged_logits(
         c: *mut XlaCtx,
         bsz: c_int,
+        adapter_modes: *const c_int,
         tokens: *const c_int,
         pos: *const c_int,
         cache_len: *const c_int,
@@ -373,6 +387,7 @@ unsafe extern "C" {
     fn xla_llama_decode_ragged_mrope_logits(
         c: *mut XlaCtx,
         bsz: c_int,
+        adapter_modes: *const c_int,
         tokens: *const c_int,
         positions: *const c_int,
         cache_len: *const c_int,
@@ -397,6 +412,22 @@ fn checked_ffi_int(value: usize, name: &str) -> Result<c_int, String> {
     c_int::try_from(value).map_err(|_| format!("{name}={value} does not fit the IREE C ABI c_int"))
 }
 
+fn validate_adapter_modes(modes: &[i32], supported: bool) -> Result<(), String> {
+    for (row, &mode) in modes.iter().enumerate() {
+        if !(0..=2).contains(&mode) {
+            return Err(format!(
+                "adapter mode {mode} for row {row} is outside the supported range 0..=2"
+            ));
+        }
+        if !supported && mode != 0 {
+            return Err(format!(
+                "adapter mode {mode} for row {row} is incompatible with this runtime bundle"
+            ));
+        }
+    }
+    Ok(())
+}
+
 fn checked_ffi_i64(value: usize, name: &str) -> Result<i64, String> {
     i64::try_from(value).map_err(|_| format!("{name}={value} does not fit the IREE C ABI i64"))
 }
@@ -521,6 +552,10 @@ impl RuntimeConfig {
         }
     }
 
+    fn has_adapter_modes(&self) -> bool {
+        matches!(self, Self::Dense(config) if config.phi4_lora.is_some())
+    }
+
     fn supports_f16_resident(&self) -> bool {
         matches!(self, Self::Dense(config) if config.supports_f16_resident())
     }
@@ -1764,6 +1799,7 @@ fn create_ctx_with_diagnostics(
             } else {
                 deepstack_target_layers.as_ptr()
             },
+            i32::from(cfg.has_adapter_modes()),
         )
     };
     // Weights are resident on the device now; free the host copy.
@@ -1784,6 +1820,8 @@ pub struct IreeLlama {
     ctx: *mut XlaCtx,
     context_capacity: usize,
     hidden_size: usize,
+    has_adapter_modes: bool,
+    adapter_mode: i32,
     dense_ple_shape: Option<[usize; 3]>,
     position_mode: PreparedPositionMode,
     decode_position: DecodePositionState,
@@ -1803,6 +1841,8 @@ impl IreeLlama {
             ctx,
             context_capacity: cfg.context_capacity(),
             hidden_size: cfg.hidden(),
+            has_adapter_modes: cfg.has_adapter_modes(),
+            adapter_mode: 0,
             dense_ple_shape: cfg.dense_ple_shape(),
             position_mode: cfg.position_mode(),
             decode_position: DecodePositionState::default(),
@@ -1857,12 +1897,18 @@ impl IreeLlama {
         let (embeddings, positions, attention_bias) = prepared_descriptors(&prepared)?;
         let real_len = i32::try_from(prepared.effective_len)
             .map_err(|_| "prepared effective length does not fit i32".to_string())?;
+        if !self.has_adapter_modes && prepared.adapter_mode != 0 {
+            return Err(
+                "prepared adapter mode is incompatible with this runtime bundle".to_string(),
+            );
+        }
         let mut out = 0i32;
         // Safety: every descriptor references a validated owned vector that
         // outlives the call; the C shim repeats dtype/rank/shape/byte checks.
         let rc = unsafe {
             xla_llama_prefill_embeddings(
                 self.ctx,
+                prepared.adapter_mode,
                 prepared.positions.mode().ffi_code(),
                 &embeddings,
                 &positions,
@@ -1872,6 +1918,7 @@ impl IreeLlama {
             )
         };
         let result = if rc == 0 {
+            self.adapter_mode = prepared.adapter_mode;
             Ok(out)
         } else {
             Err(format!("xla_llama_prefill_embeddings failed (status {rc})"))
@@ -1913,6 +1960,7 @@ impl IreeLlama {
         let rc = unsafe {
             xla_llama_prefill_embeddings_deepstack(
                 self.ctx,
+                0,
                 prepared.positions.mode().ffi_code(),
                 &embeddings,
                 &positions,
@@ -1927,6 +1975,7 @@ impl IreeLlama {
             )
         };
         let result = if rc == 0 {
+            self.adapter_mode = 0;
             Ok(out)
         } else {
             Err(format!(
@@ -1967,6 +2016,7 @@ impl IreeLlama {
         let rc = unsafe {
             xla_llama_prefill_embeddings_ple(
                 self.ctx,
+                0,
                 prepared.positions.mode().ffi_code(),
                 &embeddings,
                 &dense_ple,
@@ -1977,6 +2027,7 @@ impl IreeLlama {
             )
         };
         let result = if rc == 0 {
+            self.adapter_mode = 0;
             Ok(out)
         } else {
             Err(format!(
@@ -2008,6 +2059,7 @@ impl IreeLlama {
         let rc = unsafe {
             xla_llama_prefill(
                 self.ctx,
+                0,
                 tokens.as_ptr(),
                 capacity,
                 positions.as_ptr(),
@@ -2016,6 +2068,7 @@ impl IreeLlama {
             )
         };
         let result = if rc == 0 {
+            self.adapter_mode = 0;
             Ok(out)
         } else {
             Err(format!("xla_llama_prefill failed (status {rc})"))
@@ -2036,7 +2089,16 @@ impl IreeLlama {
         let rc = match self.position_mode {
             PreparedPositionMode::OneD => {
                 // Safety: the shim threads its own resident KV; only scalars cross here.
-                unsafe { xla_llama_decode(self.ctx, token, cache_len, cache_len, &mut out) }
+                unsafe {
+                    xla_llama_decode(
+                        self.ctx,
+                        self.adapter_mode,
+                        token,
+                        cache_len,
+                        cache_len,
+                        &mut out,
+                    )
+                }
             }
             PreparedPositionMode::Mrope3D => {
                 let coordinate = self
@@ -2046,7 +2108,14 @@ impl IreeLlama {
                 let positions = [coordinate; 3];
                 // Safety: positions is exactly the explicit rank-1 `[3]` ABI payload.
                 unsafe {
-                    xla_llama_decode_mrope(self.ctx, token, positions.as_ptr(), cache_len, &mut out)
+                    xla_llama_decode_mrope(
+                        self.ctx,
+                        self.adapter_mode,
+                        token,
+                        positions.as_ptr(),
+                        cache_len,
+                        &mut out,
+                    )
                 }
             }
         };
@@ -2096,6 +2165,7 @@ pub struct IreeRaggedLlama {
     vocab: usize,
     context_capacity: usize,
     hidden_size: usize,
+    has_adapter_modes: bool,
     #[cfg(feature = "diagnostics")]
     model_layers: usize,
     dense_ple_shape: Option<[usize; 3]>,
@@ -2238,6 +2308,7 @@ impl IreeRaggedLlama {
             vocab: cfg.vocab(),
             context_capacity: cfg.context_capacity(),
             hidden_size: cfg.hidden(),
+            has_adapter_modes: cfg.has_adapter_modes(),
             #[cfg(feature = "diagnostics")]
             model_layers: cfg.n_layers(),
             dense_ple_shape: cfg.dense_ple_shape(),
@@ -2364,6 +2435,7 @@ impl IreeRaggedLlama {
             xla_llama_prefill_diagnostics_slot(
                 self.ctx,
                 slot,
+                0,
                 tokens.as_ptr(),
                 capacity,
                 positions.as_ptr(),
@@ -2412,6 +2484,7 @@ impl IreeRaggedLlama {
             xla_llama_prefill_slot_logits(
                 self.ctx,
                 slot,
+                0,
                 tokens.as_ptr(),
                 capacity,
                 positions.as_ptr(),
@@ -2452,6 +2525,11 @@ impl IreeRaggedLlama {
                 "prepared IREE payload is incompatible with this runtime bundle".to_string(),
             );
         }
+        if !self.has_adapter_modes && prepared.adapter_mode != 0 {
+            return Err(
+                "prepared adapter mode is incompatible with this runtime bundle".to_string(),
+            );
+        }
         let (embeddings, positions, attention_bias) = prepared_descriptors(prepared)?;
         let real_len = i32::try_from(prepared.effective_len)
             .map_err(|_| "prepared effective length does not fit i32".to_string())?;
@@ -2464,6 +2542,7 @@ impl IreeRaggedLlama {
             xla_llama_prefill_embeddings_slot_logits(
                 self.ctx,
                 slot,
+                prepared.adapter_mode,
                 prepared.positions.mode().ffi_code(),
                 &embeddings,
                 &positions,
@@ -2509,6 +2588,11 @@ impl IreeRaggedLlama {
                 "prepared IREE payload is incompatible with this runtime bundle".to_string(),
             );
         }
+        if !self.has_adapter_modes && prepared.adapter_mode != 0 {
+            return Err(
+                "prepared adapter mode is incompatible with this runtime bundle".to_string(),
+            );
+        }
         if kv_width == 0 {
             return Err("diagnostic KV width must be greater than zero".to_string());
         }
@@ -2531,6 +2615,7 @@ impl IreeRaggedLlama {
             xla_llama_prefill_embeddings_slot_diagnostics(
                 self.ctx,
                 slot,
+                prepared.adapter_mode,
                 prepared.positions.mode().ffi_code(),
                 &embeddings,
                 &positions,
@@ -2598,6 +2683,7 @@ impl IreeRaggedLlama {
             xla_llama_prefill_embeddings_deepstack_slot_logits(
                 self.ctx,
                 slot,
+                0,
                 prepared.positions.mode().ffi_code(),
                 &embeddings,
                 &positions,
@@ -2654,6 +2740,7 @@ impl IreeRaggedLlama {
             xla_llama_prefill_embeddings_ple_slot_logits(
                 self.ctx,
                 slot,
+                0,
                 prepared.positions.mode().ffi_code(),
                 &embeddings,
                 &dense_ple,
@@ -2681,16 +2768,33 @@ impl IreeRaggedLlama {
         tokens: &[i32],
         pos: &[i32],
         cache_len: &[i32],
+    ) -> Result<Vec<f32>, String> {
+        self.decode_ragged_logits_with_modes(&vec![0; self.b_max], tokens, pos, cache_len)
+    }
+
+    /// Mode-aware ragged decode. Each row retains the adapter selected at
+    /// prefill; inactive rows must carry the language mode (`0`).
+    pub fn decode_ragged_logits_with_modes(
+        &mut self,
+        adapter_modes: &[i32],
+        tokens: &[i32],
+        pos: &[i32],
+        cache_len: &[i32],
     ) -> Result<Vec<f32>, String> {
         if self.position_mode != PreparedPositionMode::OneD {
             return Err("1D decode was requested from an M-RoPE runtime bundle".to_string());
         }
-        if tokens.len() != self.b_max || pos.len() != self.b_max || cache_len.len() != self.b_max {
+        if adapter_modes.len() != self.b_max
+            || tokens.len() != self.b_max
+            || pos.len() != self.b_max
+            || cache_len.len() != self.b_max
+        {
             return Err(format!(
-                "decode_ragged_logits expects per-row arrays of length b_max = {}",
+                "decode_ragged_logits expects adapter/token/position/cache arrays of length b_max = {}",
                 self.b_max
             ));
         }
+        validate_adapter_modes(adapter_modes, self.has_adapter_modes)?;
         for row in 0..self.b_max {
             if pos[row] != cache_len[row]
                 || pos[row] < 0
@@ -2711,6 +2815,7 @@ impl IreeRaggedLlama {
             xla_llama_decode_ragged_logits(
                 self.ctx,
                 b_max,
+                adapter_modes.as_ptr(),
                 tokens.as_ptr(),
                 pos.as_ptr(),
                 cache_len.as_ptr(),
@@ -2733,19 +2838,37 @@ impl IreeRaggedLlama {
         tokens: &[i32],
         positions: &[[i32; 3]],
         cache_len: &[i32],
+    ) -> Result<Vec<f32>, String> {
+        self.decode_ragged_mrope_logits_with_modes(
+            &vec![0; self.b_max],
+            tokens,
+            positions,
+            cache_len,
+        )
+    }
+
+    /// Mode-aware M-RoPE ragged decode.
+    pub fn decode_ragged_mrope_logits_with_modes(
+        &mut self,
+        adapter_modes: &[i32],
+        tokens: &[i32],
+        positions: &[[i32; 3]],
+        cache_len: &[i32],
     ) -> Result<Vec<f32>, String> {
         if self.position_mode != PreparedPositionMode::Mrope3D {
             return Err("M-RoPE decode was requested from a 1D runtime bundle".to_string());
         }
-        if tokens.len() != self.b_max
+        if adapter_modes.len() != self.b_max
+            || tokens.len() != self.b_max
             || positions.len() != self.b_max
             || cache_len.len() != self.b_max
         {
             return Err(format!(
-                "decode_ragged_mrope_logits expects per-row arrays of length b_max = {}",
+                "decode_ragged_mrope_logits expects adapter/token/position/cache arrays of length b_max = {}",
                 self.b_max
             ));
         }
+        validate_adapter_modes(adapter_modes, self.has_adapter_modes)?;
         for (row, (&physical, coordinates)) in cache_len.iter().zip(positions.iter()).enumerate() {
             if physical < 0 || physical as usize >= self.context_capacity {
                 return Err(format!(
@@ -2772,6 +2895,7 @@ impl IreeRaggedLlama {
             xla_llama_decode_ragged_mrope_logits(
                 self.ctx,
                 b_max,
+                adapter_modes.as_ptr(),
                 tokens.as_ptr(),
                 flat_positions.as_ptr(),
                 cache_len.as_ptr(),
diff --git a/src/lib/mlxcel-xla/src/lib.rs b/src/lib/mlxcel-xla/src/lib.rs
index 4cbab7716..19fd11043 100644
--- a/src/lib/mlxcel-xla/src/lib.rs
+++ b/src/lib/mlxcel-xla/src/lib.rs
@@ -65,6 +65,8 @@ mod aux_smoke;
 #[cfg(feature = "iree")]
 mod iree;
 #[cfg(feature = "iree")]
+mod phi4_audio;
+#[cfg(feature = "iree")]
 mod vision_runtime;
 
 // The continuous-batching engine (#449 M3 Stage 2b). Present under `iree` (real
@@ -139,6 +141,14 @@ pub use emitter::{
 };
 #[cfg(feature = "diagnostics")]
 pub use iree::PreparedPrefillDiagnostics;
+#[cfg(feature = "iree")]
+pub use phi4_audio::{
+    PHI4MM_AUDIO_CHECKPOINT_REVISION, PHI4MM_AUDIO_FRAME_BUCKETS, Phi4AudioOutput,
+    Phi4AudioProjectionMode, Phi4AudioRuntime, phi4_audio_bucket_for_frames,
+};
+#[cfg(feature = "iree")]
+#[doc(hidden)]
+pub use phi4_audio::{Phi4AudioCheckpoint, Phi4AudioDiagnosticRuntime, Phi4AudioDiagnostics};
 #[cfg(feature = "diagnostics")]
 pub use vision_runtime::{IreeVisionDiagnosticProjector, VisionDiagnosticProjection};
 #[cfg(feature = "iree")]
diff --git a/src/lib/mlxcel-xla/src/phi4_audio.rs b/src/lib/mlxcel-xla/src/phi4_audio.rs
new file mode 100644
index 000000000..25bf812bc
--- /dev/null
+++ b/src/lib/mlxcel-xla/src/phi4_audio.rs
@@ -0,0 +1,860 @@
+// Copyright 2025-2026 Lablup Inc. and Jeongkyu Shin
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+//! Resident IREE runtime for Phi4MM's pinned cascaded Conformer.
+//!
+//! Host code supplies the #874 SpeechLib feature frames. This module owns only
+//! the audio encoder/projection weights and the `audio.main` VMFB; it never
+//! constructs or falls back to an MLX decoder.
+
+use std::collections::BTreeMap;
+use std::fs::File;
+use std::io::Read;
+use std::mem::size_of;
+use std::path::{Path, PathBuf};
+use std::process::Command;
+
+use memmap2::Mmap;
+use safetensors::{Dtype, SafeTensors};
+use sha2::{Digest, Sha256};
+
+use crate::aux::{
+    AuxiliaryInput, AuxiliaryOutput, AuxiliaryTensorDType, AuxiliaryWeight, AuxiliaryWeightDType,
+    IreeAuxiliaryModule,
+};
+use crate::aux_manifest::{AuxiliaryArtifactContract, ensure_qualified_auxiliary_artifact};
+use crate::emitter::{
+    Phi4AudioConfig, Precision, emit_phi4_audio_diagnostic_with, emit_phi4_audio_with,
+    phi4_audio_diagnostic_specs, phi4_audio_weight_specs,
+};
+use crate::iree::{cached_vmfb_path, compile_one_to, iree_compile_bin, target_flags};
+use crate::weights::{bf16_to_f32, f16_to_f32};
+
+/// Immutable upstream checkpoint revision qualified by #874.
+pub const PHI4MM_AUDIO_CHECKPOINT_REVISION: &str = "93f923e1a7727d1c4f446756212d9d3e8fcc5d81";
+
+/// Static production buckets. The largest bucket maps to 500 Conformer rows.
+pub const PHI4MM_AUDIO_FRAME_BUCKETS: &[usize] = &[64, 128, 256, 512, 1024, 2048, 4000];
+
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub enum Phi4AudioProjectionMode {
+    Speech,
+    Vision,
+}
+
+impl Phi4AudioProjectionMode {
+    fn code(self) -> i32 {
+        match self {
+            Self::Speech => 1,
+            Self::Vision => 2,
+        }
+    }
+}
+
+#[derive(Debug, Clone, PartialEq)]
+pub struct Phi4AudioOutput {
+    /// Pinned #874 projection stage, cropped to `valid_rows`, shape
+    /// `[1, valid_rows, hidden_size]`.
+    pub projected: Vec<f32>,
+    pub valid_rows: usize,
+    pub hidden_size: usize,
+    pub frame_bucket: usize,
+}
+
+#[must_use]
+pub fn phi4_audio_bucket_for_frames(frame_len: usize) -> Option<usize> {
+    PHI4MM_AUDIO_FRAME_BUCKETS
+        .iter()
+        .copied()
+        .find(|bucket| frame_len > 0 && frame_len <= *bucket)
+}
+
+fn config_identity(config: &Phi4AudioConfig, frame_bucket: usize, precision: Precision) -> String {
+    format!(
+        "family=phi4mm-audio;revision={PHI4MM_AUDIO_CHECKPOINT_REVISION};\
+         input={};dim={};heads={};blocks={};ff={};reduction={};channels={};\
+         kernel={};relative={};projection={};bucket={frame_bucket};precision={precision:?};\
+         projection-modes=speech,vision;abi=features-mask-length-mode-to-projection-length-v1",
+        config.input_size,
+        config.attention_dim,
+        config.attention_heads,
+        config.num_blocks,
+        config.linear_units,
+        config.time_reduction,
+        config.conv_channels,
+        config.kernel_size,
+        config.relative_bias_max_distance,
+        config.projection_hidden,
+    )
+}
+
+fn audio_precision() -> Result<Precision, String> {
+    match std::env::var("MLXCEL_PHI4MM_AUDIO_PRECISION").as_deref() {
+        Ok("f32") | Err(_) => Ok(Precision::F32),
+        Ok("bf16") => Ok(Precision::Bf16),
+        Ok(value) => Err(format!(
+            "unsupported MLXCEL_PHI4MM_AUDIO_PRECISION={value}; expected f32 or bf16"
+        )),
+    }
+}
+
+fn generation_identity(compiler: &Path, flags: &[&str], graph: &str) -> Result<String, String> {
+    let version = Command::new(compiler)
+        .arg("--version")
+        .output()
+        .map_err(|error| format!("run {} --version: {error}", compiler.display()))?;
+    if !version.status.success() {
+        return Err(format!(
+            "{} --version failed: {}",
+            compiler.display(),
+            String::from_utf8_lossy(&version.stderr)
+        ));
+    }
+    generation_identity_from_version(
+        compiler,
+        flags,
+        graph,
+        String::from_utf8_lossy(&version.stdout).trim(),
+    )
+}
+
+fn generation_identity_from_version(
+    compiler: &Path,
+    flags: &[&str],
+    graph: &str,
+    version: &str,
+) -> Result<String, String> {
+    Ok(format!(
+        "compiler={};compiler_sha256={};version={version};flags={flags:?};stablehlo_sha256={}",
+        compiler.display(),
+        sha256_file(compiler)?,
+        sha256_hex(graph.as_bytes()),
+    ))
+}
+
+fn sha256_hex(bytes: &[u8]) -> String {
+    let digest = Sha256::digest(bytes);
+    hex_bytes(&digest)
+}
+
+fn hex_bytes(bytes: &[u8]) -> String {
+    let mut output = String::with_capacity(bytes.len() * 2);
+    for byte in bytes {
+        use std::fmt::Write;
+        write!(output, "{byte:02x}").expect("writing to String cannot fail");
+    }
+    output
+}
+
+fn sha256_file(path: &Path) -> Result<String, String> {
+    let mut file = File::open(path).map_err(|error| format!("open {}: {error}", path.display()))?;
+    let mut digest = Sha256::new();
+    let mut buffer = [0u8; 64 * 1024];
+    loop {
+        let read = file
+            .read(&mut buffer)
+            .map_err(|error| format!("read {}: {error}", path.display()))?;
+        if read == 0 {
+            break;
+        }
+        digest.update(&buffer[..read]);
+    }
+    Ok(hex_bytes(&digest.finalize()))
+}
+
+fn validate_finite_values(label: &str, values: &[f32]) -> Result<(), String> {
+    if let Some((index, value)) = values
+        .iter()
+        .enumerate()
+        .find(|(_, value)| !value.is_finite())
+    {
+        return Err(format!(
+            "{label} contains non-finite value {value} at flat index {index}"
+        ));
+    }
+    Ok(())
+}
+
+fn resolve_weight_shards(model_dir: &Path, names: &[String]) -> Result<Vec<PathBuf>, String> {
+    let single = model_dir.join("model.safetensors");
+    if single.is_file() {
+        return Ok(vec![single; names.len()]);
+    }
+    let index = model_dir.join("model.safetensors.index.json");
+    let text = std::fs::read_to_string(&index)
+        .map_err(|error| format!("read {}: {error}", index.display()))?;
+    let value: serde_json::Value = serde_json::from_str(&text)
+        .map_err(|error| format!("parse {}: {error}", index.display()))?;
+    let map = value
+        .get("weight_map")
+        .and_then(serde_json::Value::as_object)
+        .ok_or_else(|| format!("{} is missing object `weight_map`", index.display()))?;
+    names
+        .iter()
+        .map(|name| {
+            map.get(name)
+                .and_then(serde_json::Value::as_str)
+                .map(|file| model_dir.join(file))
+                .ok_or_else(|| format!("{} has no entry for `{name}`", index.display()))
+        })
+        .collect()
+}
+
+fn f32_bytes(values: &[f32]) -> Vec<u8> {
+    values
+        .iter()
+        .flat_map(|value| value.to_ne_bytes())
+        .collect()
+}
+
+fn i32_bytes(values: &[i32]) -> Vec<u8> {
+    values
+        .iter()
+        .flat_map(|value| value.to_ne_bytes())
+        .collect()
+}
+
+fn decode_f32(bytes: &[u8]) -> Vec<f32> {
+    bytes
+        .chunks_exact(4)
+        .map(|chunk| f32::from_ne_bytes(chunk.try_into().expect("four-byte f32 chunk")))
+        .collect()
+}
+
+fn decode_i32_scalar(bytes: &[u8]) -> Result<i32, String> {
+    let bytes: [u8; 4] = bytes
+        .try_into()
+        .map_err(|_| format!("Phi4MM audio scalar output has {} bytes", bytes.len()))?;
+    Ok(i32::from_ne_bytes(bytes))
+}
+
+fn load_audio_weights(
+    model_dir: &Path,
+    config: &Phi4AudioConfig,
+) -> Result<Vec<AuxiliaryWeight>, String> {
+    let specs = phi4_audio_weight_specs(config);
+    let names = specs
+        .iter()
+        .map(|spec| spec.name.clone())
+        .collect::<Vec<_>>();
+    let shards = resolve_weight_shards(model_dir, &names)?;
+    let mut by_shard = BTreeMap::<&Path, Vec<usize>>::new();
+    for (index, shard) in shards.iter().enumerate() {
+        by_shard.entry(shard).or_default().push(index);
+    }
+    let mut loaded = (0..specs.len())
+        .map(|_| None)
+        .collect::<Vec<Option<AuxiliaryWeight>>>();
+    for (shard, indices) in by_shard {
+        let file =
+            File::open(shard).map_err(|error| format!("open {}: {error}", shard.display()))?;
+        // Safety: the mapping is read-only and cannot outlive `file` within
+        // this block; every selected tensor is copied before the mapping drops.
+        let mmap = unsafe { Mmap::map(&file) }
+            .map_err(|error| format!("mmap {}: {error}", shard.display()))?;
+        let tensors = SafeTensors::deserialize(&mmap)
+            .map_err(|error| format!("parse {}: {error}", shard.display()))?;
+        for index in indices {
+            let spec = &specs[index];
+            let tensor = tensors.tensor(&spec.name).map_err(|error| {
+                format!("read `{}` from {}: {error}", spec.name, shard.display())
+            })?;
+            if tensor.shape() != spec.shape {
+                return Err(format!(
+                    "Phi4MM audio weight `{}` has shape {:?}, expected {:?}",
+                    spec.name,
+                    tensor.shape(),
+                    spec.shape
+                ));
+            }
+            let values = match tensor.dtype() {
+                Dtype::BF16 => bf16_to_f32(tensor.data()),
+                Dtype::F16 => f16_to_f32(tensor.data()),
+                Dtype::F32 => tensor
+                    .data()
+                    .chunks_exact(4)
+                    .map(|chunk| {
+                        f32::from_le_bytes(chunk.try_into().expect("four-byte safetensors f32"))
+                    })
+                    .collect(),
+                dtype => {
+                    return Err(format!(
+                        "Phi4MM audio weight `{}` has unsupported dtype {dtype:?}",
+                        spec.name
+                    ));
+                }
+            };
+            let expected_elements = spec.shape.iter().product::<usize>();
+            if values.len() != expected_elements {
+                return Err(format!(
+                    "Phi4MM audio weight `{}` has {} elements, expected {expected_elements}",
+                    spec.name,
+                    values.len()
+                ));
+            }
+            validate_finite_values(&format!("Phi4MM audio weight `{}`", spec.name), &values)?;
+            loaded[index] = Some(AuxiliaryWeight {
+                name: spec.name.clone(),
+                bytes: f32_bytes(&values),
+                dtype: AuxiliaryWeightDType::Float32,
+                shape: spec.shape.clone(),
+            });
+        }
+    }
+    loaded
+        .into_iter()
+        .enumerate()
+        .map(|(index, weight)| {
+            weight
+                .ok_or_else(|| format!("Phi4MM audio weight {} was not loaded", specs[index].name))
+        })
+        .collect()
+}
+
+fn load_audio_module(
+    model_dir: &Path,
+    device: &str,
+    config: &Phi4AudioConfig,
+    frame_bucket: usize,
+    precision: Precision,
+    graph: &str,
+    artifact_name: &str,
+    entry_name: &str,
+) -> Result<IreeAuxiliaryModule, String> {
+    let compiler = iree_compile_bin()?;
+    let flags = target_flags(device)?;
+    let cache = std::env::temp_dir().join("mlxcel-xla-vmfb");
+    std::fs::create_dir_all(&cache)
+        .map_err(|error| format!("mkdir {}: {error}", cache.display()))?;
+    let vmfb = cached_vmfb_path(&compiler, graph, flags, &cache, artifact_name, frame_bucket);
+    let weights = load_audio_weights(model_dir, config)?;
+    let contract = AuxiliaryArtifactContract::new(
+        entry_name,
+        config_identity(config, frame_bucket, precision),
+        generation_identity(&compiler, flags, graph)?,
+    )?;
+    ensure_qualified_auxiliary_artifact(&vmfb, &contract, &weights, |temporary| {
+        compile_one_to(
+            &compiler,
+            graph,
+            flags,
+            &cache,
+            artifact_name,
+            frame_bucket,
+            temporary,
+        )
+    })?;
+    IreeAuxiliaryModule::load(device, &vmfb, &contract, weights)
+}
+
+pub struct Phi4AudioRuntime {
+    module: IreeAuxiliaryModule,
+    config: Phi4AudioConfig,
+    frame_bucket: usize,
+    encoded_bucket: usize,
+}
+
+impl Phi4AudioRuntime {
+    /// Compile/load one static audio bucket and upload all 887 encoder and
+    /// projection tensors as immutable resident buffers.
+    pub fn load(model_dir: &Path, device: &str, frame_bucket: usize) -> Result<Self, String> {
+        if !PHI4MM_AUDIO_FRAME_BUCKETS.contains(&frame_bucket) {
+            return Err(format!(
+                "unsupported Phi4MM audio frame bucket {frame_bucket}; expected one of {PHI4MM_AUDIO_FRAME_BUCKETS:?}"
+            ));
+        }
+        let config_text = std::fs::read_to_string(model_dir.join("config.json"))
+            .map_err(|error| format!("read Phi4MM config.json: {error}"))?;
+        let config = Phi4AudioConfig::from_json_str(&config_text)?;
+        let encoded_bucket = config.encoded_bucket_len(frame_bucket)?;
+        // Mirror #874's real mixed boundary: BF16 feature normalization,
+        // followed by F32 Conformer and projection stages.
+        let precision = audio_precision()?;
+        let graph = emit_phi4_audio_with(&config, frame_bucket, precision)?;
+        let module = load_audio_module(
+            model_dir,
+            device,
+            &config,
+            frame_bucket,
+            precision,
+            &graph,
+            "phi4mm-audio",
+            "audio.main",
+        )?;
+        Ok(Self {
+            module,
+            config,
+            frame_bucket,
+            encoded_bucket,
+        })
+    }
+
+    #[must_use]
+    pub fn frame_bucket(&self) -> usize {
+        self.frame_bucket
+    }
+
+    #[must_use]
+    pub fn fingerprint(&self) -> u64 {
+        self.module.fingerprint()
+    }
+
+    pub fn project(
+        &mut self,
+        features: &[f32],
+        frame_len: usize,
+        mode: Phi4AudioProjectionMode,
+    ) -> Result<Phi4AudioOutput, String> {
+        if frame_len == 0 || frame_len > self.frame_bucket {
+            return Err(format!(
+                "Phi4MM audio frame length {frame_len} is outside static bucket 1..={}",
+                self.frame_bucket
+            ));
+        }
+        let expected = frame_len
+            .checked_mul(self.config.input_size)
+            .ok_or_else(|| "Phi4MM audio feature element count overflowed".to_string())?;
+        if features.len() != expected {
+            return Err(format!(
+                "Phi4MM audio features have {} elements, expected {expected} for [{frame_len}, {}]",
+                features.len(),
+                self.config.input_size
+            ));
+        }
+        if features.iter().any(|value| !value.is_finite()) {
+            return Err("Phi4MM audio features contain non-finite values".to_string());
+        }
+        let mut padded = vec![0.0f32; self.frame_bucket * self.config.input_size];
+        padded[..features.len()].copy_from_slice(features);
+        let mut mask = vec![0i32; self.frame_bucket];
+        mask[..frame_len].fill(1);
+        let feature_bytes = f32_bytes(&padded);
+        let mask_bytes = i32_bytes(&mask);
+        let length_bytes = i32_bytes(&[i32::try_from(frame_len)
+            .map_err(|_| format!("Phi4MM audio frame length {frame_len} does not fit i32"))?]);
+        let mode_bytes = i32_bytes(&[mode.code()]);
+        let feature_shape = [1, self.frame_bucket, self.config.input_size];
+        let mask_shape = [1, self.frame_bucket];
+        let scalar_shape: [usize; 0] = [];
+        let output_shape = [1, self.encoded_bucket, self.config.projection_hidden];
+        let mut projected_bytes =
+            vec![0u8; self.encoded_bucket * self.config.projection_hidden * size_of::<f32>()];
+        let mut output_length_bytes = vec![0u8; size_of::<i32>()];
+        self.module.invoke(
+            &[
+                AuxiliaryInput {
+                    bytes: &feature_bytes,
+                    dtype: AuxiliaryTensorDType::Float32,
+                    shape: &feature_shape,
+                },
+                AuxiliaryInput {
+                    bytes: &mask_bytes,
+                    dtype: AuxiliaryTensorDType::Int32,
+                    shape: &mask_shape,
+                },
+                AuxiliaryInput {
+                    bytes: &length_bytes,
+                    dtype: AuxiliaryTensorDType::Int32,
+                    shape: &scalar_shape,
+                },
+                AuxiliaryInput {
+                    bytes: &mode_bytes,
+                    dtype: AuxiliaryTensorDType::Int32,
+                    shape: &scalar_shape,
+                },
+            ],
+            &mut [
+                AuxiliaryOutput {
+                    bytes: &mut projected_bytes,
+                    dtype: AuxiliaryTensorDType::Float32,
+                    shape: &output_shape,
+                },
+                AuxiliaryOutput {
+                    bytes: &mut output_length_bytes,
+                    dtype: AuxiliaryTensorDType::Int32,
+                    shape: &scalar_shape,
+                },
+            ],
+        )?;
+        let valid_rows = usize::try_from(decode_i32_scalar(&output_length_bytes)?)
+            .map_err(|_| "Phi4MM audio graph returned a negative valid length".to_string())?;
+        let expected_rows = self.config.encoded_valid_len(frame_len)?;
+        if valid_rows != expected_rows || valid_rows > self.encoded_bucket {
+            return Err(format!(
+                "Phi4MM audio graph returned valid length {valid_rows}, expected {expected_rows} within bucket {}",
+                self.encoded_bucket
+            ));
+        }
+        let mut projected = decode_f32(&projected_bytes);
+        projected.truncate(valid_rows * self.config.projection_hidden);
+        if projected.iter().any(|value| !value.is_finite()) {
+            return Err("Phi4MM audio graph returned non-finite projection values".to_string());
+        }
+        Ok(Phi4AudioOutput {
+            projected,
+            valid_rows,
+            hidden_size: self.config.projection_hidden,
+            frame_bucket: self.frame_bucket,
+        })
+    }
+}
+
+/// One exact tensor boundary returned by the real-checkpoint diagnostic entry.
+///
+/// This API exists only to localize parity failures; production requests use
+/// [`Phi4AudioRuntime`] and never allocate these intermediate tensors.
+#[doc(hidden)]
+#[derive(Debug, Clone, PartialEq)]
+pub struct Phi4AudioCheckpoint {
+    pub name: &'static str,
+    pub shape: Vec<usize>,
+    pub values: Vec<f32>,
+}
+
+#[doc(hidden)]
+#[derive(Debug, Clone, PartialEq)]
+pub struct Phi4AudioDiagnostics {
+    pub checkpoints: Vec<Phi4AudioCheckpoint>,
+    pub valid_rows: usize,
+}
+
+#[doc(hidden)]
+pub struct Phi4AudioDiagnosticRuntime {
+    module: IreeAuxiliaryModule,
+    config: Phi4AudioConfig,
+    frame_bucket: usize,
+    encoded_bucket: usize,
+}
+
+#[doc(hidden)]
+impl Phi4AudioDiagnosticRuntime {
+    pub fn load(model_dir: &Path, device: &str, frame_bucket: usize) -> Result<Self, String> {
+        if !PHI4MM_AUDIO_FRAME_BUCKETS.contains(&frame_bucket) {
+            return Err(format!(
+                "unsupported Phi4MM audio frame bucket {frame_bucket}; expected one of {PHI4MM_AUDIO_FRAME_BUCKETS:?}"
+            ));
+        }
+        let config_text = std::fs::read_to_string(model_dir.join("config.json"))
+            .map_err(|error| format!("read Phi4MM config.json: {error}"))?;
+        let config = Phi4AudioConfig::from_json_str(&config_text)?;
+        let encoded_bucket = config.encoded_bucket_len(frame_bucket)?;
+        let precision = audio_precision()?;
+        let graph = emit_phi4_audio_diagnostic_with(&config, frame_bucket, precision)?;
+        let module = load_audio_module(
+            model_dir,
+            device,
+            &config,
+            frame_bucket,
+            precision,
+            &graph,
+            "phi4mm-audio-diagnostic",
+            "audio.diagnostic",
+        )?;
+        Ok(Self {
+            module,
+            config,
+            frame_bucket,
+            encoded_bucket,
+        })
+    }
+
+    pub fn capture(
+        &mut self,
+        features: &[f32],
+        frame_len: usize,
+    ) -> Result<Phi4AudioDiagnostics, String> {
+        if frame_len == 0 || frame_len > self.frame_bucket {
+            return Err(format!(
+                "Phi4MM audio frame length {frame_len} is outside static bucket 1..={}",
+                self.frame_bucket
+            ));
+        }
+        let expected = frame_len
+            .checked_mul(self.config.input_size)
+            .ok_or_else(|| "Phi4MM audio feature element count overflowed".to_string())?;
+        if features.len() != expected {
+            return Err(format!(
+                "Phi4MM audio features have {} elements, expected {expected} for [{frame_len}, {}]",
+                features.len(),
+                self.config.input_size
+            ));
+        }
+        if features.iter().any(|value| !value.is_finite()) {
+            return Err("Phi4MM audio features contain non-finite values".to_string());
+        }
+        let mut padded = vec![0.0f32; self.frame_bucket * self.config.input_size];
+        padded[..features.len()].copy_from_slice(features);
+        let mut mask = vec![0i32; self.frame_bucket];
+        mask[..frame_len].fill(1);
+        let feature_bytes = f32_bytes(&padded);
+        let mask_bytes = i32_bytes(&mask);
+        let length_bytes = i32_bytes(&[i32::try_from(frame_len)
+            .map_err(|_| format!("Phi4MM audio frame length {frame_len} does not fit i32"))?]);
+        // The diagnostic returns both projection branches; this mode input is
+        // still supplied to keep its argument ABI identical to audio.main.
+        let mode_bytes = i32_bytes(&[Phi4AudioProjectionMode::Speech.code()]);
+        let feature_shape = [1, self.frame_bucket, self.config.input_size];
+        let mask_shape = [1, self.frame_bucket];
+        let scalar_shape: [usize; 0] = [];
+        let specs = phi4_audio_diagnostic_specs(&self.config, self.frame_bucket)?;
+        let mut buffers = specs
+            .iter()
+            .map(|spec| {
+                spec.shape
+                    .iter()
+                    .try_fold(size_of::<f32>(), |bytes: usize, dimension| {
+                        bytes.checked_mul(*dimension).ok_or_else(|| {
+                            format!(
+                                "Phi4MM audio diagnostic {} byte count overflowed",
+                                spec.name
+                            )
+                        })
+                    })
+                    .map(|bytes| vec![0u8; bytes])
+            })
+            .collect::<Result<Vec<_>, String>>()?;
+        let mut output_length_bytes = vec![0u8; size_of::<i32>()];
+        let mut outputs = buffers
+            .iter_mut()
+            .zip(&specs)
+            .map(|(bytes, spec)| AuxiliaryOutput {
+                bytes,
+                dtype: AuxiliaryTensorDType::Float32,
+                shape: &spec.shape,
+            })
+            .collect::<Vec<_>>();
+        outputs.push(AuxiliaryOutput {
+            bytes: &mut output_length_bytes,
+            dtype: AuxiliaryTensorDType::Int32,
+            shape: &scalar_shape,
+        });
+        self.module.invoke(
+            &[
+                AuxiliaryInput {
+                    bytes: &feature_bytes,
+                    dtype: AuxiliaryTensorDType::Float32,
+                    shape: &feature_shape,
+                },
+                AuxiliaryInput {
+                    bytes: &mask_bytes,
+                    dtype: AuxiliaryTensorDType::Int32,
+                    shape: &mask_shape,
+                },
+                AuxiliaryInput {
+                    bytes: &length_bytes,
+                    dtype: AuxiliaryTensorDType::Int32,
+                    shape: &scalar_shape,
+                },
+                AuxiliaryInput {
+                    bytes: &mode_bytes,
+                    dtype: AuxiliaryTensorDType::Int32,
+                    shape: &scalar_shape,
+                },
+            ],
+            &mut outputs,
+        )?;
+        drop(outputs);
+        let valid_rows = usize::try_from(decode_i32_scalar(&output_length_bytes)?)
+            .map_err(|_| "Phi4MM audio diagnostic returned a negative valid length".to_string())?;
+        let expected_rows = self.config.encoded_valid_len(frame_len)?;
+        if valid_rows != expected_rows || valid_rows > self.encoded_bucket {
+            return Err(format!(
+                "Phi4MM audio diagnostic returned valid length {valid_rows}, expected {expected_rows} within bucket {}",
+                self.encoded_bucket
+            ));
+        }
+        let checkpoints = specs
+            .into_iter()
+            .zip(buffers)
+            .map(|(spec, bytes)| {
+                let values = decode_f32(&bytes);
+                if values.iter().any(|value| !value.is_finite()) {
+                    return Err(format!(
+                        "Phi4MM audio diagnostic {} contains non-finite values",
+                        spec.name
+                    ));
+                }
+                Ok(Phi4AudioCheckpoint {
+                    name: spec.name,
+                    shape: spec.shape,
+                    values,
+                })
+            })
+            .collect::<Result<Vec<_>, String>>()?;
+        Ok(Phi4AudioDiagnostics {
+            checkpoints,
+            valid_rows,
+        })
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use crate::aux_manifest::{auxiliary_manifest_path, verify_auxiliary_manifest};
+
+    fn published_config() -> Phi4AudioConfig {
+        Phi4AudioConfig {
+            input_size: 80,
+            attention_dim: 1024,
+            attention_heads: 16,
+            num_blocks: 24,
+            linear_units: 1536,
+            time_reduction: 8,
+            conv_channels: 1024,
+            kernel_size: 3,
+            relative_bias_max_distance: 500,
+            projection_hidden: 3072,
+        }
+    }
+
+    fn temporary_audio_path(suffix: &str) -> PathBuf {
+        let nonce = std::time::SystemTime::now()
+            .duration_since(std::time::UNIX_EPOCH)
+            .expect("system clock must be after the Unix epoch")
+            .as_nanos();
+        std::env::temp_dir().join(format!(
+            "mlxcel-phi4mm-audio-cache-test-{}-{nonce}.{suffix}",
+            std::process::id(),
+        ))
+    }
+
+    #[test]
+    fn bucket_selection_is_static_and_fail_closed() {
+        assert_eq!(phi4_audio_bucket_for_frames(0), None);
+        assert_eq!(phi4_audio_bucket_for_frames(1), Some(64));
+        assert_eq!(phi4_audio_bucket_for_frames(64), Some(64));
+        assert_eq!(phi4_audio_bucket_for_frames(65), Some(128));
+        assert_eq!(phi4_audio_bucket_for_frames(351), Some(512));
+        assert_eq!(phi4_audio_bucket_for_frames(4000), Some(4000));
+        assert_eq!(phi4_audio_bucket_for_frames(4001), None);
+    }
+
+    #[test]
+    fn projection_modes_match_the_pinned_abi_codes() {
+        assert_eq!(Phi4AudioProjectionMode::Speech.code(), 1);
+        assert_eq!(Phi4AudioProjectionMode::Vision.code(), 2);
+    }
+
+    #[test]
+    fn audio_contract_reuses_qualified_cache_and_rebuilds_stale_pair_once() {
+        let vmfb = temporary_audio_path("vmfb");
+        let weights = vec![AuxiliaryWeight {
+            name: "model.embed_tokens_extend.audio_embed.weight".to_string(),
+            bytes: 1.0f32.to_ne_bytes().to_vec(),
+            dtype: AuxiliaryWeightDType::Float32,
+            shape: vec![1],
+        }];
+        let config = published_config();
+        let contract = AuxiliaryArtifactContract::new(
+            "audio.main",
+            config_identity(&config, 512, Precision::F32),
+            "compiler=test;flags=cpu;stablehlo_sha256=v1",
+        )
+        .unwrap();
+        let mut compile_count = 0usize;
+        ensure_qualified_auxiliary_artifact(&vmfb, &contract, &weights, |temporary| {
+            compile_count += 1;
+            std::fs::write(temporary, b"audio-vmfb-v1")
+                .map_err(|error| format!("write test audio VMFB: {error}"))
+        })
+        .unwrap();
+        assert_eq!(compile_count, 1);
+
+        ensure_qualified_auxiliary_artifact(&vmfb, &contract, &weights, |_| {
+            compile_count += 1;
+            Err("qualified Phi4MM audio cache must not compile".to_string())
+        })
+        .unwrap();
+        assert_eq!(compile_count, 1);
+
+        let stale_replacement = AuxiliaryArtifactContract::new(
+            "audio.main",
+            config_identity(&config, 512, Precision::F32),
+            "compiler=test;flags=cpu;stablehlo_sha256=v2",
+        )
+        .unwrap();
+        ensure_qualified_auxiliary_artifact(&vmfb, &stale_replacement, &weights, |temporary| {
+            compile_count += 1;
+            std::fs::write(temporary, b"audio-vmfb-v2")
+                .map_err(|error| format!("write replacement audio VMFB: {error}"))
+        })
+        .unwrap();
+        assert_eq!(compile_count, 2);
+        assert_eq!(std::fs::read(&vmfb).unwrap(), b"audio-vmfb-v2");
+        verify_auxiliary_manifest(&vmfb, &stale_replacement, &weights).unwrap();
+
+        std::fs::remove_file(auxiliary_manifest_path(&vmfb)).ok();
+        std::fs::remove_file(vmfb).ok();
+    }
+
+    #[test]
+    fn compiler_binary_replacement_at_same_path_and_version_rebuilds_audio_cache() {
+        let compiler = temporary_audio_path("compiler");
+        let vmfb = temporary_audio_path("compiler-digest.vmfb");
+        let weights = vec![AuxiliaryWeight {
+            name: "weight".to_string(),
+            bytes: 1.0f32.to_ne_bytes().to_vec(),
+            dtype: AuxiliaryWeightDType::Float32,
+            shape: vec![1],
+        }];
+        std::fs::write(&compiler, b"compiler-build-a").unwrap();
+        let first_generation =
+            generation_identity_from_version(&compiler, &["--target=cpu"], "mlir", "v1").unwrap();
+        assert!(
+            first_generation.contains(
+                "compiler_sha256=c998e735c573e64551765a30b12a0cc63d1255c2229f5943995ccdfef4939b7a"
+            ),
+            "generation identity should contain the compiler executable's raw SHA-256 digest"
+        );
+        let first_contract =
+            AuxiliaryArtifactContract::new("audio.main", "config=v1", &first_generation).unwrap();
+        let mut compile_count = 0usize;
+        ensure_qualified_auxiliary_artifact(&vmfb, &first_contract, &weights, |temporary| {
+            compile_count += 1;
+            std::fs::write(temporary, b"vmfb-a")
+                .map_err(|error| format!("write test VMFB: {error}"))
+        })
+        .unwrap();
+
+        std::fs::write(&compiler, b"compiler-build-b").unwrap();
+        let second_generation =
+            generation_identity_from_version(&compiler, &["--target=cpu"], "mlir", "v1").unwrap();
+        assert_ne!(first_generation, second_generation);
+        let second_contract =
+            AuxiliaryArtifactContract::new("audio.main", "config=v1", second_generation).unwrap();
+        ensure_qualified_auxiliary_artifact(&vmfb, &second_contract, &weights, |temporary| {
+            compile_count += 1;
+            std::fs::write(temporary, b"vmfb-b")
+                .map_err(|error| format!("write test VMFB: {error}"))
+        })
+        .unwrap();
+
+        assert_eq!(compile_count, 2);
+        assert_eq!(std::fs::read(&vmfb).unwrap(), b"vmfb-b");
+        std::fs::remove_file(auxiliary_manifest_path(&vmfb)).ok();
+        std::fs::remove_file(vmfb).ok();
+        std::fs::remove_file(compiler).ok();
+    }
+
+    #[test]
+    fn audio_weight_values_reject_non_finite_before_native_upload() {
+        assert!(validate_finite_values("audio weight", &[0.0, -1.0, 3.0]).is_ok());
+        for value in [f32::NAN, f32::INFINITY, f32::NEG_INFINITY] {
+            let error = validate_finite_values("audio weight", &[0.0, value]).unwrap_err();
+            assert!(error.contains("audio weight"));
+            assert!(error.contains("flat index 1"));
+        }
+    }
+}
diff --git a/src/lib/mlxcel-xla/src/prepared.rs b/src/lib/mlxcel-xla/src/prepared.rs
index c393d9ee2..bc1b4eaad 100644
--- a/src/lib/mlxcel-xla/src/prepared.rs
+++ b/src/lib/mlxcel-xla/src/prepared.rs
@@ -80,6 +80,7 @@ impl PreparedIreePositions {
 #[derive(Debug)]
 pub(crate) struct PreparedIreePrefill {
     pub(crate) token_ids: Vec<i32>,
+    pub(crate) adapter_mode: i32,
     pub(crate) embeddings: Vec<f32>,
     pub(crate) positions: PreparedIreePositions,
     pub(crate) attention_bias: Vec<f32>,
@@ -581,6 +582,7 @@ impl PreparedIreePrefill {
 
         Ok(Self {
             token_ids: value.token_ids.clone(),
+            adapter_mode: value.adapter_mode.code(),
             embeddings,
             positions,
             attention_bias,
diff --git a/src/lib/mlxcel-xla/src/weights.rs b/src/lib/mlxcel-xla/src/weights.rs
index cefaf588c..22338ab7e 100644
--- a/src/lib/mlxcel-xla/src/weights.rs
+++ b/src/lib/mlxcel-xla/src/weights.rs
@@ -125,6 +125,21 @@ fn push_proj(out: &mut Vec<WeightSpec>, name: String, quant: bool) {
     }
 }
 
+fn phi4_base_projection(cfg: &Config, name: String) -> String {
+    if cfg.phi4_lora.is_some() {
+        name.strip_suffix(".weight")
+            .map(|stem| format!("{stem}.base_layer.weight"))
+            .unwrap_or(name)
+    } else {
+        name
+    }
+}
+
+fn push_phi4_lora_pair(out: &mut Vec<WeightSpec>, stem: &str, adapter: &str) {
+    out.push(WeightSpec::Proj(format!("{stem}.lora_A.{adapter}.weight")));
+    out.push(WeightSpec::Proj(format!("{stem}.lora_B.{adapter}.weight")));
+}
+
 /// [`weight_specs`] with the packed-path decision passed explicitly (so it is
 /// testable without the `MLXCEL_XLA_QUANT` env). `quant` already folds in
 /// [`Config::supports_packed_quant`], so it is `true` only for the standard Llama
@@ -160,18 +175,22 @@ fn weight_specs_q(cfg: &Config, quant: bool) -> Vec<WeightSpec> {
             } else {
                 format!("{p}{}", s.down)
             };
-            push_proj(&mut out, name, quant);
+            push_proj(&mut out, phi4_base_projection(cfg, name), quant);
         }
         // gate (gated MLP only; the first half of gate_up_proj for a fused Phi3).
         if gated && !moe_layer {
             if cfg.fused_gate_up {
                 out.push(WeightSpec::Rows {
-                    name: format!("{p}mlp.gate_up_proj.weight"),
+                    name: phi4_base_projection(cfg, format!("{p}mlp.gate_up_proj.weight")),
                     start: 0,
                     end: inter,
                 });
             } else {
-                push_proj(&mut out, format!("{p}{}", s.gate), quant);
+                push_proj(
+                    &mut out,
+                    phi4_base_projection(cfg, format!("{p}{}", s.gate)),
+                    quant,
+                );
             }
         }
         // input_layernorm: present unless the OLMo reordered post-norm drops it.
@@ -190,25 +209,32 @@ fn weight_specs_q(cfg: &Config, quant: bool) -> Vec<WeightSpec> {
         if !moe_layer {
             if cfg.fused_gate_up {
                 out.push(WeightSpec::Rows {
-                    name: format!("{p}mlp.gate_up_proj.weight"),
+                    name: phi4_base_projection(cfg, format!("{p}mlp.gate_up_proj.weight")),
                     start: inter,
                     end: 2 * inter,
                 });
             } else if cfg.dense_mlp {
                 out.push(WeightSpec::Whole(format!("{p}mlp.c_fc.weight")));
             } else {
-                push_proj(&mut out, format!("{p}{}", s.up), quant);
+                push_proj(
+                    &mut out,
+                    phi4_base_projection(cfg, format!("{p}{}", s.up)),
+                    quant,
+                );
             }
         }
         // wk, wo, wq, wv (JAX-alphabetical; a fused Phi3 qkv_proj is [Q|K|V] rows).
         if cfg.fused_qkv {
-            let qkv = format!("{p}self_attn.qkv_proj.weight");
+            let qkv = phi4_base_projection(cfg, format!("{p}self_attn.qkv_proj.weight"));
             out.push(WeightSpec::Rows {
                 name: qkv.clone(),
                 start: nq,
                 end: nq + nkv,
             }); // wk
-            out.push(WeightSpec::Whole(format!("{p}{}", s.o_proj))); // wo
+            out.push(WeightSpec::Proj(phi4_base_projection(
+                cfg,
+                format!("{p}{}", s.o_proj),
+            ))); // wo
             out.push(WeightSpec::Rows {
                 name: qkv.clone(),
                 start: 0,
@@ -220,10 +246,26 @@ fn weight_specs_q(cfg: &Config, quant: bool) -> Vec<WeightSpec> {
                 end: nq + 2 * nkv,
             }); // wv
         } else {
-            push_proj(&mut out, format!("{p}{}", s.k_proj), quant);
-            push_proj(&mut out, format!("{p}{}", s.o_proj), quant);
-            push_proj(&mut out, format!("{p}{}", s.q_proj), quant);
-            push_proj(&mut out, format!("{p}{}", s.v_proj), quant);
+            push_proj(
+                &mut out,
+                phi4_base_projection(cfg, format!("{p}{}", s.k_proj)),
+                quant,
+            );
+            push_proj(
+                &mut out,
+                phi4_base_projection(cfg, format!("{p}{}", s.o_proj)),
+                quant,
+            );
+            push_proj(
+                &mut out,
+                phi4_base_projection(cfg, format!("{p}{}", s.q_proj)),
+                quant,
+            );
+            push_proj(
+                &mut out,
+                phi4_base_projection(cfg, format!("{p}{}", s.v_proj)),
+                quant,
+            );
         }
         // q/k/v projection biases (k, q, v order).
         if cfg.qkv_bias {
@@ -305,6 +347,17 @@ fn weight_specs_q(cfg: &Config, quant: bool) -> Vec<WeightSpec> {
                 }
             }
         }
+        if cfg.phi4_lora.is_some() {
+            for stem in [
+                format!("{p}self_attn.qkv_proj"),
+                format!("{p}self_attn.o_proj"),
+                format!("{p}mlp.gate_up_proj"),
+                format!("{p}mlp.down_proj"),
+            ] {
+                push_phi4_lora_pair(&mut out, &stem, "speech");
+                push_phi4_lora_pair(&mut out, &stem, "vision");
+            }
+        }
     }
     out
 }
@@ -686,6 +739,110 @@ pub(crate) fn dequantize_affine_stacked(
 #[cfg(test)]
 mod tests {
     use super::*;
+    use crate::emitter::Phi4LoraConfig;
+
+    fn phi4_config() -> Config {
+        Config {
+            hidden: 8,
+            inter: 16,
+            n_layers: 2,
+            n_q: 2,
+            n_kv: 1,
+            head_dim: 4,
+            rotary_dim: Some(4),
+            vocab: 32,
+            tie_word_embeddings: true,
+            fused_qkv: true,
+            fused_gate_up: true,
+            phi4_lora: Some(Phi4LoraConfig {
+                speech_rank: 3,
+                speech_scale: 2.0,
+                vision_rank: 2,
+                vision_scale: 2.5,
+            }),
+            ..Config::llama_3_2_1b()
+        }
+    }
+
+    #[test]
+    fn phi4mm_weight_schema_uses_base_layers_and_complete_adapter_pairs() {
+        let config = phi4_config();
+        let specs = weight_specs_q(&config, false);
+        let names = specs
+            .iter()
+            .map(WeightSpec::tensor_name)
+            .collect::<Vec<_>>();
+        for layer in 0..config.n_layers {
+            let prefix = format!("model.layers.{layer}");
+            for projection in [
+                "self_attn.qkv_proj",
+                "self_attn.o_proj",
+                "mlp.gate_up_proj",
+                "mlp.down_proj",
+            ] {
+                let stem = format!("{prefix}.{projection}");
+                assert!(
+                    names.contains(&format!("{stem}.base_layer.weight").as_str()),
+                    "missing immutable base projection {stem}"
+                );
+                for adapter in ["speech", "vision"] {
+                    assert!(names.contains(&format!("{stem}.lora_A.{adapter}.weight").as_str()));
+                    assert!(names.contains(&format!("{stem}.lora_B.{adapter}.weight").as_str()));
+                }
+            }
+        }
+        assert_eq!(
+            names.iter().filter(|name| name.contains(".lora_")).count(),
+            config.n_layers * 4 * 2 * 2,
+            "four projections x two adapters x complete A/B pairs per layer"
+        );
+        assert!(
+            names
+                .iter()
+                .filter(|name| {
+                    name.contains("qkv_proj")
+                        || name.contains("o_proj")
+                        || name.contains("gate_up_proj")
+                        || name.contains("down_proj")
+                })
+                .all(|name| name.contains(".base_layer.") || name.contains(".lora_")),
+            "Phi4MM must never resolve a mutable unwrapped projection"
+        );
+    }
+
+    #[test]
+    #[ignore = "requires PHI4MM_MODEL_DIR pointing at the pinned official checkpoint"]
+    fn real_phi4mm_index_covers_decoder_and_complete_lora_schema() {
+        let model_dir = std::path::PathBuf::from(
+            std::env::var("PHI4MM_MODEL_DIR").expect("set PHI4MM_MODEL_DIR"),
+        );
+        let config = Config::from_json(&model_dir).expect("parse pinned Phi4MM config");
+        let lora = config.phi4_lora.expect("Phi4MM LoRA contract");
+        assert_eq!((lora.speech_rank, lora.vision_rank), (320, 256));
+        assert_eq!((lora.speech_scale, lora.vision_scale), (2.0, 2.0));
+
+        let index_path = model_dir.join("model.safetensors.index.json");
+        let index: serde_json::Value = serde_json::from_str(
+            &std::fs::read_to_string(&index_path).expect("read checkpoint index"),
+        )
+        .expect("parse checkpoint index");
+        let weight_map = index["weight_map"].as_object().expect("weight_map object");
+        let specs = weight_specs_q(&config, false);
+        let missing = specs
+            .iter()
+            .map(WeightSpec::tensor_name)
+            .filter(|name| !weight_map.contains_key(*name))
+            .collect::<Vec<_>>();
+        assert!(missing.is_empty(), "missing graph weights: {missing:?}");
+        assert_eq!(
+            specs
+                .iter()
+                .filter(|spec| spec.tensor_name().contains(".lora_"))
+                .count(),
+            512,
+            "32 layers x 4 projections x 2 modes x A/B"
+        );
+    }
 
     /// The Llama family weight order is embed + norm (f32-resident `Whole`) then, per
     /// layer, the 7 linear projections as `Proj` (f16-resident-capable, issue #572)
diff --git a/src/loading/mod.rs b/src/loading/mod.rs
index 2615bb32a..288f1f02c 100644
--- a/src/loading/mod.rs
+++ b/src/loading/mod.rs
@@ -60,6 +60,10 @@ pub(crate) use self::vlm::load_llava_host_preprocessor;
 #[cfg(feature = "xla-iree")]
 pub(crate) use self::vlm::load_llava_iree_host_preprocessor;
 pub use self::vlm::load_qwen3_omni_speech;
+#[cfg(feature = "xla-iree")]
+pub(crate) use self::vlm::{
+    Phi4MMXlaVisionComponents, load_phi4mm_xla_media_components, load_phi4mm_xla_text_embeddings,
+};
 
 /// Resolve model path: if a file is given, use its parent directory.
 ///
diff --git a/src/loading/vlm.rs b/src/loading/vlm.rs
index 234b538de..e026b054b 100644
--- a/src/loading/vlm.rs
+++ b/src/loading/vlm.rs
@@ -123,6 +123,10 @@ pub(crate) use qwen::{
 pub use qwen::load_qwen3_omni_speech;
 pub(crate) use siglip::{load_aya_vision_vlm, load_paligemma_vlm};
 pub(crate) use smolvlm::load_smolvlm_vlm;
+#[cfg(feature = "xla-iree")]
+pub(crate) use special::{
+    Phi4MMXlaVisionComponents, load_phi4mm_xla_media_components, load_phi4mm_xla_text_embeddings,
+};
 pub(crate) use special::{
     load_llama4_vlm, load_minicpmo_vlm, load_minicpmv4_6_vlm, load_molmo_point_vlm, load_molmo_vlm,
     load_molmo2_vlm, load_moondream2_vlm, load_moondream3_vlm, load_phi3_vlm, load_phi4_siglip_vlm,
diff --git a/src/loading/vlm_special.rs b/src/loading/vlm_special.rs
index f93fe9ed1..ea8353ddb 100644
--- a/src/loading/vlm_special.rs
+++ b/src/loading/vlm_special.rs
@@ -1327,6 +1327,171 @@ pub(crate) fn load_phi4mm_vlm(model_path: &Path) -> Result<LoadedModel> {
     }))
 }
 
+/// Load only Phi4MM's canonical text embedding table for the XLA audio
+/// producer. Decoder layers, LM head, vision tower, and MLX audio encoder are
+/// deliberately excluded.
+#[cfg(feature = "xla-iree")]
+pub(crate) fn load_phi4mm_xla_text_embeddings(
+    model_path: &Path,
+) -> Result<(mlxcel_core::layers::UnifiedEmbedding, usize, usize)> {
+    let (_config_str, full_config) = read_sanitized_vlm_config(model_path)?;
+    if full_config.get("model_type").and_then(Value::as_str) != Some("phi4mm") {
+        return Err(anyhow::anyhow!(
+            "{} is not a Phi4MM checkpoint",
+            model_path.display()
+        ));
+    }
+    let mut text_config_value = phi4mm_text_config_value(&full_config)?;
+    inherit_quantization_if_missing(&mut text_config_value, &full_config)?;
+    let text_config: models::phi4mm::ModelArgs = serde_json::from_value(text_config_value)
+        .map_err(|error| anyhow::anyhow!("Failed to parse Phi4MM text config: {error}"))?;
+    let weights = super::load_vlm_weights_common_filtered_canonical(model_path, |name| {
+        name.starts_with("model.embed_tokens.")
+    })?;
+    let embeddings = mlxcel_core::layers::UnifiedEmbedding::from_weights(
+        &weights,
+        "model.embed_tokens",
+        text_config.group_size(),
+        text_config.bits(),
+    )
+    .map_err(|error| anyhow::anyhow!("invalid Phi4MM text embedding table: {error}"))?;
+    Ok((
+        embeddings,
+        text_config.hidden_size,
+        text_config.max_position_embeddings,
+    ))
+}
+
+/// Filtered Phi4MM host components used by the XLA prepared-prefill producer.
+///
+/// The struct deliberately contains no text decoder, LM head, audio encoder, or
+/// decoder LoRA tensors. The latter are owned by the IREE runtime bundle.
+#[cfg(feature = "xla-iree")]
+pub(crate) struct Phi4MMXlaVisionComponents {
+    pub vision_tower: vision::encoders::phi4_siglip::Phi4SigLipVisionEncoder,
+    pub mm_projector_linear1: mlxcel_core::layers::UnifiedLinear,
+    pub mm_projector_linear2: mlxcel_core::layers::UnifiedLinear,
+    pub processor: vision::processors::phi4mm::Phi4MMProcessor,
+    pub select_layer: isize,
+    pub glb_gn: mlxcel_core::UniquePtr<mlxcel_core::MlxArray>,
+    pub sub_gn: mlxcel_core::UniquePtr<mlxcel_core::MlxArray>,
+    pub hd_transform_order: String,
+}
+
+/// Load the canonical Phi4MM embedding lookup and image encoder/projector
+/// without constructing any MLX language or audio inference layers.
+#[cfg(feature = "xla-iree")]
+pub(crate) fn load_phi4mm_xla_media_components(
+    model_path: &Path,
+) -> Result<(
+    mlxcel_core::layers::UnifiedEmbedding,
+    usize,
+    usize,
+    Phi4MMXlaVisionComponents,
+)> {
+    use vision::encoders::phi4_siglip::{Phi4SigLipVisionConfig, Phi4SigLipVisionEncoder};
+    use vision::processors::phi4mm::Phi4MMProcessor;
+
+    let (_config_str, full_config) = read_sanitized_vlm_config(model_path)?;
+    if full_config.get("model_type").and_then(Value::as_str) != Some("phi4mm") {
+        return Err(anyhow::anyhow!(
+            "{} is not a Phi4MM checkpoint",
+            model_path.display()
+        ));
+    }
+    let mut text_config_value = phi4mm_text_config_value(&full_config)?;
+    inherit_quantization_if_missing(&mut text_config_value, &full_config)?;
+    let text_config: models::phi4mm::ModelArgs = serde_json::from_value(text_config_value)
+        .map_err(|error| anyhow::anyhow!("Failed to parse Phi4MM text config: {error}"))?;
+    let vision_config: Phi4SigLipVisionConfig =
+        serde_json::from_value(phi4mm_vision_config_value(&full_config))
+            .map_err(|error| anyhow::anyhow!("Failed to parse Phi4MM vision config: {error}"))?;
+
+    let image_prefix = "model.embed_tokens_extend.image_embed.";
+    let raw_weights = super::load_vlm_weights_common_filtered_canonical(model_path, |name| {
+        name.starts_with("model.embed_tokens.") || name.starts_with(image_prefix)
+    })?;
+    let (mut weights, lora_pairs) = remap_phi4mm_weights(raw_weights)?;
+    if !lora_pairs.is_empty() {
+        return Err(anyhow::anyhow!(
+            "filtered Phi4MM XLA media loader unexpectedly retained decoder LoRA tensors"
+        ));
+    }
+    models::sanitize_tied_embeddings(&mut weights, &full_config);
+
+    let text_embeddings = mlxcel_core::layers::UnifiedEmbedding::from_weights(
+        &weights,
+        "model.embed_tokens",
+        text_config.group_size(),
+        text_config.bits(),
+    )
+    .map_err(|error| anyhow::anyhow!("invalid Phi4MM text embedding table: {error}"))?;
+    let vision_tower = Phi4SigLipVisionEncoder::from_weights(
+        &weights,
+        &vision_config,
+        "vision_tower.vision_tower.vision_model",
+        text_config.group_size(),
+        text_config.bits(),
+    )
+    .map_err(|error| anyhow::anyhow!("Failed to load Phi4MM vision tower: {error}"))?;
+    let mm_projector_linear1 = mlxcel_core::layers::UnifiedLinear::from_weights(
+        &weights,
+        "mm_projector_linear1",
+        text_config.group_size(),
+        text_config.bits(),
+    )
+    .map_err(|error| anyhow::anyhow!("Failed to load Phi4MM mm_projector_linear1: {error}"))?;
+    let mm_projector_linear2 = mlxcel_core::layers::UnifiedLinear::from_weights(
+        &weights,
+        "mm_projector_linear2",
+        text_config.group_size(),
+        text_config.bits(),
+    )
+    .map_err(|error| anyhow::anyhow!("Failed to load Phi4MM mm_projector_linear2: {error}"))?;
+    let glb_gn = weights
+        .get("glb_GN")
+        .map(|weight| mlxcel_core::copy(weight))
+        .ok_or_else(|| anyhow::anyhow!("Phi4MM glb_GN weight not found"))?;
+    let sub_gn = weights
+        .get("sub_GN")
+        .map(|weight| mlxcel_core::copy(weight))
+        .ok_or_else(|| anyhow::anyhow!("Phi4MM sub_GN weight not found"))?;
+
+    let image_embd_layer = full_config
+        .get("embd_layer")
+        .and_then(|layer| layer.get("image_embd_layer"));
+    let crop_size = image_embd_layer
+        .and_then(|config| config.get("crop_size"))
+        .and_then(Value::as_u64)
+        .unwrap_or(448) as usize;
+    let hd_transform_order = image_embd_layer
+        .and_then(|config| config.get("hd_transform_order"))
+        .and_then(Value::as_str)
+        .unwrap_or("glb_sub")
+        .to_string();
+    let dynamic_hd = read_optional_model_json(model_path, "preprocessor_config.json")
+        .as_ref()
+        .and_then(|config| config.get("dynamic_hd"))
+        .and_then(Value::as_u64)
+        .unwrap_or(36) as usize;
+
+    Ok((
+        text_embeddings,
+        text_config.hidden_size,
+        text_config.max_position_embeddings,
+        Phi4MMXlaVisionComponents {
+            vision_tower,
+            mm_projector_linear1,
+            mm_projector_linear2,
+            processor: Phi4MMProcessor::new(crop_size, vision_config.patch_size, dynamic_hd),
+            select_layer: -2,
+            glb_gn,
+            sub_gn,
+            hd_transform_order,
+        },
+    ))
+}
+
 pub(crate) fn load_phi4_siglip_vlm(model_path: &Path) -> Result<LoadedModel> {
     use vision::encoders::phi4_siglip::{Phi4SigLipVisionConfig, Phi4SigLipVisionEncoder};
     use vision::processors::phi4_siglip::Phi4SigLipProcessor;
diff --git a/src/multimodal/mod.rs b/src/multimodal/mod.rs
index cdb483942..fdf3e0e7a 100644
--- a/src/multimodal/mod.rs
+++ b/src/multimodal/mod.rs
@@ -51,6 +51,9 @@ pub mod moondream3_prompt;
 pub mod phi3v_prompt;
 pub mod phi4_siglip_prompt;
 pub mod phi4mm_prompt;
+#[cfg(feature = "xla-iree")]
+#[doc(hidden)]
+pub mod phi4mm_xla_audio;
 pub mod pixtral_prompt;
 pub mod qwen_vl;
 pub mod smolvlm_prompt;
diff --git a/src/multimodal/phi4mm_xla_audio.rs b/src/multimodal/phi4mm_xla_audio.rs
new file mode 100644
index 000000000..e3d2fa9e9
--- /dev/null
+++ b/src/multimodal/phi4mm_xla_audio.rs
@@ -0,0 +1,704 @@
+// Copyright 2025-2026 Lablup Inc. and Jeongkyu Shin
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+//! Phi4MM host media/text-embedding producer paired with the IREE audio and
+//! language runtimes.
+//!
+//! SpeechLib, the qualified image tower/projector, and the checkpoint embedding
+//! table execute through MLX. The cascaded Conformer, audio projection,
+//! modality-LoRA decoder, KV, and logits stay on the IREE path; no MLX decoder
+//! or MLX audio encoder is constructed or retained.
+
+use std::path::{Path, PathBuf};
+
+use mlxcel_core::layers::UnifiedEmbedding;
+use mlxcel_core::session::{
+    OwnedTensor, PreparedAdapterMode, PreparedAttentionBias, PreparedModality, PreparedPositions,
+    PreparedPrefill, PreparedTensorDType,
+};
+use mlxcel_xla::{Phi4AudioProjectionMode, Phi4AudioRuntime, phi4_audio_bucket_for_frames};
+
+use crate::audio::phi4mm::Phi4MMAudioFeatureExtractor;
+use crate::audio::{AudioPreprocessCheckpoint, AudioWaveformBatch};
+use crate::loading::Phi4MMXlaVisionComponents;
+use crate::multimodal::phi4_siglip_prompt::PHI4_SIGLIP_IMAGE_TOKEN_INDEX;
+use crate::multimodal::phi4mm_prompt::{PHI4MM_AUDIO_TOKEN_ID, expand_phi4mm_placeholders};
+use crate::vision::processors::phi4mm::Phi4MMImageInput;
+
+#[doc(hidden)]
+pub fn load_phi4mm_audio_policy(
+    model_path: &Path,
+) -> Result<crate::audio::AudioFamilyPolicy, String> {
+    let config_text = std::fs::read_to_string(model_path.join("config.json"))
+        .map_err(|error| format!("read Phi4MM config.json: {error}"))?;
+    let config: serde_json::Value = serde_json::from_str(&config_text)
+        .map_err(|error| format!("parse Phi4MM config.json: {error}"))?;
+    let preprocessor_path = model_path.join("preprocessor_config.json");
+    let preprocessor = if preprocessor_path.is_file() {
+        let text = std::fs::read_to_string(&preprocessor_path)
+            .map_err(|error| format!("read {}: {error}", preprocessor_path.display()))?;
+        Some(
+            serde_json::from_str::<serde_json::Value>(&text)
+                .map_err(|error| format!("parse {}: {error}", preprocessor_path.display()))?,
+        )
+    } else {
+        None
+    };
+    crate::audio::AudioFamilyPolicy::from_phi4mm_configs(&config, preprocessor.as_ref())
+        .map_err(|error| error.to_string())
+}
+
+/// Thread-confined host producer for Phi4MM audio requests.
+#[doc(hidden)]
+pub struct Phi4MmXlaAudioProducer {
+    model_path: PathBuf,
+    device: String,
+    extractor: Phi4MMAudioFeatureExtractor,
+    text_embeddings: UnifiedEmbedding,
+    hidden_size: usize,
+    max_sequence_len: usize,
+    runtime: Option<Phi4AudioRuntime>,
+    vision: Option<Phi4MMXlaVisionComponents>,
+}
+
+impl Phi4MmXlaAudioProducer {
+    #[doc(hidden)]
+    pub fn load(model_path: &Path, device: &str, context_capacity: usize) -> Result<Self, String> {
+        let (text_embeddings, hidden_size, checkpoint_capacity) =
+            crate::loading::load_phi4mm_xla_text_embeddings(model_path)
+                .map_err(|error| error.to_string())?;
+        let max_sequence_len = context_capacity.min(checkpoint_capacity);
+        if max_sequence_len == 0 {
+            return Err("Phi4MM XLA audio context capacity must be positive".to_string());
+        }
+        Ok(Self {
+            model_path: model_path.to_path_buf(),
+            device: device.to_string(),
+            extractor: Phi4MMAudioFeatureExtractor::new(),
+            text_embeddings,
+            hidden_size,
+            max_sequence_len,
+            runtime: None,
+            vision: None,
+        })
+    }
+
+    /// Load the same audio path plus the filtered #874 image tower/projector.
+    ///
+    /// This constructor is used by serving because one worker must be able to
+    /// prepare image-only, audio-only, and mixed requests without changing
+    /// process-global model state.
+    #[doc(hidden)]
+    pub fn load_multimodal(
+        model_path: &Path,
+        device: &str,
+        context_capacity: usize,
+    ) -> Result<Self, String> {
+        let (text_embeddings, hidden_size, checkpoint_capacity, vision) =
+            crate::loading::load_phi4mm_xla_media_components(model_path)
+                .map_err(|error| error.to_string())?;
+        let max_sequence_len = context_capacity.min(checkpoint_capacity);
+        if max_sequence_len == 0 {
+            return Err("Phi4MM XLA media context capacity must be positive".to_string());
+        }
+        Ok(Self {
+            model_path: model_path.to_path_buf(),
+            device: device.to_string(),
+            extractor: Phi4MMAudioFeatureExtractor::new(),
+            hidden_size,
+            max_sequence_len,
+            runtime: None,
+            text_embeddings,
+            vision: Some(vision),
+        })
+    }
+
+    #[doc(hidden)]
+    pub fn prepare_audio(
+        &mut self,
+        waveforms: AudioWaveformBatch,
+        token_ids: Vec<i32>,
+        cancelled: &std::sync::atomic::AtomicBool,
+    ) -> Result<PreparedPrefill, String> {
+        self.prepare_media(waveforms, token_ids, &[], cancelled)
+    }
+
+    #[doc(hidden)]
+    pub fn prepare_images(
+        &mut self,
+        token_ids: Vec<i32>,
+        images: &[image::DynamicImage],
+        cancelled: &std::sync::atomic::AtomicBool,
+    ) -> Result<PreparedPrefill, String> {
+        self.prepare_media(
+            AudioWaveformBatch {
+                family: "phi4mm",
+                clips: Vec::new(),
+                boundaries: Vec::new(),
+                total_source_samples: 0,
+                total_samples: 0,
+                total_source_duration_micros: 0,
+                estimated_frames: 0,
+                effective_audio_tokens: 0,
+            },
+            token_ids,
+            images,
+            cancelled,
+        )
+    }
+
+    /// Prepare image-only, audio-only, or mixed Phi4MM embeddings.
+    ///
+    /// Any request containing an image selects the Vision decoder adapter and
+    /// the vision branch of the audio projector. Audio-only requests select the
+    /// Speech branch. This exactly matches the request mode contract qualified
+    /// by #874.
+    #[doc(hidden)]
+    pub fn prepare_media(
+        &mut self,
+        waveforms: AudioWaveformBatch,
+        token_ids: Vec<i32>,
+        images: &[image::DynamicImage],
+        cancelled: &std::sync::atomic::AtomicBool,
+    ) -> Result<PreparedPrefill, String> {
+        if waveforms.family != "phi4mm" {
+            return Err(format!(
+                "Phi4MM XLA producer received `{}` audio policy",
+                waveforms.family
+            ));
+        }
+        if !images.is_empty() && self.vision.is_none() {
+            return Err(
+                "Phi4MM image input requires the filtered XLA multimodal producer".to_string(),
+            );
+        }
+        if cancelled.load(std::sync::atomic::Ordering::Acquire) {
+            return Err("Phi4MM media preparation was cancelled before features".to_string());
+        }
+
+        let processed_images = self
+            .vision
+            .as_ref()
+            .map(|vision| vision.processor.preprocess(images))
+            .unwrap_or_default();
+        let image_sizes = processed_images
+            .iter()
+            .map(|image| image.num_img_tokens)
+            .collect::<Vec<_>>();
+
+        let clips = waveforms
+            .clips
+            .into_iter()
+            .map(|clip| (clip.samples, clip.sample_rate))
+            .collect::<Vec<_>>();
+        let (audio_features, audio_sizes, audio_frame_lengths) = if clips.is_empty() {
+            (Vec::new(), Vec::new(), Vec::new())
+        } else {
+            let batch = self
+                .extractor
+                .extract_batch_cancellable(&clips, cancelled)?;
+            (batch.clips, batch.embed_sizes, batch.frame_lengths)
+        };
+
+        let vision_mode = !processed_images.is_empty();
+        let projection_mode = if vision_mode {
+            Phi4AudioProjectionMode::Vision
+        } else {
+            Phi4AudioProjectionMode::Speech
+        };
+        let mut projected_audio = Vec::with_capacity(audio_features.len());
+        for (index, features) in audio_features.iter().enumerate() {
+            if cancelled.load(std::sync::atomic::Ordering::Acquire) {
+                return Err(format!(
+                    "Phi4MM audio preparation was cancelled at {:?}",
+                    AudioPreprocessCheckpoint::Feature
+                ));
+            }
+            let frame_len = audio_frame_lengths[index];
+            let bucket = phi4_audio_bucket_for_frames(frame_len).ok_or_else(|| {
+                format!(
+                    "Phi4MM audio clip {} has unsupported {frame_len} frames",
+                    index + 1
+                )
+            })?;
+            if self
+                .runtime
+                .as_ref()
+                .is_none_or(|runtime| runtime.frame_bucket() != bucket)
+            {
+                self.runtime = Some(Phi4AudioRuntime::load(
+                    &self.model_path,
+                    &self.device,
+                    bucket,
+                )?);
+            }
+            let features = mlxcel_core::astype(features, mlxcel_core::dtype::FLOAT32);
+            let feature_values = raw_f32(&features, "Phi4MM SpeechLib features")?;
+            let output = self
+                .runtime
+                .as_mut()
+                .expect("runtime loaded for selected bucket")
+                .project(&feature_values, frame_len, projection_mode)?;
+            if output.valid_rows != audio_sizes[index] {
+                return Err(format!(
+                    "Phi4MM audio clip {} produced {} projection rows, expected {}",
+                    index + 1,
+                    output.valid_rows,
+                    audio_sizes[index]
+                ));
+            }
+            if output.hidden_size != self.hidden_size {
+                return Err(format!(
+                    "Phi4MM audio projection hidden size {} does not match text hidden size {}",
+                    output.hidden_size, self.hidden_size
+                ));
+            }
+            projected_audio.push(output.projected);
+        }
+
+        let logical_tokens = expand_phi4mm_placeholders(&token_ids, &image_sizes, &audio_sizes)?;
+        if logical_tokens.len() > self.max_sequence_len {
+            return Err(format!(
+                "Phi4MM prepared sequence length {} exceeds XLA context capacity {}",
+                logical_tokens.len(),
+                self.max_sequence_len
+            ));
+        }
+        let safe_tokens = logical_tokens
+            .iter()
+            .map(|&token| {
+                if token == PHI4MM_AUDIO_TOKEN_ID || token == PHI4_SIGLIP_IMAGE_TOKEN_INDEX {
+                    0
+                } else {
+                    token
+                }
+            })
+            .collect::<Vec<_>>();
+        let input_ids =
+            mlxcel_core::from_slice_i32(&safe_tokens, &[1, logical_tokens.len() as i32]);
+        let text = self.text_embeddings.forward(&input_ids);
+        let shape = mlxcel_core::array_shape(&text);
+        if shape != [1, logical_tokens.len() as i32, self.hidden_size as i32] {
+            return Err(format!(
+                "Phi4MM text embeddings have shape {shape:?}, expected [1, {}, {}]",
+                logical_tokens.len(),
+                self.hidden_size
+            ));
+        }
+        let text_dtype = mlxcel_core::array_dtype(&text);
+        let mut projected_images = Vec::with_capacity(processed_images.len());
+        if let Some(vision) = self.vision.as_ref() {
+            for (index, processed) in processed_images.iter().enumerate() {
+                if cancelled.load(std::sync::atomic::Ordering::Acquire) {
+                    return Err(format!(
+                        "Phi4MM image preparation was cancelled before image {}",
+                        index + 1
+                    ));
+                }
+                let projected = vision.hd_transform(processed, text_dtype)?;
+                let shape = mlxcel_core::array_shape(&projected);
+                if shape != [1, image_sizes[index] as i32, self.hidden_size as i32] {
+                    return Err(format!(
+                        "Phi4MM image {} projection has shape {shape:?}, expected [1, {}, {}]",
+                        index + 1,
+                        image_sizes[index],
+                        self.hidden_size
+                    ));
+                }
+                let projected = mlxcel_core::astype(&projected, mlxcel_core::dtype::FLOAT32);
+                projected_images.push(raw_f32(
+                    &projected,
+                    &format!("Phi4MM projected image {}", index + 1),
+                )?);
+            }
+        }
+        let text = mlxcel_core::astype(&text, mlxcel_core::dtype::FLOAT32);
+        let mut merged = raw_f32(&text, "Phi4MM text embeddings")?;
+        merge_media_rows(
+            &logical_tokens,
+            &projected_images,
+            &projected_audio,
+            self.hidden_size,
+            &mut merged,
+        )?;
+        let embeddings = OwnedTensor::new(
+            f32_bytes(&merged),
+            PreparedTensorDType::Float32,
+            vec![1, logical_tokens.len(), self.hidden_size],
+        )
+        .map_err(|error| error.to_string())?;
+        let attention_bias = PreparedAttentionBias {
+            tensor: OwnedTensor::new(
+                vec![0; logical_tokens.len() * std::mem::size_of::<f32>()],
+                PreparedTensorDType::Float32,
+                vec![1, 1, 1, logical_tokens.len()],
+            )
+            .map_err(|error| error.to_string())?,
+            causal: true,
+        };
+        let mut modalities = Vec::with_capacity(2);
+        if !projected_images.is_empty() {
+            modalities.push(PreparedModality {
+                family: "phi4mm-image".to_string(),
+                item_count: projected_images.len(),
+                token_count: image_sizes.iter().sum(),
+            });
+        }
+        if !projected_audio.is_empty() {
+            modalities.push(PreparedModality {
+                family: "phi4mm-audio".to_string(),
+                item_count: projected_audio.len(),
+                token_count: audio_sizes.iter().sum(),
+            });
+        }
+        let adapter_mode = if vision_mode {
+            PreparedAdapterMode::Vision
+        } else if !projected_audio.is_empty() {
+            PreparedAdapterMode::Speech
+        } else {
+            PreparedAdapterMode::Language
+        };
+        PreparedPrefill::new(
+            logical_tokens,
+            embeddings,
+            PreparedPositions::Sequential {
+                start: 0,
+                length: safe_tokens.len(),
+            },
+            attention_bias,
+            modalities,
+        )
+        .map(|prepared| prepared.with_adapter_mode(adapter_mode))
+        .map_err(|error| error.to_string())
+    }
+}
+
+impl Phi4MMXlaVisionComponents {
+    /// Execute the exact #874 HD image transform while retaining only the
+    /// filtered image tower/projector weights.
+    fn hd_transform(
+        &self,
+        processed: &Phi4MMImageInput,
+        target_dtype: i32,
+    ) -> Result<mlxcel_core::UniquePtr<mlxcel_core::MlxArray>, String> {
+        let (h_crops, w_crops) = processed.image_grid;
+        let expected_crops = h_crops
+            .checked_mul(w_crops)
+            .and_then(|count| count.checked_add(1))
+            .ok_or("Phi4MM image crop count overflow")?;
+        if processed.crops.len() != expected_crops {
+            return Err(format!(
+                "Phi4MM image processor produced {} crops, expected {expected_crops}",
+                processed.crops.len()
+            ));
+        }
+        let pooled_grid = processed.pooled_grid_size;
+        if pooled_grid == 0 {
+            return Err("Phi4MM image pooled grid must be positive".to_string());
+        }
+        let spatial_per_crop = (
+            (self.processor.crop_size / self.processor.patch_size) as i32,
+            (self.processor.crop_size / self.processor.patch_size) as i32,
+        );
+        let mut crop_features = Vec::with_capacity(processed.crops.len());
+        for crop in &processed.crops {
+            let pixels = mlxcel_core::astype(&crop.pixel_values, target_dtype);
+            let mut hidden_states = self
+                .vision_tower
+                .forward_hidden_states(&pixels, spatial_per_crop);
+            let layer_count = hidden_states.len() as isize;
+            let selected = if self.select_layer < 0 {
+                layer_count.checked_add(self.select_layer)
+            } else {
+                Some(self.select_layer)
+            }
+            .filter(|&index| index >= 0 && index < layer_count)
+            .ok_or_else(|| {
+                format!(
+                    "Phi4MM vision select layer {} is outside {layer_count} hidden states",
+                    self.select_layer
+                )
+            })? as usize;
+            crop_features.push(hidden_states.swap_remove(selected));
+        }
+
+        let original_grid = self.processor.crop_size / self.processor.patch_size;
+        let pooled_features = crop_features
+            .iter()
+            .map(|features| avg_pool_2d(features, original_grid, pooled_grid))
+            .collect::<Vec<_>>();
+        let vision_dim = mlxcel_core::array_shape(&pooled_features[0])[2];
+        let global_grid = mlxcel_core::reshape(
+            &pooled_features[0],
+            &[1, pooled_grid as i32, pooled_grid as i32, vision_dim],
+        );
+        let sub_gn_col = self.make_sub_gn_column(pooled_grid as i32, vision_dim, target_dtype);
+        let global_with_sep = mlxcel_core::concatenate(&global_grid, &sub_gn_col, 2);
+        let global_tokens = mlxcel_core::reshape(
+            &global_with_sep,
+            &[1, (pooled_grid * (pooled_grid + 1)) as i32, vision_dim],
+        );
+        let sub_tokens =
+            self.assemble_sub_features(&pooled_features[1..], processed, vision_dim, target_dtype)?;
+        let global_separator = mlxcel_core::astype(&self.glb_gn, target_dtype);
+        let global_separator = mlxcel_core::reshape(&global_separator, &[1, 1, vision_dim]);
+        let combined = if self.hd_transform_order == "sub_glb" {
+            concat_3arrays(&sub_tokens, &global_separator, &global_tokens)
+        } else {
+            concat_3arrays(&global_tokens, &global_separator, &sub_tokens)
+        };
+        let projected = self.mm_projector_linear1.forward(&combined);
+        let projected = mlxcel_core::gelu_approx(&projected);
+        Ok(self.mm_projector_linear2.forward(&projected))
+    }
+
+    fn make_sub_gn_column(
+        &self,
+        rows: i32,
+        dim: i32,
+        dtype: i32,
+    ) -> mlxcel_core::UniquePtr<mlxcel_core::MlxArray> {
+        let separator = mlxcel_core::astype(&self.sub_gn, dtype);
+        mlxcel_core::broadcast_to(&separator, &[1, rows, 1, dim])
+    }
+
+    fn assemble_sub_features(
+        &self,
+        features: &[mlxcel_core::UniquePtr<mlxcel_core::MlxArray>],
+        processed: &Phi4MMImageInput,
+        vision_dim: i32,
+        dtype: i32,
+    ) -> Result<mlxcel_core::UniquePtr<mlxcel_core::MlxArray>, String> {
+        let (h_crops, w_crops) = processed.image_grid;
+        let pooled_grid = processed.pooled_grid_size;
+        if features.len() != h_crops.saturating_mul(w_crops) {
+            return Err(format!(
+                "Phi4MM sub-crop feature count is {}, expected {}",
+                features.len(),
+                h_crops.saturating_mul(w_crops)
+            ));
+        }
+        let grid = pooled_grid as i32;
+        let total_h = h_crops * pooled_grid;
+        let total_w = w_crops * pooled_grid;
+        let mut row_segments = Vec::with_capacity(h_crops);
+        for crop_row in 0..h_crops {
+            let mut columns = Vec::with_capacity(w_crops);
+            for crop_col in 0..w_crops {
+                let index = crop_row * w_crops + crop_col;
+                columns.push(mlxcel_core::reshape(
+                    &features[index],
+                    &[grid, grid, vision_dim],
+                ));
+            }
+            row_segments.push(concat_arrays_axis(&columns, 1));
+        }
+        let spatial = concat_arrays_axis(&row_segments, 0);
+        let useful_h = processed.active_rows.min(total_h) as i32;
+        let useful_w = processed.active_cols.min(total_w) as i32;
+        if useful_h <= 0 || useful_w <= 0 {
+            return Err("Phi4MM image active region must be positive".to_string());
+        }
+        let cropped = if useful_h < total_h as i32 || useful_w < total_w as i32 {
+            mlxcel_core::slice(&spatial, &[0, 0, 0], &[useful_h, useful_w, vision_dim])
+        } else {
+            spatial
+        };
+        let cropped = mlxcel_core::reshape(&cropped, &[1, useful_h, useful_w, vision_dim]);
+        let separator = self.make_sub_gn_column(useful_h, vision_dim, dtype);
+        let with_separator = mlxcel_core::concatenate(&cropped, &separator, 2);
+        Ok(mlxcel_core::reshape(
+            &with_separator,
+            &[1, useful_h * (useful_w + 1), vision_dim],
+        ))
+    }
+}
+
+fn avg_pool_2d(
+    features: &mlxcel_core::MlxArray,
+    original_grid: usize,
+    target_grid: usize,
+) -> mlxcel_core::UniquePtr<mlxcel_core::MlxArray> {
+    let shape = mlxcel_core::array_shape(features);
+    let dim = shape[2];
+    let original_grid = original_grid as i32;
+    let target_grid = target_grid as i32;
+    let pool_size = original_grid / target_grid;
+    let spatial = mlxcel_core::reshape(features, &[original_grid, original_grid, dim]);
+    let transposed = mlxcel_core::transpose_axes(&spatial, &[2, 0, 1]);
+    let batched = mlxcel_core::reshape(&transposed, &[dim, 1, original_grid, original_grid]);
+    let blocked = mlxcel_core::reshape(
+        &batched,
+        &[dim, 1, target_grid, pool_size, target_grid, pool_size],
+    );
+    let mean = mlxcel_core::mean_axis(&blocked, 5, true);
+    let mean = mlxcel_core::mean_axis(&mean, 3, true);
+    let squeezed = mlxcel_core::reshape(&mean, &[dim, target_grid, target_grid]);
+    let result = mlxcel_core::transpose_axes(&squeezed, &[1, 2, 0]);
+    mlxcel_core::reshape(&result, &[1, target_grid * target_grid, dim])
+}
+
+fn concat_3arrays(
+    first: &mlxcel_core::MlxArray,
+    second: &mlxcel_core::MlxArray,
+    third: &mlxcel_core::MlxArray,
+) -> mlxcel_core::UniquePtr<mlxcel_core::MlxArray> {
+    let first_two = mlxcel_core::concatenate(first, second, 1);
+    mlxcel_core::concatenate(&first_two, third, 1)
+}
+
+fn concat_arrays_axis(
+    arrays: &[mlxcel_core::UniquePtr<mlxcel_core::MlxArray>],
+    axis: i32,
+) -> mlxcel_core::UniquePtr<mlxcel_core::MlxArray> {
+    let mut output = mlxcel_core::copy(&arrays[0]);
+    for array in &arrays[1..] {
+        output = mlxcel_core::concatenate(&output, array, axis);
+    }
+    output
+}
+
+fn raw_f32(array: &mlxcel_core::MlxArray, label: &str) -> Result<Vec<f32>, String> {
+    let bytes = mlxcel_core::try_array_to_raw_bytes(array)
+        .map_err(|error| format!("export {label}: {error}"))?;
+    if bytes.len() % std::mem::size_of::<f32>() != 0 {
+        return Err(format!("{label} has a non-f32 byte length {}", bytes.len()));
+    }
+    Ok(bytes
+        .chunks_exact(4)
+        .map(|chunk| f32::from_ne_bytes(chunk.try_into().expect("four-byte f32")))
+        .collect())
+}
+
+fn f32_bytes(values: &[f32]) -> Vec<u8> {
+    values
+        .iter()
+        .flat_map(|value| value.to_ne_bytes())
+        .collect()
+}
+
+fn merge_media_rows(
+    logical_tokens: &[i32],
+    projected_images: &[Vec<f32>],
+    projected_audio: &[Vec<f32>],
+    hidden_size: usize,
+    merged: &mut [f32],
+) -> Result<(), String> {
+    let expected = logical_tokens
+        .len()
+        .checked_mul(hidden_size)
+        .ok_or("Phi4MM merged embedding size overflow")?;
+    if merged.len() != expected {
+        return Err(format!(
+            "Phi4MM merged embedding buffer has {} elements, expected {expected}",
+            merged.len()
+        ));
+    }
+    let mut token_index = 0usize;
+    let mut image_index = 0usize;
+    let mut audio_index = 0usize;
+    while token_index < logical_tokens.len() {
+        let kind = logical_tokens[token_index];
+        if kind != PHI4_SIGLIP_IMAGE_TOKEN_INDEX && kind != PHI4MM_AUDIO_TOKEN_ID {
+            token_index += 1;
+            continue;
+        }
+        let (rows, ordinal, label) = if kind == PHI4_SIGLIP_IMAGE_TOKEN_INDEX {
+            let rows = projected_images
+                .get(image_index)
+                .ok_or("Phi4MM prompt has more image placeholder rows than image inputs")?;
+            image_index += 1;
+            (rows, image_index, "image")
+        } else {
+            let rows = projected_audio
+                .get(audio_index)
+                .ok_or("Phi4MM prompt has more audio placeholder rows than audio inputs")?;
+            audio_index += 1;
+            (rows, audio_index, "audio")
+        };
+        if rows.len() % hidden_size != 0 {
+            return Err(format!(
+                "Phi4MM projected {label} {ordinal} has {} elements not divisible by hidden size {hidden_size}",
+                rows.len()
+            ));
+        }
+        let row_count = rows.len() / hidden_size;
+        let end = token_index
+            .checked_add(row_count)
+            .ok_or("Phi4MM media placeholder range overflow")?;
+        if row_count == 0
+            || end > logical_tokens.len()
+            || logical_tokens[token_index..end]
+                .iter()
+                .any(|&token| token != kind)
+        {
+            return Err(format!(
+                "Phi4MM {label} placeholder at position {token_index} does not contain {row_count} contiguous rows"
+            ));
+        }
+        merged[token_index * hidden_size..end * hidden_size].copy_from_slice(rows);
+        token_index = end;
+    }
+    if image_index != projected_images.len() || audio_index != projected_audio.len() {
+        return Err(format!(
+            "Phi4MM prompt consumed {image_index} image/{audio_index} audio inputs but {}/{} were prepared",
+            projected_images.len(),
+            projected_audio.len()
+        ));
+    }
+    Ok(())
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn audio_rows_replace_exact_adjacent_placeholder_segments() {
+        let tokens = vec![
+            7,
+            PHI4MM_AUDIO_TOKEN_ID,
+            PHI4MM_AUDIO_TOKEN_ID,
+            8,
+            PHI4MM_AUDIO_TOKEN_ID,
+            9,
+        ];
+        let projected_audio = vec![vec![1.0, 2.0, 3.0, 4.0], vec![5.0, 6.0]];
+        let mut merged = vec![0.0; tokens.len() * 2];
+        merge_media_rows(&tokens, &[], &projected_audio, 2, &mut merged).unwrap();
+        assert_eq!(
+            merged,
+            vec![0.0, 0.0, 1.0, 2.0, 3.0, 4.0, 0.0, 0.0, 5.0, 6.0, 0.0, 0.0]
+        );
+    }
+
+    #[test]
+    fn image_and_audio_rows_preserve_mixed_prompt_order() {
+        let tokens = vec![
+            PHI4_SIGLIP_IMAGE_TOKEN_INDEX,
+            PHI4_SIGLIP_IMAGE_TOKEN_INDEX,
+            7,
+            PHI4MM_AUDIO_TOKEN_ID,
+            8,
+        ];
+        let images = vec![vec![1.0, 2.0, 3.0, 4.0]];
+        let audio = vec![vec![5.0, 6.0]];
+        let mut merged = vec![0.0; tokens.len() * 2];
+        merge_media_rows(&tokens, &images, &audio, 2, &mut merged).unwrap();
+        assert_eq!(
+            merged,
+            vec![1.0, 2.0, 3.0, 4.0, 0.0, 0.0, 5.0, 6.0, 0.0, 0.0]
+        );
+    }
+}
diff --git a/src/server/batch/xla_audio_preprocess.rs b/src/server/batch/xla_audio_preprocess.rs
index 8078d77eb..12e2401eb 100644
--- a/src/server/batch/xla_audio_preprocess.rs
+++ b/src/server/batch/xla_audio_preprocess.rs
@@ -12,11 +12,12 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-//! Bounded host audio preprocessing outside the XLA scheduling loop.
+//! Bounded Phi4MM media preprocessing outside the XLA scheduling loop.
 //!
-//! This stage owns encoded request buffers, waveform decode/resampling, and a
-//! future family feature producer. It intentionally has no production
-//! Phi4MM/Gemma3n producer yet; therefore XLA admission remains audio-false.
+//! This stage owns encoded image/audio request buffers, image and waveform
+//! decoding, resampling, host feature extraction, and prepared-prefill export.
+//! A single thread-confined producer handles image-only, audio-only, and mixed
+//! Phi4MM requests so request mode selection cannot split across workers.
 
 use std::mem::size_of;
 use std::panic::{AssertUnwindSafe, catch_unwind};
@@ -38,6 +39,8 @@ pub(crate) struct AudioPreprocessJob {
     pub job_id: u64,
     pub token_ids: Vec<i32>,
     pub max_prefill_tokens: usize,
+    pub expected_image_count: usize,
+    pub images: Vec<Vec<u8>>,
     pub clips: Vec<AudioEncodedClip>,
     pub policy: AudioFamilyPolicy,
     pub cancelled: Arc<AtomicBool>,
@@ -66,6 +69,11 @@ pub(crate) enum AudioStageError {
     },
     #[error("audio feature preprocessing panicked for {family}")]
     FeaturePanic { family: &'static str },
+    #[error("media image preprocessing failed for {family}: {reason}")]
+    Image {
+        family: &'static str,
+        reason: String,
+    },
     #[error(
         "audio prepared-prefill context limit exceeded for {family}: actual {actual}, maximum {maximum}"
     )]
@@ -114,15 +122,29 @@ pub(crate) enum AudioQueueError {
     Overflow,
 }
 
-pub(crate) trait AudioFeatureProducer: Send + 'static {
+pub(crate) trait AudioFeatureProducer: 'static {
     fn prepare(
         &mut self,
         waveforms: AudioWaveformBatch,
         token_ids: Vec<i32>,
+        images: Vec<image::DynamicImage>,
         cancelled: &AtomicBool,
     ) -> Result<PreparedPrefill, String>;
 }
 
+#[cfg(feature = "xla-iree")]
+impl AudioFeatureProducer for crate::multimodal::phi4mm_xla_audio::Phi4MmXlaAudioProducer {
+    fn prepare(
+        &mut self,
+        waveforms: AudioWaveformBatch,
+        token_ids: Vec<i32>,
+        images: Vec<image::DynamicImage>,
+        cancelled: &AtomicBool,
+    ) -> Result<PreparedPrefill, String> {
+        self.prepare_media(waveforms, token_ids, &images, cancelled)
+    }
+}
+
 pub(crate) struct AudioPreprocessLimits {
     pub queue_depth: usize,
     pub result_queue_depth: usize,
@@ -394,13 +416,29 @@ pub(crate) struct AudioPreprocessStage {
 }
 
 impl AudioPreprocessStage {
-    pub fn spawn<P: AudioFeatureProducer>(
-        mut producer: P,
+    pub fn spawn<P: AudioFeatureProducer + Send>(
+        producer: P,
         limits: AudioPreprocessLimits,
         observability: Arc<BatchObservability>,
     ) -> Result<Self, String> {
+        Self::spawn_with_loader(limits, observability, move || Ok(producer))
+    }
+
+    /// Construct MLX-backed feature producers inside their owning worker
+    /// thread. This keeps MLX handles thread-confined without imposing an
+    /// artificial `Send` contract on the producer.
+    pub fn spawn_with_loader<P, F>(
+        limits: AudioPreprocessLimits,
+        observability: Arc<BatchObservability>,
+        loader: F,
+    ) -> Result<Self, String>
+    where
+        P: AudioFeatureProducer,
+        F: FnOnce() -> Result<P, String> + Send + 'static,
+    {
         let (sender, receiver) = mpsc::sync_channel::<QueuedJob>(limits.queue_depth.max(1));
         let (result_tx, result_rx) = mpsc::sync_channel(limits.result_queue_depth.max(1));
+        let (ready_tx, ready_rx) = mpsc::sync_channel(1);
         let metrics = Arc::new(AudioPreprocessMetrics::new(observability));
         let worker_metrics = metrics.clone();
         let healthy = Arc::new(AtomicBool::new(true));
@@ -408,6 +446,24 @@ impl AudioPreprocessStage {
         let worker = thread::Builder::new()
             .name("mlxcel-xla-audio-preprocess".to_string())
             .spawn(move || {
+                let mut producer = match catch_unwind(AssertUnwindSafe(loader)) {
+                    Ok(Ok(producer)) => {
+                        let _ = ready_tx.send(Ok(()));
+                        producer
+                    }
+                    Ok(Err(error)) => {
+                        let _ = ready_tx.send(Err(error));
+                        worker_healthy.store(false, Ordering::Release);
+                        return;
+                    }
+                    Err(_) => {
+                        let _ = ready_tx.send(Err(
+                            "audio feature producer panicked during startup".to_string(),
+                        ));
+                        worker_healthy.store(false, Ordering::Release);
+                        return;
+                    }
+                };
                 while let Ok(queued) = receiver.recv() {
                     let Some((job, reservation)) = queued.dequeue() else {
                         continue;
@@ -420,6 +476,19 @@ impl AudioPreprocessStage {
                 worker_healthy.store(false, Ordering::Release);
             })
             .map_err(|error| format!("failed to spawn audio preprocessing worker: {error}"))?;
+        match ready_rx.recv() {
+            Ok(Ok(())) => {}
+            Ok(Err(error)) => {
+                let _ = worker.join();
+                return Err(error);
+            }
+            Err(_) => {
+                let _ = worker.join();
+                return Err(
+                    "audio preprocessing worker exited before reporting startup status".to_string(),
+                );
+            }
+        }
         Ok(Self {
             sender: Some(sender),
             result_rx: Some(result_rx),
@@ -442,11 +511,11 @@ impl AudioPreprocessStage {
                 checkpoint: AudioPreprocessCheckpoint::Queue,
             });
         }
-        let request_bytes = job.clips.iter().try_fold(0usize, |sum, clip| {
+        let audio_bytes = job.clips.iter().try_fold(0usize, |sum, clip| {
             sum.checked_add(clip.bytes.len())
                 .ok_or(AudioQueueError::Overflow)
         });
-        let request_bytes = match request_bytes {
+        let audio_bytes = match audio_bytes {
             Ok(bytes) => bytes,
             Err(error) => {
                 self.metrics
@@ -454,15 +523,27 @@ impl AudioPreprocessStage {
                 return Err(error);
             }
         };
-        if request_bytes > job.policy.max_encoded_bytes_per_request {
+        if audio_bytes > job.policy.max_encoded_bytes_per_request {
             self.metrics
                 .record_rejection(AudioRejectionReason::MemoryLimit);
             return Err(AudioQueueError::MemoryLimit {
-                request_bytes,
+                request_bytes: audio_bytes,
                 queued_bytes: self.metrics.queued_encoded_bytes.load(Ordering::Acquire),
                 maximum_bytes: job.policy.max_encoded_bytes_per_request,
             });
         }
+        let request_bytes = job.images.iter().try_fold(audio_bytes, |sum, image| {
+            sum.checked_add(image.len())
+                .ok_or(AudioQueueError::Overflow)
+        });
+        let request_bytes = match request_bytes {
+            Ok(bytes) => bytes,
+            Err(error) => {
+                self.metrics
+                    .record_rejection(AudioRejectionReason::Overflow);
+                return Err(error);
+            }
+        };
         let host_reservation_bytes = job_envelope_host_bytes(&job)
             .and_then(|bytes| bytes.checked_add(job.policy.max_waveform_working_bytes_per_request))
             .and_then(|bytes| bytes.checked_add(job.policy.max_prepared_result_bytes_per_request));
@@ -587,7 +668,21 @@ fn process_job<P: AudioFeatureProducer>(
     let family = job.policy.family;
     let max_prefill_tokens = job.max_prefill_tokens;
     let max_result_bytes = job.policy.max_prepared_result_bytes_per_request;
-    let outcome = match preprocess_wav_batch(&job.clips, job.policy, job.cancelled.as_ref()) {
+    let waveforms = if job.clips.is_empty() {
+        Ok(AudioWaveformBatch {
+            family: job.policy.family,
+            clips: Vec::new(),
+            boundaries: Vec::new(),
+            total_source_samples: 0,
+            total_samples: 0,
+            total_source_duration_micros: 0,
+            estimated_frames: 0,
+            effective_audio_tokens: 0,
+        })
+    } else {
+        preprocess_wav_batch(&job.clips, job.policy, job.cancelled.as_ref())
+    };
+    let outcome = match waveforms {
         Err(AudioPreprocessError::Cancelled { checkpoint, .. }) => {
             AudioPreprocessOutcome::Cancelled(checkpoint)
         }
@@ -597,40 +692,77 @@ fn process_job<P: AudioFeatureProducer>(
             if job.cancelled.load(Ordering::Acquire) {
                 AudioPreprocessOutcome::Cancelled(AudioPreprocessCheckpoint::Feature)
             } else {
-                let prepared = catch_unwind(AssertUnwindSafe(|| {
-                    producer.prepare(waveforms, job.token_ids, job.cancelled.as_ref())
-                }));
-                match prepared {
-                    Ok(Ok(_prepared)) if job.cancelled.load(Ordering::Acquire) => {
-                        AudioPreprocessOutcome::Cancelled(AudioPreprocessCheckpoint::Feature)
-                    }
-                    Ok(Ok(prepared)) if prepared.sequence_len > max_prefill_tokens => {
-                        AudioPreprocessOutcome::Failed(AudioStageError::ContextLimit {
+                let decoded_images = if job.images.is_empty() && job.expected_image_count == 0 {
+                    Ok(Vec::new())
+                } else {
+                    crate::server::model_provider::model_worker::decode_request_images(&job.images)
+                };
+                match decoded_images {
+                    Ok(images) if images.len() != job.expected_image_count => {
+                        AudioPreprocessOutcome::Failed(AudioStageError::Image {
                             family,
-                            actual: prepared.sequence_len,
-                            maximum: max_prefill_tokens,
+                            reason: format!(
+                                "decoded image cardinality mismatch: expected {}, decoded {}",
+                                job.expected_image_count,
+                                images.len()
+                            ),
                         })
                     }
-                    Ok(Ok(prepared)) => {
-                        let result_bytes =
-                            prepared_prefill_host_bytes(&prepared).unwrap_or(usize::MAX);
-                        if result_bytes > max_result_bytes {
-                            AudioPreprocessOutcome::Failed(AudioStageError::ResultMemoryLimit {
-                                family,
-                                actual: result_bytes,
-                                maximum: max_result_bytes,
-                            })
-                        } else {
-                            metrics.record_effective_prefill(prepared.sequence_len);
-                            AudioPreprocessOutcome::Prepared(prepared)
+                    Err(error) => AudioPreprocessOutcome::Failed(AudioStageError::Image {
+                        family,
+                        reason: error.to_string(),
+                    }),
+                    Ok(images) => {
+                        let prepared = catch_unwind(AssertUnwindSafe(|| {
+                            producer.prepare(
+                                waveforms,
+                                job.token_ids,
+                                images,
+                                job.cancelled.as_ref(),
+                            )
+                        }));
+                        match prepared {
+                            Ok(Ok(_prepared)) if job.cancelled.load(Ordering::Acquire) => {
+                                AudioPreprocessOutcome::Cancelled(
+                                    AudioPreprocessCheckpoint::Feature,
+                                )
+                            }
+                            Ok(Ok(prepared)) if prepared.sequence_len > max_prefill_tokens => {
+                                AudioPreprocessOutcome::Failed(AudioStageError::ContextLimit {
+                                    family,
+                                    actual: prepared.sequence_len,
+                                    maximum: max_prefill_tokens,
+                                })
+                            }
+                            Ok(Ok(prepared)) => {
+                                let result_bytes =
+                                    prepared_prefill_host_bytes(&prepared).unwrap_or(usize::MAX);
+                                if result_bytes > max_result_bytes {
+                                    AudioPreprocessOutcome::Failed(
+                                        AudioStageError::ResultMemoryLimit {
+                                            family,
+                                            actual: result_bytes,
+                                            maximum: max_result_bytes,
+                                        },
+                                    )
+                                } else {
+                                    metrics.record_effective_prefill(prepared.sequence_len);
+                                    AudioPreprocessOutcome::Prepared(prepared)
+                                }
+                            }
+                            Ok(Err(reason)) => {
+                                AudioPreprocessOutcome::Failed(AudioStageError::Feature {
+                                    family,
+                                    reason,
+                                })
+                            }
+                            Err(_) => {
+                                AudioPreprocessOutcome::Failed(AudioStageError::FeaturePanic {
+                                    family,
+                                })
+                            }
                         }
                     }
-                    Ok(Err(reason)) => {
-                        AudioPreprocessOutcome::Failed(AudioStageError::Feature { family, reason })
-                    }
-                    Err(_) => {
-                        AudioPreprocessOutcome::Failed(AudioStageError::FeaturePanic { family })
-                    }
                 }
             }
         }
@@ -659,6 +791,7 @@ fn process_job<P: AudioFeatureProducer>(
                 AudioStageError::Waveform(_) => AudioRejectionReason::Waveform,
                 AudioStageError::Feature { .. } => AudioRejectionReason::Feature,
                 AudioStageError::FeaturePanic { .. } => AudioRejectionReason::FeaturePanic,
+                AudioStageError::Image { .. } => AudioRejectionReason::Feature,
                 AudioStageError::ContextLimit { .. } => AudioRejectionReason::ContextLimit,
                 AudioStageError::ResultMemoryLimit { .. } => AudioRejectionReason::MemoryLimit,
             });
@@ -715,6 +848,10 @@ fn job_envelope_host_bytes(job: &AudioPreprocessJob) -> Option<usize> {
     for clip in &job.clips {
         total = total.checked_add(clip.bytes.capacity())?;
     }
+    total = total.checked_add(job.images.capacity().checked_mul(size_of::<Vec<u8>>())?)?;
+    for image in &job.images {
+        total = total.checked_add(image.capacity())?;
+    }
     Some(total)
 }
 
diff --git a/src/server/batch/xla_audio_preprocess_tests.rs b/src/server/batch/xla_audio_preprocess_tests.rs
index ac29b96f4..262147747 100644
--- a/src/server/batch/xla_audio_preprocess_tests.rs
+++ b/src/server/batch/xla_audio_preprocess_tests.rs
@@ -41,6 +41,8 @@ fn job(id: u64, samples: usize) -> AudioPreprocessJob {
         job_id: id,
         token_ids: vec![1, id as i32, 2],
         max_prefill_tokens: 4_096,
+        expected_image_count: 0,
+        images: Vec::new(),
         clips: vec![AudioEncodedClip {
             bytes: wav_pcm16(&vec![0; samples]),
             source: AudioSourceKind::ServerInline,
@@ -90,6 +92,7 @@ impl AudioFeatureProducer for RecordingProducer {
         &mut self,
         _waveforms: AudioWaveformBatch,
         token_ids: Vec<i32>,
+        _images: Vec<image::DynamicImage>,
         _cancelled: &AtomicBool,
     ) -> Result<PreparedPrefill, String> {
         self.order.lock().unwrap().push(token_ids[1] as u64);
@@ -147,6 +150,7 @@ impl AudioFeatureProducer for BlockingProducer {
         &mut self,
         _waveforms: AudioWaveformBatch,
         token_ids: Vec<i32>,
+        _images: Vec<image::DynamicImage>,
         _cancelled: &AtomicBool,
     ) -> Result<PreparedPrefill, String> {
         let _ = self.started.send(());
@@ -224,6 +228,7 @@ impl AudioFeatureProducer for CancelAfterFeature {
         &mut self,
         _waveforms: AudioWaveformBatch,
         token_ids: Vec<i32>,
+        _images: Vec<image::DynamicImage>,
         cancelled: &AtomicBool,
     ) -> Result<PreparedPrefill, String> {
         cancelled.store(true, Ordering::Release);
@@ -270,6 +275,7 @@ impl AudioFeatureProducer for PanicOnce {
         &mut self,
         _waveforms: AudioWaveformBatch,
         token_ids: Vec<i32>,
+        _images: Vec<image::DynamicImage>,
         _cancelled: &AtomicBool,
     ) -> Result<PreparedPrefill, String> {
         if !self.0 {
@@ -370,6 +376,7 @@ impl AudioFeatureProducer for StartedProducer {
         &mut self,
         _waveforms: AudioWaveformBatch,
         token_ids: Vec<i32>,
+        _images: Vec<image::DynamicImage>,
         _cancelled: &AtomicBool,
     ) -> Result<PreparedPrefill, String> {
         self.started.send(token_ids[1] as u64).unwrap();
diff --git a/src/server/batch/xla_worker.rs b/src/server/batch/xla_worker.rs
index dde268bba..81105bd6e 100644
--- a/src/server/batch/xla_worker.rs
+++ b/src/server/batch/xla_worker.rs
@@ -55,6 +55,7 @@ use mlxcel_xla::{EngineEvent, FinishReason as XlaFinishReason, SampleParams, Xla
 use super::BatchEngine;
 use super::observability::BatchObservability;
 use super::stop_matcher::StopMatcher;
+use super::xla_audio_preprocess::{AudioPreprocessLimits, AudioPreprocessStage};
 use super::xla_preprocess::ImagePreprocessStage;
 use crate::server::ServerGenerateOptions;
 use crate::server::media::MediaRequestMetadata;
@@ -167,6 +168,16 @@ struct PendingImageState {
     start: Instant,
 }
 
+struct PendingAudioState {
+    response_tx: mpsc::Sender<GenerateEvent>,
+    cancelled: Arc<AtomicBool>,
+    prompt_tokens: Vec<i32>,
+    params: SampleParams,
+    stop_sequences: Vec<String>,
+    max_tokens: usize,
+    start: Instant,
+}
+
 /// Server-side worker that serves requests through the OpenXLA continuous-batching
 /// engine. Built and run on a single worker thread (see
 /// `model_worker::spawn_xla_model_worker`).
@@ -187,6 +198,14 @@ pub(crate) struct XlaServeWorker<E = XlaBatchEngine> {
     image_preprocessor: Option<ImagePreprocessStage>,
     pending_images: HashMap<u64, PendingImageState>,
     next_image_job_id: u64,
+    audio_preprocessor: Option<AudioPreprocessStage>,
+    /// The audio stage is the unified Phi4MM image/audio producer rather than
+    /// an audio-only family producer.
+    phi4mm_media_preprocessor: bool,
+    audio_policy: Option<crate::audio::AudioFamilyPolicy>,
+    pending_audio: HashMap<u64, PendingAudioState>,
+    next_audio_job_id: u64,
+    context_capacity: usize,
     shutdown: bool,
 }
 
@@ -195,11 +214,36 @@ impl XlaServeWorker<XlaBatchEngine> {
         engine: XlaBatchEngine,
         tokenizer: MlxcelTokenizer,
         model_path: std::path::PathBuf,
+        device: String,
         request_rx: mpsc::Receiver<ModelRequest>,
         batch_metrics: Arc<BatchMetrics>,
         batch_observability: Arc<BatchObservability>,
     ) -> Result<Self, String> {
-        let image_preprocessor = ImagePreprocessStage::spawn_for_model(model_path, engine.b_max())?;
+        let context_capacity = engine.context_capacity();
+        let image_preprocessor =
+            ImagePreprocessStage::spawn_for_model(model_path.clone(), engine.b_max())?;
+        let phi4mm_media_preprocessor = crate::models::get_model_type(&model_path)
+            .map_err(|error| error.to_string())?
+            == crate::models::ModelType::Phi4MMVLM;
+        let (audio_preprocessor, audio_policy) = if phi4mm_media_preprocessor {
+            let policy =
+                crate::multimodal::phi4mm_xla_audio::load_phi4mm_audio_policy(&model_path)?;
+            let producer_path = model_path.clone();
+            let stage = AudioPreprocessStage::spawn_with_loader(
+                AudioPreprocessLimits::default(),
+                batch_observability.clone(),
+                move || {
+                    crate::multimodal::phi4mm_xla_audio::Phi4MmXlaAudioProducer::load_multimodal(
+                        &producer_path,
+                        &device,
+                        context_capacity,
+                    )
+                },
+            )?;
+            (Some(stage), Some(policy))
+        } else {
+            (None, None)
+        };
         Ok(Self {
             engine,
             tokenizer,
@@ -210,6 +254,12 @@ impl XlaServeWorker<XlaBatchEngine> {
             image_preprocessor,
             pending_images: HashMap::new(),
             next_image_job_id: 0,
+            audio_preprocessor,
+            phi4mm_media_preprocessor,
+            audio_policy,
+            pending_audio: HashMap::new(),
+            next_audio_job_id: 0,
+            context_capacity,
             shutdown: false,
         })
     }
@@ -248,6 +298,15 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
         for job_id in jobs {
             self.pending_images.remove(&job_id);
         }
+        let audio_jobs: Vec<u64> = self
+            .pending_audio
+            .iter()
+            .filter(|(_, state)| state.cancelled.load(Ordering::Relaxed))
+            .map(|(&job_id, _)| job_id)
+            .collect();
+        for job_id in audio_jobs {
+            self.pending_audio.remove(&job_id);
+        }
     }
 
     /// Map one engine step's events onto the per-request channels.
@@ -375,6 +434,11 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
                 .response_tx
                 .send(GenerateEvent::Error(msg.to_string()));
         }
+        for (_, state) in self.pending_audio.drain() {
+            let _ = state
+                .response_tx
+                .send(GenerateEvent::Error(msg.to_string()));
+        }
     }
 }
 
diff --git a/src/server/batch/xla_worker_admission.rs b/src/server/batch/xla_worker_admission.rs
index 3bf566d7b..567132af8 100644
--- a/src/server/batch/xla_worker_admission.rs
+++ b/src/server/batch/xla_worker_admission.rs
@@ -16,6 +16,9 @@
 
 use std::time::Duration;
 
+use super::super::xla_audio_preprocess::{
+    AudioPreprocessJob, AudioPreprocessOutcome, AudioPreprocessResult, AudioQueueError,
+};
 use super::super::xla_preprocess::{
     ImagePreprocessJob, ImagePreprocessOutcome, ImagePreprocessResult,
 };
@@ -51,7 +54,12 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
         response_tx: mpsc::Sender<GenerateEvent>,
         cancelled: Arc<AtomicBool>,
     ) {
-        if let Err(error) = media.validate_xla_raw_counts(images.len(), audio.len(), videos.len()) {
+        if let Err(error) = media.validate_xla_raw_counts_with_audio(
+            images.len(),
+            audio.len(),
+            videos.len(),
+            self.audio_preprocessor.is_some(),
+        ) {
             let _ = response_tx.send(GenerateEvent::Error(error));
             return;
         }
@@ -61,22 +69,74 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
             let _ = response_tx.send(GenerateEvent::Error(error));
             return;
         }
-        let prompt_tokens = match prompt_token_ids {
-            Some(tokens) => tokens,
-            None => {
-                let add_special = !prompt.starts_with("<bos>") && !prompt.starts_with("<s>");
-                match self.tokenizer.encode(&prompt, add_special) {
-                    Ok(ids) => ids.into_iter().map(|token| token as i32).collect(),
+        let uses_phi4mm_media =
+            self.phi4mm_media_preprocessor && (!images.is_empty() || !audio.is_empty());
+        let prompt_tokens = if !uses_phi4mm_media {
+            match prompt_token_ids {
+                Some(tokens) => tokens,
+                None => {
+                    let add_special = !prompt.starts_with("<bos>") && !prompt.starts_with("<s>");
+                    match self.tokenizer.encode(&prompt, add_special) {
+                        Ok(ids) => ids.into_iter().map(|token| token as i32).collect(),
+                        Err(error) => {
+                            let _ = response_tx
+                                .send(GenerateEvent::Error(format!("Tokenization error: {error}")));
+                            return;
+                        }
+                    }
+                }
+            }
+        } else if let Some(tokens) = prompt_token_ids {
+            let image_placeholders = tokens
+                .iter()
+                .filter(|&&token| token == crate::phi4_siglip_prompt::PHI4_SIGLIP_IMAGE_TOKEN_INDEX)
+                .count();
+            let audio_placeholders = tokens
+                .iter()
+                .filter(|&&token| token == crate::phi4mm_prompt::PHI4MM_AUDIO_TOKEN_ID)
+                .count();
+            if image_placeholders != images.len() || audio_placeholders != audio.len() {
+                let _ = response_tx.send(GenerateEvent::Error(format!(
+                    "Phi4MM token ids contain {image_placeholders} image/{audio_placeholders} \
+                     audio placeholders for {}/{} inputs",
+                    images.len(),
+                    audio.len(),
+                )));
+                return;
+            }
+            tokens
+        } else {
+            let mut tokenization_error = None;
+            let prepared = crate::phi4mm_prompt::prepare_phi4mm_prompt_tokens(
+                &prompt,
+                images.len(),
+                audio.len(),
+                |text, add_special| match self.tokenizer.encode(text, add_special) {
+                    Ok(tokens) => tokens.into_iter().map(|token| token as i32).collect(),
                     Err(error) => {
-                        let _ = response_tx
-                            .send(GenerateEvent::Error(format!("Tokenization error: {error}")));
-                        return;
+                        tokenization_error = Some(error.to_string());
+                        Vec::new()
                     }
+                },
+            );
+            if let Some(error) = tokenization_error {
+                let _ =
+                    response_tx.send(GenerateEvent::Error(format!("Tokenization error: {error}")));
+                return;
+            }
+            match prepared {
+                Ok(prepared) => prepared.tokens,
+                Err(error) => {
+                    let _ = response_tx.send(GenerateEvent::Error(format!(
+                        "Phi4MM prompt normalization failed: {error}"
+                    )));
+                    return;
                 }
             }
         };
         #[cfg(feature = "xla-diagnostics")]
-        if !images.is_empty()
+        if !uses_phi4mm_media
+            && !images.is_empty()
             && let Err(error) = validate_reference_prompt_tokens_from_env(&prompt_tokens)
         {
             let _ = response_tx.send(GenerateEvent::Error(error));
@@ -86,6 +146,72 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
         let stop_sequences = options.stop_sequences.unwrap_or_default();
         let start = Instant::now();
 
+        if uses_phi4mm_media {
+            let Some(stage) = self.audio_preprocessor.as_ref() else {
+                let _ = response_tx.send(GenerateEvent::Error(
+                    "the loaded OpenXLA model/runtime bundle does not support Phi4MM media input"
+                        .to_string(),
+                ));
+                return;
+            };
+            let Some(policy) = self.audio_policy else {
+                let _ = response_tx.send(GenerateEvent::Error(
+                    "the OpenXLA audio preprocessing policy is unavailable".to_string(),
+                ));
+                return;
+            };
+            let Some(next_job_id) = self.next_audio_job_id.checked_add(1) else {
+                let _ = response_tx.send(GenerateEvent::Error(
+                    "OpenXLA audio preprocessing request id overflowed".to_string(),
+                ));
+                return;
+            };
+            let job_id = self.next_audio_job_id;
+            self.next_audio_job_id = next_job_id;
+            let clips = audio
+                .into_iter()
+                .enumerate()
+                .map(|(index, bytes)| crate::audio::AudioEncodedClip {
+                    bytes,
+                    source: crate::audio::AudioSourceKind::ServerInline,
+                    placeholder_ordinal: index + 1,
+                })
+                .collect();
+            let job = AudioPreprocessJob {
+                job_id,
+                token_ids: prompt_tokens.clone(),
+                max_prefill_tokens: self.context_capacity,
+                expected_image_count: media.declared_images,
+                images,
+                clips,
+                policy,
+                cancelled: cancelled.clone(),
+            };
+            match stage.try_submit(job) {
+                Ok(()) => {
+                    self.pending_audio.insert(
+                        job_id,
+                        PendingAudioState {
+                            response_tx,
+                            cancelled,
+                            prompt_tokens,
+                            params,
+                            stop_sequences,
+                            max_tokens: options.max_tokens,
+                            start,
+                        },
+                    );
+                }
+                Err(AudioQueueError::Cancelled { .. }) => {}
+                Err(error) => {
+                    let _ = response_tx.send(GenerateEvent::Error(format!(
+                        "OpenXLA audio preprocessing admission failed: {error}"
+                    )));
+                }
+            }
+            return;
+        }
+
         if images.is_empty() {
             self.submit_text(
                 prompt_tokens,
@@ -257,15 +383,68 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
         }
     }
 
+    fn handle_audio_preprocessed(&mut self, result: AudioPreprocessResult) {
+        let Some(state) = self.pending_audio.remove(&result.job_id) else {
+            return;
+        };
+        if state.cancelled.load(Ordering::Acquire) {
+            return;
+        }
+        let prepared = match result.outcome {
+            AudioPreprocessOutcome::Prepared(prepared) => prepared,
+            AudioPreprocessOutcome::Cancelled(_) => return,
+            AudioPreprocessOutcome::Failed(error) => {
+                let _ = state.response_tx.send(GenerateEvent::Error(format!(
+                    "OpenXLA audio preprocessing failed: {error}"
+                )));
+                return;
+            }
+        };
+        if state.max_tokens == 0 {
+            self.finish_zero_budget(state.prompt_tokens, state.start, state.response_tx);
+            return;
+        }
+        let effective_prefill_len = prepared.sequence_len;
+        let prompt_token_count = state.prompt_tokens.len();
+        match self
+            .engine
+            .submit_prepared(prepared, state.max_tokens, state.params)
+        {
+            Ok(req_id) => {
+                self.batch_observability
+                    .record_prefill_start(effective_prefill_len);
+                self.states.insert(
+                    req_id,
+                    ServeState {
+                        response_tx: state.response_tx,
+                        cancelled: state.cancelled,
+                        detok: StreamingDecodeState::new(&self.tokenizer, &state.prompt_tokens),
+                        stop: StopMatcher::new(state.stop_sequences),
+                        start: state.start,
+                        prompt_token_count,
+                        effective_prefill_len,
+                        max_tokens: state.max_tokens,
+                        generated_tokens: 0,
+                    },
+                );
+            }
+            Err(error) => {
+                let _ = state.response_tx.send(GenerateEvent::Error(format!(
+                    "OpenXLA audio prepared-prefill admission failed: {error}"
+                )));
+            }
+        }
+    }
+
     pub(super) fn drain_preprocessed(&mut self) {
         loop {
             let received = match self.image_preprocessor.as_ref() {
                 Some(stage) => stage.try_recv(),
-                None => return,
+                None => break,
             };
             match received {
                 Ok(result) => self.handle_preprocessed(result),
-                Err(mpsc::TryRecvError::Empty) => return,
+                Err(mpsc::TryRecvError::Empty) => break,
                 Err(mpsc::TryRecvError::Disconnected) => {
                     self.image_preprocessor = None;
                     for (_, state) in self.pending_images.drain() {
@@ -273,6 +452,26 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
                             "OpenXLA image preprocessor stopped unexpectedly".to_string(),
                         ));
                     }
+                    break;
+                }
+            }
+        }
+        loop {
+            let received = match self.audio_preprocessor.as_ref() {
+                Some(stage) => stage.try_recv(),
+                None => return,
+            };
+            match received {
+                Ok(result) => self.handle_audio_preprocessed(result),
+                Err(mpsc::TryRecvError::Empty) => return,
+                Err(mpsc::TryRecvError::Disconnected) => {
+                    self.audio_preprocessor = None;
+                    self.audio_policy = None;
+                    for (_, state) in self.pending_audio.drain() {
+                        let _ = state.response_tx.send(GenerateEvent::Error(
+                            "OpenXLA audio preprocessor stopped unexpectedly".to_string(),
+                        ));
+                    }
                     return;
                 }
             }
@@ -308,6 +507,10 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
                     state.cancelled.store(true, Ordering::Release);
                 }
                 self.pending_images.clear();
+                for state in self.pending_audio.values() {
+                    state.cancelled.store(true, Ordering::Release);
+                }
+                self.pending_audio.clear();
             }
         }
     }
@@ -317,7 +520,7 @@ impl<E: XlaServingEngine> XlaServeWorker<E> {
     /// without delaying completion or active decode rows.
     pub(super) fn drain_incoming(&mut self, block: bool) {
         if block {
-            let request = if self.pending_images.is_empty() {
+            let request = if self.pending_images.is_empty() && self.pending_audio.is_empty() {
                 match self.request_rx.recv() {
                     Ok(request) => request,
                     Err(_) => {
diff --git a/src/server/batch/xla_worker_tests.rs b/src/server/batch/xla_worker_tests.rs
index 8aa3c3c1e..39307d3da 100644
--- a/src/server/batch/xla_worker_tests.rs
+++ b/src/server/batch/xla_worker_tests.rs
@@ -19,9 +19,16 @@ use std::sync::{Arc, mpsc};
 use std::time::Duration;
 
 use image::{DynamicImage, ImageFormat};
+use mlxcel_core::session::{
+    OwnedTensor, PreparedAdapterMode, PreparedAttentionBias, PreparedPositions, PreparedTensorDType,
+};
 
 use super::*;
 use crate::FakeHostMultimodalPreprocessor;
+use crate::audio::{AudioFamilyPolicy, AudioWaveformBatch};
+use crate::server::batch::xla_audio_preprocess::{
+    AudioFeatureProducer, AudioPreprocessLimits, AudioPreprocessStage,
+};
 use crate::server::request_options::{RequestOptionOverrides, build_server_generate_options};
 use crate::server::{GenerationResult, ServerConfig};
 
@@ -31,6 +38,7 @@ struct FakeServingEngine {
     active: BTreeSet<u64>,
     text_submissions: Vec<Vec<i32>>,
     prepared_submissions: Vec<(Vec<i32>, usize)>,
+    prepared_modes: Vec<PreparedAdapterMode>,
 }
 
 impl XlaServingEngine for FakeServingEngine {
@@ -66,6 +74,7 @@ impl XlaServingEngine for FakeServingEngine {
         _max_new_tokens: usize,
         _params: SampleParams,
     ) -> Result<u64, String> {
+        self.prepared_modes.push(prepared.adapter_mode);
         self.prepared_submissions
             .push((prepared.token_ids, prepared.sequence_len));
         Ok(self.activate())
@@ -107,6 +116,28 @@ fn png_bytes() -> Vec<u8> {
     bytes
 }
 
+fn wav_bytes() -> Vec<u8> {
+    let samples = [0i16; 320];
+    let data_len = samples.len() * 2;
+    let mut wav = Vec::with_capacity(44 + data_len);
+    wav.extend_from_slice(b"RIFF");
+    wav.extend_from_slice(&(36u32 + data_len as u32).to_le_bytes());
+    wav.extend_from_slice(b"WAVEfmt ");
+    wav.extend_from_slice(&16u32.to_le_bytes());
+    wav.extend_from_slice(&1u16.to_le_bytes());
+    wav.extend_from_slice(&1u16.to_le_bytes());
+    wav.extend_from_slice(&16_000u32.to_le_bytes());
+    wav.extend_from_slice(&32_000u32.to_le_bytes());
+    wav.extend_from_slice(&2u16.to_le_bytes());
+    wav.extend_from_slice(&16u16.to_le_bytes());
+    wav.extend_from_slice(b"data");
+    wav.extend_from_slice(&(data_len as u32).to_le_bytes());
+    for sample in samples {
+        wav.extend_from_slice(&sample.to_le_bytes());
+    }
+    wav
+}
+
 fn options(max_tokens: usize) -> ServerGenerateOptions {
     build_server_generate_options(
         &ServerConfig::default(),
@@ -146,6 +177,89 @@ fn worker(image_preprocessor: Option<ImagePreprocessStage>) -> XlaServeWorker<Fa
         image_preprocessor,
         pending_images: HashMap::new(),
         next_image_job_id: 0,
+        audio_preprocessor: None,
+        phi4mm_media_preprocessor: false,
+        audio_policy: None,
+        pending_audio: HashMap::new(),
+        next_audio_job_id: 0,
+        context_capacity: 32,
+        shutdown: false,
+    }
+}
+
+struct FakePhi4MediaProducer;
+
+impl AudioFeatureProducer for FakePhi4MediaProducer {
+    fn prepare(
+        &mut self,
+        waveforms: AudioWaveformBatch,
+        token_ids: Vec<i32>,
+        images: Vec<DynamicImage>,
+        _cancelled: &AtomicBool,
+    ) -> Result<PreparedPrefill, String> {
+        let sequence = token_ids.len();
+        let embeddings = OwnedTensor::new(
+            vec![0; sequence * 2 * std::mem::size_of::<f32>()],
+            PreparedTensorDType::Float32,
+            vec![1, sequence, 2],
+        )
+        .map_err(|error| error.to_string())?;
+        let bias = OwnedTensor::new(
+            vec![0; sequence * std::mem::size_of::<f32>()],
+            PreparedTensorDType::Float32,
+            vec![1, 1, 1, sequence],
+        )
+        .map_err(|error| error.to_string())?;
+        let mode = if !images.is_empty() {
+            PreparedAdapterMode::Vision
+        } else if !waveforms.clips.is_empty() {
+            PreparedAdapterMode::Speech
+        } else {
+            PreparedAdapterMode::Language
+        };
+        PreparedPrefill::new(
+            token_ids,
+            embeddings,
+            PreparedPositions::Sequential {
+                start: 0,
+                length: sequence,
+            },
+            PreparedAttentionBias {
+                tensor: bias,
+                causal: true,
+            },
+            Vec::new(),
+        )
+        .map(|prepared| prepared.with_adapter_mode(mode))
+        .map_err(|error| error.to_string())
+    }
+}
+
+fn phi4mm_worker() -> XlaServeWorker<FakeServingEngine> {
+    let (_request_tx, request_rx) = mpsc::channel();
+    let observability = Arc::new(BatchObservability::new());
+    let stage = AudioPreprocessStage::spawn(
+        FakePhi4MediaProducer,
+        AudioPreprocessLimits::default(),
+        observability.clone(),
+    )
+    .unwrap();
+    XlaServeWorker {
+        engine: FakeServingEngine::default(),
+        tokenizer: MlxcelTokenizer::stub(),
+        request_rx,
+        states: HashMap::new(),
+        batch_metrics: Arc::new(BatchMetrics::new()),
+        batch_observability: observability,
+        image_preprocessor: None,
+        pending_images: HashMap::new(),
+        next_image_job_id: 0,
+        audio_preprocessor: Some(stage),
+        phi4mm_media_preprocessor: true,
+        audio_policy: Some(AudioFamilyPolicy::phi4mm()),
+        pending_audio: HashMap::new(),
+        next_audio_job_id: 0,
+        context_capacity: 32,
         shutdown: false,
     }
 }
@@ -153,7 +267,7 @@ fn worker(image_preprocessor: Option<ImagePreprocessStage>) -> XlaServeWorker<Fa
 fn wait_for_preprocessing(worker: &mut XlaServeWorker<FakeServingEngine>) {
     for _ in 0..500 {
         worker.drain_preprocessed();
-        if worker.pending_images.is_empty() {
+        if worker.pending_images.is_empty() && worker.pending_audio.is_empty() {
             return;
         }
         std::thread::sleep(Duration::from_millis(1));
@@ -161,6 +275,61 @@ fn wait_for_preprocessing(worker: &mut XlaServeWorker<FakeServingEngine>) {
     panic!("timed out waiting for image preprocessing");
 }
 
+#[test]
+fn phi4mm_media_admission_selects_speech_or_vision_per_request() {
+    let mut worker = phi4mm_worker();
+    let requests = [
+        (
+            vec![crate::phi4mm_prompt::PHI4MM_AUDIO_TOKEN_ID],
+            Vec::new(),
+            vec![wav_bytes()],
+            PreparedAdapterMode::Speech,
+        ),
+        (
+            vec![crate::phi4_siglip_prompt::PHI4_SIGLIP_IMAGE_TOKEN_INDEX],
+            vec![png_bytes()],
+            Vec::new(),
+            PreparedAdapterMode::Vision,
+        ),
+        (
+            vec![
+                crate::phi4_siglip_prompt::PHI4_SIGLIP_IMAGE_TOKEN_INDEX,
+                crate::phi4mm_prompt::PHI4MM_AUDIO_TOKEN_ID,
+            ],
+            vec![png_bytes()],
+            vec![wav_bytes()],
+            PreparedAdapterMode::Vision,
+        ),
+    ];
+    let mut receivers = Vec::new();
+    for (tokens, images, audio, _mode) in &requests {
+        let (tx, rx) = mpsc::channel();
+        worker.admit(
+            String::new(),
+            Some(tokens.clone()),
+            options(1),
+            images.clone(),
+            audio.clone(),
+            Vec::new(),
+            media(images.len(), audio.len(), 0),
+            tx,
+            Arc::new(AtomicBool::new(false)),
+        );
+        wait_for_preprocessing(&mut worker);
+        receivers.push(rx);
+    }
+    assert_eq!(
+        worker.engine.prepared_modes,
+        requests.iter().map(|request| request.3).collect::<Vec<_>>()
+    );
+    assert_eq!(worker.states.len(), 3);
+    let events = worker.engine.pump().unwrap();
+    worker.dispatch(events);
+    for receiver in receivers {
+        assert_eq!(receive_done(&receiver).completion_tokens, 1);
+    }
+}
+
 fn receive_done(rx: &mpsc::Receiver<GenerateEvent>) -> GenerationResult {
     loop {
         match rx.recv_timeout(Duration::from_secs(1)).unwrap() {
@@ -470,6 +639,12 @@ fn pending_preprocess_poll_timeout_does_not_shutdown_worker() {
         image_preprocessor: Some(image_stage()),
         pending_images: HashMap::new(),
         next_image_job_id: 0,
+        audio_preprocessor: None,
+        phi4mm_media_preprocessor: false,
+        audio_policy: None,
+        pending_audio: HashMap::new(),
+        next_audio_job_id: 0,
+        context_capacity: 32,
         shutdown: false,
     };
     let (response_tx, _response_rx) = mpsc::channel();
diff --git a/src/server/media.rs b/src/server/media.rs
index 8b685f9d8..088528adf 100644
--- a/src/server/media.rs
+++ b/src/server/media.rs
@@ -180,7 +180,18 @@ impl MediaRequestMetadata {
         audio: usize,
         videos: usize,
     ) -> Result<(), String> {
-        if self.declared_audio > 0 {
+        self.validate_xla_raw_counts_with_audio(images, audio, videos, false)
+    }
+
+    #[cfg_attr(not(feature = "xla-iree"), allow(dead_code))]
+    pub(crate) fn validate_xla_raw_counts_with_audio(
+        self,
+        images: usize,
+        audio: usize,
+        videos: usize,
+        supports_audio: bool,
+    ) -> Result<(), String> {
+        if self.declared_audio > 0 && !supports_audio {
             return Err(format!(
                 "the OpenXLA backend does not support audio input yet ({} declared)",
                 self.declared_audio
@@ -211,6 +222,13 @@ impl MediaRequestMetadata {
                 self.declared_images, self.resolved_images
             ));
         }
+        if self.declared_audio != self.resolved_audio {
+            return Err(format!(
+                "OpenXLA audio resolution cardinality mismatch: {} audio input(s) declared, \
+                 {} raw payload(s) resolved; refusing text fallback",
+                self.declared_audio, self.resolved_audio
+            ));
+        }
         Ok(())
     }
 }
diff --git a/src/server/model_worker.rs b/src/server/model_worker.rs
index a79317938..f0ae3f56c 100644
--- a/src/server/model_worker.rs
+++ b/src/server/model_worker.rs
@@ -1020,6 +1020,7 @@ pub(crate) fn spawn_xla_model_worker(
                 engine,
                 tokenizer,
                 model_path,
+                device.clone(),
                 request_rx,
                 batch_metrics,
                 batch_observability,