diff --git a/.github/workflows/stdarch.yml b/.github/workflows/stdarch.yml
index 17d6449c85e..fdede2bb1f8 100644
--- a/.github/workflows/stdarch.yml
+++ b/.github/workflows/stdarch.yml
@@ -95,15 +95,15 @@ jobs:
     - name: Run stdarch tests
       if: ${{ !matrix.cargo_runner }}
       run: |
-        # FIXME: remove --skip test_tile_ and --skip --skip test__tile when it's implemented.
-        ./y.sh test --release --stdarch-tests -- --skip test_tile_ --skip test__tile
+        # Compiler-allocated __tile_* APIs are not supported yet.
+        ./y.sh test --release --stdarch-tests -- --skip test__tile
 
     - name: Run stdarch tests
       if: ${{ matrix.cargo_runner }}
       run: |
         # FIXME: these tests fail when the sysroot is compiled with LTO because of a missing symbol in proc-macro.
-        # FIXME: remove --skip test_tile_ and --skip --skip test__tile when it's implemented.
-        STDARCH_TEST_SKIP_FUNCTION="xsave,xsaveopt,xsave64,xsaveopt64" STDARCH_TEST_EVERYTHING=1 CHANNEL=release CARGO_TARGET_X86_64_UNKNOWN_LINUX_GNU_RUNNER="${{ matrix.cargo_runner }}" TARGET=x86_64-unknown-linux-gnu CG_RUSTFLAGS="-Ainternal_features" ./y.sh cargo test --manifest-path build/build_sysroot/sysroot_src/library/stdarch/Cargo.toml -- --skip rtm --skip tbm --skip sse4a --skip test_tile_ --skip test__tile
+        # Compiler-allocated __tile_* APIs are not supported yet.
+        STDARCH_TEST_SKIP_FUNCTION="xsave,xsaveopt,xsave64,xsaveopt64" STDARCH_TEST_EVERYTHING=1 CHANNEL=release CARGO_TARGET_X86_64_UNKNOWN_LINUX_GNU_RUNNER="${{ matrix.cargo_runner }}" TARGET=x86_64-unknown-linux-gnu CG_RUSTFLAGS="-Ainternal_features" ./y.sh cargo test --manifest-path build/build_sysroot/sysroot_src/library/stdarch/Cargo.toml -- --skip rtm --skip tbm --skip sse4a --skip test__tile
 
   # Summary job for the merge queue.
   # ALL THE PREVIOUS JOBS NEED TO BE ADDED TO THE `needs` SECTION OF THIS JOB!
diff --git a/src/intrinsic/llvm.rs b/src/intrinsic/llvm.rs
index ef381715c1e..5a6073a46c1 100644
--- a/src/intrinsic/llvm.rs
+++ b/src/intrinsic/llvm.rs
@@ -1,12 +1,140 @@
 use std::borrow::Cow;
 
 use gccjit::{CType, Context, Field, Function, FunctionPtrType, RValue, ToRValue, Type};
+#[cfg(feature = "master")]
+use rustc_codegen_ssa::mir::operand::OperandRef;
 use rustc_codegen_ssa::traits::BuilderMethods;
+#[cfg(feature = "master")]
+use rustc_codegen_ssa::traits::LayoutTypeCodegenMethods;
+#[cfg(feature = "master")]
+use rustc_middle::ty;
+#[cfg(feature = "master")]
+use rustc_middle::ty::layout::LayoutOf;
 
 use crate::builder::Builder;
 use crate::context::{CodegenCx, new_array_type};
 use crate::type_::{StructAttribute, apply_struct_attributes};
 
+/// Lower AMX operations whose tile operands name architectural registers, not SSA values.
+#[cfg(feature = "master")]
+pub(super) fn codegen_x86_amx<'a, 'gcc, 'tcx>(
+    builder: &Builder<'a, 'gcc, 'tcx>,
+    instance: ty::Instance<'tcx>,
+    name: &str,
+    args: &[OperandRef<'tcx, RValue<'gcc>>],
+) -> Option<RValue<'gcc>> {
+    macro_rules! tile_load {
+        ($mnemonic:literal) => {
+            (
+                concat!($mnemonic, " {(%1,%2,1), %%tmm%c0|tmm%c0, [%1+%2*1]}"),
+                &["i", "r", "r"][..],
+                true,
+                false,
+            )
+        };
+    }
+    macro_rules! tile_dot_product {
+        ($mnemonic:literal) => {
+            (
+                concat!($mnemonic, " {%%tmm%c2, %%tmm%c1, %%tmm%c0|tmm%c0, tmm%c1, tmm%c2}"),
+                &["i", "i", "i"][..],
+                false,
+                false,
+            )
+        };
+    }
+    macro_rules! tile_row {
+        ($mnemonic:literal, $row:literal, $constraint:literal) => {
+            (
+                concat!($mnemonic, " {", $row, ", %%tmm%c1, %0|%0, tmm%c1, ", $row, "}"),
+                &["i", $constraint][..],
+                false,
+                true,
+            )
+        };
+    }
+
+    let (template, constraints, accesses_memory, returns_vector): (&str, &[&str], bool, bool) =
+        match name {
+            "llvm.x86.tileloadd64" => tile_load!("tileloadd"),
+            "llvm.x86.tileloaddt164" => tile_load!("tileloaddt1"),
+            "llvm.x86.tileloaddrs64" => tile_load!("tileloaddrs"),
+            "llvm.x86.tileloaddrst164" => tile_load!("tileloaddrst1"),
+            "llvm.x86.tilestored64" => (
+                "tilestored {%%tmm%c0, (%1,%2,1)|[%1+%2*1], tmm%c0}",
+                &["i", "r", "r"],
+                true,
+                false,
+            ),
+            "llvm.x86.tilezero" => ("tilezero {%%tmm%c0|tmm%c0}", &["i"], false, false),
+            "llvm.x86.tilerelease" => ("tilerelease", &[], false, false),
+            "llvm.x86.tdpbf16ps" => tile_dot_product!("tdpbf16ps"),
+            "llvm.x86.tdpbssd" => tile_dot_product!("tdpbssd"),
+            "llvm.x86.tdpbsud" => tile_dot_product!("tdpbsud"),
+            "llvm.x86.tdpbusd" => tile_dot_product!("tdpbusd"),
+            "llvm.x86.tdpbuud" => tile_dot_product!("tdpbuud"),
+            "llvm.x86.tdpfp16ps" => tile_dot_product!("tdpfp16ps"),
+            "llvm.x86.tcmmimfp16ps" => tile_dot_product!("tcmmimfp16ps"),
+            "llvm.x86.tcmmrlfp16ps" => tile_dot_product!("tcmmrlfp16ps"),
+            "llvm.x86.tdpbf8ps" => tile_dot_product!("tdpbf8ps"),
+            "llvm.x86.tdpbhf8ps" => tile_dot_product!("tdpbhf8ps"),
+            "llvm.x86.tdphbf8ps" => tile_dot_product!("tdphbf8ps"),
+            "llvm.x86.tdphf8ps" => tile_dot_product!("tdphf8ps"),
+            // The output occupies %0, shifting the tile and row inputs to %1 and %2.
+            // %k prints a 32-bit row register; plain %2 preserves the immediate prefix for each dialect.
+            "llvm.x86.tilemovrow" => tile_row!("tilemovrow", "%k2", "r"),
+            "llvm.x86.tilemovrowi" => tile_row!("tilemovrow", "%2", "i"),
+            "llvm.x86.tcvtrowd2ps" => tile_row!("tcvtrowd2ps", "%k2", "r"),
+            "llvm.x86.tcvtrowd2psi" => tile_row!("tcvtrowd2ps", "%2", "i"),
+            "llvm.x86.tcvtrowps2phh" => tile_row!("tcvtrowps2phh", "%k2", "r"),
+            "llvm.x86.tcvtrowps2phhi" => tile_row!("tcvtrowps2phh", "%2", "i"),
+            "llvm.x86.tcvtrowps2phl" => tile_row!("tcvtrowps2phl", "%k2", "r"),
+            "llvm.x86.tcvtrowps2phli" => tile_row!("tcvtrowps2phl", "%2", "i"),
+            "llvm.x86.tcvtrowps2bf16h" => tile_row!("tcvtrowps2bf16h", "%k2", "r"),
+            "llvm.x86.tcvtrowps2bf16hi" => tile_row!("tcvtrowps2bf16h", "%2", "i"),
+            "llvm.x86.tcvtrowps2bf16l" => tile_row!("tcvtrowps2bf16l", "%k2", "r"),
+            "llvm.x86.tcvtrowps2bf16li" => tile_row!("tcvtrowps2bf16l", "%2", "i"),
+            _ => return None,
+        };
+
+    let result = if returns_vector {
+        // LLVM intrinsics have no ordinary call ABI. Query the declared result layout directly.
+        let sig = builder
+            .tcx
+            .fn_sig(instance.def_id())
+            .instantiate(builder.tcx, instance.args)
+            .skip_norm_wip();
+        let sig = builder.tcx.instantiate_bound_regions_with_erased(sig);
+        let result_type = builder.backend_type(builder.layout_of(sig.output()));
+        Some(builder.current_func().new_local(builder.location, result_type, "amx_row"))
+    } else {
+        None
+    };
+    // Create the output before recording the asm: libgccjit replays nodes in creation order.
+    let asm = builder.llbb().add_extended_asm(builder.location, template);
+    // GCC does not allocate these tile registers. Keep their implicit state changes, including
+    // operations with no memory effects, ordered with the other AMX operations.
+    asm.set_volatile_flag(true);
+    if let Some(result) = result {
+        asm.add_output_operand(None, "=v", result);
+    }
+    assert_eq!(args.len(), constraints.len());
+    for (arg, constraint) in args.iter().zip(constraints) {
+        asm.add_input_operand(None, constraint, arg.immediate());
+    }
+    if accesses_memory {
+        // A register operand for the base address does not describe the memory being accessed.
+        // The extent depends on TILECFG and the runtime stride, so a fixed-size memory operand
+        // would be incorrect. Loads also need this barrier to retain preceding buffer writes.
+        asm.add_clobber("memory");
+    }
+    Some(match result {
+        Some(result) => result.to_rvalue(),
+        // Match the value returned for a void builtin without querying its ABI.
+        None => builder.context.new_rvalue_zero(builder.isize_type),
+    })
+}
+
 fn encode_key_128_type<'a, 'gcc, 'tcx>(
     builder: &Builder<'a, 'gcc, 'tcx>,
 ) -> (Type<'gcc>, Field<'gcc>, Field<'gcc>) {
@@ -120,7 +248,7 @@ pub fn adjust_intrinsic_arguments<'a, 'b, 'gcc, 'tcx>(
     mut args: Cow<'b, [RValue<'gcc>]>,
     func_name: &str,
 ) -> Cow<'b, [RValue<'gcc>]> {
-    // FIXME: this might not be a good way to workaround the missing tile builtins.
+    // Discard arguments when an unsupported intrinsic is lowered to a trap.
     if func_name == "__builtin_trap" {
         return vec![].into();
     }
@@ -1655,67 +1783,40 @@ pub fn intrinsic<'gcc, 'tcx>(name: &str, cx: &CodegenCx<'gcc, 'tcx>) -> Function
         "llvm.x86.avx512.fpclass.pd.512" => "__builtin_ia32_fpclasspd512_mask",
         "llvm.x86.avx512.fpclass.ps.512" => "__builtin_ia32_fpclassps512_mask",
 
-        // FIXME: support the tile builtins:
-        "llvm.x86.ldtilecfg" => "__builtin_trap",
-        "llvm.x86.sttilecfg" => "__builtin_trap",
-        "llvm.x86.tileloadd64" => "__builtin_trap",
+        // GCC's configuration builtins model the full 64-byte memory operand.
+        "llvm.x86.ldtilecfg" => "__builtin_ia32_ldtilecfg",
+        "llvm.x86.sttilecfg" => "__builtin_ia32_sttilecfg",
+
+        // FIXME: support compiler-allocated tiles (.internal), used by Rust's __tile_* APIs.
         "llvm.x86.tileloadd64.internal" => "__builtin_trap",
-        "llvm.x86.tilerelease" => "__builtin_trap",
-        "llvm.x86.tilestored64" => "__builtin_trap",
         "llvm.x86.tilestored64.internal" => "__builtin_trap",
-        "llvm.x86.tileloaddrs64" => "__builtin_trap",
         "llvm.x86.tileloaddrs64.internal" => "__builtin_trap",
-        "llvm.x86.tileloaddt164" => "__builtin_trap",
         "llvm.x86.tileloaddt164.internal" => "__builtin_trap",
-        "llvm.x86.tileloaddrst164" => "__builtin_trap",
         "llvm.x86.tileloaddrst164.internal" => "__builtin_trap",
-        "llvm.x86.tilezero" => "__builtin_trap",
         "llvm.x86.tilezero.internal" => "__builtin_trap",
-        "llvm.x86.tilemovrow" => "__builtin_trap",
         "llvm.x86.tilemovrow.internal" => "__builtin_trap",
-        "llvm.x86.tilemovrowi" => "__builtin_trap",
-        "llvm.x86.tdpbhf8ps" => "__builtin_trap",
         "llvm.x86.tdpbhf8ps.internal" => "__builtin_trap",
-        "llvm.x86.tdphbf8ps" => "__builtin_trap",
         "llvm.x86.tdphbf8ps.internal" => "__builtin_trap",
-        "llvm.x86.tdpbf8ps" => "__builtin_trap",
         "llvm.x86.tdpbf8ps.internal" => "__builtin_trap",
-        "llvm.x86.tdphf8ps" => "__builtin_trap",
         "llvm.x86.tdphf8ps.internal" => "__builtin_trap",
-        "llvm.x86.tdpbf16ps" => "__builtin_trap",
         "llvm.x86.tdpbf16ps.internal" => "__builtin_trap",
-        "llvm.x86.tdpbssd" => "__builtin_trap",
         "llvm.x86.tdpbssd.internal" => "__builtin_trap",
-        "llvm.x86.tdpbsud" => "__builtin_trap",
         "llvm.x86.tdpbsud.internal" => "__builtin_trap",
-        "llvm.x86.tdpbusd" => "__builtin_trap",
         "llvm.x86.tdpbusd.internal" => "__builtin_trap",
-        "llvm.x86.tdpbuud" => "__builtin_trap",
         "llvm.x86.tdpbuud.internal" => "__builtin_trap",
-        "llvm.x86.tdpfp16ps" => "__builtin_trap",
         "llvm.x86.tdpfp16ps.internal" => "__builtin_trap",
-        "llvm.x86.tmmultf32ps" => "__builtin_trap",
-        "llvm.x86.tmmultf32ps.internal" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2phh" => "__builtin_trap",
         "llvm.x86.tcvtrowps2phh.internal" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2phl" => "__builtin_trap",
         "llvm.x86.tcvtrowps2phl.internal" => "__builtin_trap",
-        "llvm.x86.tcvtrowd2ps" => "__builtin_trap",
         "llvm.x86.tcvtrowd2ps.internal" => "__builtin_trap",
-        "llvm.x86.tcvtrowd2psi" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2phhi" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2phli" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2bf16h" => "__builtin_trap",
         "llvm.x86.tcvtrowps2bf16h.internal" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2bf16hi" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2bf16l" => "__builtin_trap",
         "llvm.x86.tcvtrowps2bf16l.internal" => "__builtin_trap",
-        "llvm.x86.tcvtrowps2bf16li" => "__builtin_trap",
-        "llvm.x86.tcmmimfp16ps" => "__builtin_trap",
         "llvm.x86.tcmmimfp16ps.internal" => "__builtin_trap",
-        "llvm.x86.tcmmrlfp16ps" => "__builtin_trap",
         "llvm.x86.tcmmrlfp16ps.internal" => "__builtin_trap",
 
+        // AMX-TF32 support was removed in GCC 17 and is absent from current stdarch.
+        "llvm.x86.tmmultf32ps" => "__builtin_trap",
+        "llvm.x86.tmmultf32ps.internal" => "__builtin_trap",
+
         // NOTE: this file is generated by https://github.com/GuillaumeGomez/llvmint/blob/master/generate_list.py
         _ => map_arch_intrinsic(name),
     };
diff --git a/src/intrinsic/mod.rs b/src/intrinsic/mod.rs
index b6015a74d9e..6e40b3ae2a9 100644
--- a/src/intrinsic/mod.rs
+++ b/src/intrinsic/mod.rs
@@ -591,6 +591,10 @@ impl<'a, 'gcc, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'a, 'gcc, 'tc
             func
         } else {
             let sym = self.tcx.symbol_name(instance).name;
+            #[cfg(feature = "master")]
+            if let Some(result) = llvm::codegen_x86_amx(self, instance, sym, args) {
+                return result;
+            }
 
             let func = if let Some(func) = self.intrinsics.borrow().get(sym) {
                 *func
diff --git a/tests/asm/x86_64-amx.rs b/tests/asm/x86_64-amx.rs
new file mode 100644
index 00000000000..37412d64af9
--- /dev/null
+++ b/tests/asm/x86_64-amx.rs
@@ -0,0 +1,237 @@
+//@ assembly-output: emit-asm
+//@ only-x86_64
+//@ compile-flags: -Copt-level=3
+//@ revisions: att intel
+//@[att] compile-flags: -Cllvm-args=-masm=att
+// cg_gcc defaults to -masm=intel (gcc_util::new_context). It appends -Cllvm-args
+// afterwards, so the att revision overrides that default; intel uses it unchanged.
+
+#![crate_type = "lib"]
+#![no_std]
+#![feature(avx10_target_feature, x86_amx_intrinsics)]
+
+use core::arch::x86_64::*;
+
+// These use the fixed-register APIs, not the __tile* APIs with compiler-allocated tiles.
+// Callers provide a palette-1 configuration with 16 rows and 64 bytes per row for
+// each used tile, and buffers large enough for 16 rows at the supplied byte stride.
+// Functions are alphabetical to match the assembly runner's emission order.
+
+// CHECK-LABEL: {{^"?_?}}amx_float{{"?}}:
+// CHECK: ldtilecfg
+// att: tileloadd {{\(%[a-z0-9]+,}}[[FLOAT_STRIDE:%[a-z0-9]+]]{{(,1)?\)}}, %tmm0
+// intel: tileloadd tmm0, {{\[[a-z0-9]+ *\+ *}}[[FLOAT_STRIDE:[a-z0-9]+]]{{(\*1)?\]}}
+// att: tileloaddt1 {{\(%[a-z0-9]+,}}[[FLOAT_STRIDE]]{{(,1)?\)}}, %tmm3
+// intel: tileloaddt1 tmm3, {{\[[a-z0-9]+ *\+ *}}[[FLOAT_STRIDE]]{{(\*1)?\]}}
+// CHECK: tilezero {{%?}}tmm7
+// att: tdpbf16ps %tmm3, %tmm0, %tmm7
+// intel: tdpbf16ps tmm7, tmm0, tmm3
+// att: tilestored %tmm7, {{\(%[a-z0-9]+,}}[[FLOAT_STRIDE]]{{(,1)?\)}}
+// intel: tilestored {{\[[a-z0-9]+ *\+ *}}[[FLOAT_STRIDE]]{{(\*1)?\]}}, tmm7
+// CHECK: tilezero {{%?}}tmm6
+// att: tdpfp16ps %tmm0, %tmm3, %tmm6
+// intel: tdpfp16ps tmm6, tmm3, tmm0
+// att: tilestored %tmm6, {{\(%[a-z0-9]+,}}[[FLOAT_STRIDE]]{{(,1)?\)}}
+// intel: tilestored {{\[[a-z0-9]+ *\+ *}}[[FLOAT_STRIDE]]{{(\*1)?\]}}, tmm6
+// CHECK: tilezero {{%?}}tmm7
+// att: tcmmimfp16ps %tmm3, %tmm0, %tmm7
+// intel: tcmmimfp16ps tmm7, tmm0, tmm3
+// att: tilestored %tmm7, {{\(%[a-z0-9]+,}}[[FLOAT_STRIDE]]{{(,1)?\)}}
+// intel: tilestored {{\[[a-z0-9]+ *\+ *}}[[FLOAT_STRIDE]]{{(\*1)?\]}}, tmm7
+// CHECK: tilezero {{%?}}tmm6
+// att: tcmmrlfp16ps %tmm3, %tmm0, %tmm6
+// intel: tcmmrlfp16ps tmm6, tmm0, tmm3
+// att: tilestored %tmm6, {{\(%[a-z0-9]+,}}[[FLOAT_STRIDE]]{{(,1)?\)}}
+// intel: tilestored {{\[[a-z0-9]+ *\+ *}}[[FLOAT_STRIDE]]{{(\*1)?\]}}, tmm6
+// CHECK: tilerelease
+#[no_mangle]
+#[target_feature(enable = "amx-tile,amx-bf16,amx-fp16,amx-complex")]
+pub unsafe fn amx_float(
+    config: *const u8,
+    a: *const u8,
+    b: *const u8,
+    bf16: *mut u8,
+    fp16: *mut u8,
+    imaginary: *mut u8,
+    real: *mut u8,
+    stride: usize,
+) {
+    _tile_loadconfig(config);
+    _tile_loadd::<0>(a, stride);
+    _tile_stream_loadd::<3>(b, stride);
+
+    // Interpret the input bits as BF16, FP16 and complex FP16 independently.
+    _tile_zero::<7>();
+    _tile_dpbf16ps::<7, 0, 3>();
+    _tile_stored::<7>(bf16, stride);
+    _tile_zero::<6>();
+    _tile_dpfp16ps::<6, 3, 0>();
+    _tile_stored::<6>(fp16, stride);
+    _tile_zero::<7>();
+    _tile_cmmimfp16ps::<7, 0, 3>();
+    _tile_stored::<7>(imaginary, stride);
+    _tile_zero::<6>();
+    _tile_cmmrlfp16ps::<6, 0, 3>();
+    _tile_stored::<6>(real, stride);
+    _tile_release();
+}
+
+// CHECK-LABEL: {{^"?_?}}amx_fp8{{"?}}:
+// att: tdpbf8ps %tmm2, %tmm1, %tmm7
+// intel: tdpbf8ps tmm7, tmm1, tmm2
+// att: tdpbhf8ps %tmm1, %tmm2, %tmm7
+// intel: tdpbhf8ps tmm7, tmm2, tmm1
+// att: tdphbf8ps %tmm2, %tmm1, %tmm7
+// intel: tdphbf8ps tmm7, tmm1, tmm2
+// att: tdphf8ps %tmm1, %tmm2, %tmm7
+// intel: tdphf8ps tmm7, tmm2, tmm1
+#[no_mangle]
+#[target_feature(enable = "amx-fp8")]
+pub unsafe fn amx_fp8() {
+    _tile_dpbf8ps::<7, 1, 2>();
+    _tile_dpbhf8ps::<7, 2, 1>();
+    _tile_dphbf8ps::<7, 1, 2>();
+    _tile_dphf8ps::<7, 2, 1>();
+}
+
+// CHECK-LABEL: {{^"?_?}}amx_int8{{"?}}:
+// CHECK: ldtilecfg
+// CHECK: sttilecfg
+// att: tileloadd {{\(%[a-z0-9]+,}}[[INT_STRIDE:%[a-z0-9]+]]{{(,1)?\)}}, %tmm1
+// intel: tileloadd tmm1, {{\[[a-z0-9]+ *\+ *}}[[INT_STRIDE:[a-z0-9]+]]{{(\*1)?\]}}
+// att: tileloaddt1 {{\(%[a-z0-9]+,}}[[INT_STRIDE]]{{(,1)?\)}}, %tmm2
+// intel: tileloaddt1 tmm2, {{\[[a-z0-9]+ *\+ *}}[[INT_STRIDE]]{{(\*1)?\]}}
+// CHECK: tilezero {{%?}}tmm7
+// att: tdpbssd %tmm2, %tmm1, %tmm7
+// intel: tdpbssd tmm7, tmm1, tmm2
+// att: tdpbsud %tmm1, %tmm2, %tmm7
+// intel: tdpbsud tmm7, tmm2, tmm1
+// att: tdpbusd %tmm2, %tmm1, %tmm7
+// intel: tdpbusd tmm7, tmm1, tmm2
+// att: tdpbuud %tmm1, %tmm2, %tmm7
+// intel: tdpbuud tmm7, tmm2, tmm1
+// att: tilestored %tmm7, {{\(%[a-z0-9]+,}}[[INT_STRIDE]]{{(,1)?\)}}
+// intel: tilestored {{\[[a-z0-9]+ *\+ *}}[[INT_STRIDE]]{{(\*1)?\]}}, tmm7
+// CHECK: tilerelease
+#[no_mangle]
+#[target_feature(enable = "amx-tile,amx-int8")]
+pub unsafe fn amx_int8(
+    config: *const u8,
+    saved_config: *mut u8,
+    a: *const u8,
+    b: *const u8,
+    output: *mut u8,
+    stride: usize,
+) {
+    _tile_loadconfig(config);
+    _tile_storeconfig(saved_config);
+    _tile_loadd::<1>(a, stride);
+    _tile_stream_loadd::<2>(b, stride);
+    _tile_zero::<7>();
+    // Accumulate all signedness combinations; reversing inputs must not swap
+    // which source is signed in the two mixed-signedness instructions.
+    _tile_dpbssd::<7, 1, 2>();
+    _tile_dpbsud::<7, 2, 1>();
+    _tile_dpbusd::<7, 1, 2>();
+    _tile_dpbuud::<7, 2, 1>();
+    _tile_stored::<7>(output, stride);
+    _tile_release();
+}
+
+// Here config enables only tile 7, with one row of four bytes. Input and output
+// point to separate u32s. No volatile Rust accesses hide missing memory effects:
+// the first input store is dead unless tileloadd reads memory, and the result is
+// the sentinel unless tilestored invalidates the compiler's cached output value.
+// Capture addresses rather than pinning register allocation or stack offsets.
+// CHECK-LABEL: {{^"?_?}}amx_memory{{"?}}:
+// CHECK: ldtilecfg
+// att: movl {{%[a-z0-9]+}}, ([[INPUT:%[a-z0-9]+]])
+// intel: mov DWORD PTR {{\[}}[[INPUT:[a-z0-9]+]]{{\]}}, {{[a-z0-9]+}}
+// att: tileloadd ([[INPUT]],[[MEM_STRIDE:%[a-z0-9]+]]{{(,1)?}}), %tmm7
+// intel: tileloadd tmm7, {{\[}}[[INPUT]]{{ *\+ *}}[[MEM_STRIDE:[a-z0-9]+]]{{(\*1)?\]}}
+// att: movl $0, ([[INPUT]])
+// intel: mov DWORD PTR {{\[}}[[INPUT]]{{\]}}, 0
+// att: movl $305419896, ([[OUTPUT:%[a-z0-9]+]])
+// intel: mov DWORD PTR {{\[}}[[OUTPUT:[a-z0-9]+]]{{\]}}, 305419896
+// att: tilestored %tmm7, ([[OUTPUT]],[[MEM_STRIDE]]{{(,1)?}})
+// intel: tilestored {{\[}}[[OUTPUT]]{{ *\+ *}}[[MEM_STRIDE]]{{(\*1)?\]}}, tmm7
+// att-DAG: movl ([[OUTPUT]]), %eax
+// intel-DAG: mov eax, DWORD PTR {{\[}}[[OUTPUT]]{{\]}}
+// CHECK-DAG: tilerelease
+// CHECK: ret
+#[no_mangle]
+#[target_feature(enable = "amx-tile")]
+pub unsafe extern "C" fn amx_memory(
+    config: *const u8,
+    input: *mut u32,
+    output: *mut u32,
+    value: u32,
+    stride: usize,
+) -> u32 {
+    _tile_loadconfig(config);
+    input.write(value);
+    _tile_loadd::<7>(input.cast(), stride);
+    input.write(0);
+    output.write(0x12345678);
+    _tile_stored::<7>(output.cast(), stride);
+    let result = output.read();
+    _tile_release();
+    result
+}
+
+// CHECK-LABEL: {{^"?_?}}amx_movrs{{"?}}:
+// att: tileloaddrs {{\(%[a-z0-9]+,}}[[RS_STRIDE:%[a-z0-9]+]]{{(,1)?\)}}, %tmm7
+// intel: tileloaddrs tmm7, {{\[[a-z0-9]+ *\+ *}}[[RS_STRIDE:[a-z0-9]+]]{{(\*1)?\]}}
+// att: tileloaddrst1 {{\(%[a-z0-9]+,}}[[RS_STRIDE]]{{(,1)?\)}}, %tmm0
+// intel: tileloaddrst1 tmm0, {{\[[a-z0-9]+ *\+ *}}[[RS_STRIDE]]{{(\*1)?\]}}
+#[no_mangle]
+#[target_feature(enable = "amx-movrs")]
+pub unsafe fn amx_movrs(input: *const u8, stride: usize) {
+    _tile_loaddrs::<7>(input, stride);
+    _tile_stream_loaddrs::<0>(input, stride);
+}
+
+// Row operations return vectors, unlike the other fixed-register intrinsics. Preserve
+// all results in caller-provided storage. Exercise both row encodings and both tile
+// register boundaries; row 15 is the last row in a fully configured tile.
+// CHECK-LABEL: {{^"?_?}}amx_rows{{"?}}:
+// att: tilemovrow [[ROW:%(e[a-z]+|r[0-9]+d)]], %tmm7, {{%zmm[0-9]+}}
+// intel: tilemovrow {{zmm[0-9]+}}, tmm7, [[ROW:(e[a-z]+|r[0-9]+d)]]
+// att: tilemovrow $15, %tmm0, {{%zmm[0-9]+}}
+// intel: tilemovrow {{zmm[0-9]+}}, tmm0, 15
+// att: tcvtrowd2ps [[ROW]], %tmm7, {{%zmm[0-9]+}}
+// intel: tcvtrowd2ps {{zmm[0-9]+}}, tmm7, [[ROW]]
+// att: tcvtrowd2ps $15, %tmm0, {{%zmm[0-9]+}}
+// intel: tcvtrowd2ps {{zmm[0-9]+}}, tmm0, 15
+// att: tcvtrowps2phh [[ROW]], %tmm7, {{%zmm[0-9]+}}
+// intel: tcvtrowps2phh {{zmm[0-9]+}}, tmm7, [[ROW]]
+// att: tcvtrowps2phh $15, %tmm0, {{%zmm[0-9]+}}
+// intel: tcvtrowps2phh {{zmm[0-9]+}}, tmm0, 15
+// att: tcvtrowps2phl [[ROW]], %tmm7, {{%zmm[0-9]+}}
+// intel: tcvtrowps2phl {{zmm[0-9]+}}, tmm7, [[ROW]]
+// att: tcvtrowps2phl $15, %tmm0, {{%zmm[0-9]+}}
+// intel: tcvtrowps2phl {{zmm[0-9]+}}, tmm0, 15
+// att: tcvtrowps2bf16h [[ROW]], %tmm7, {{%zmm[0-9]+}}
+// intel: tcvtrowps2bf16h {{zmm[0-9]+}}, tmm7, [[ROW]]
+// att: tcvtrowps2bf16h $15, %tmm0, {{%zmm[0-9]+}}
+// intel: tcvtrowps2bf16h {{zmm[0-9]+}}, tmm0, 15
+// att: tcvtrowps2bf16l [[ROW]], %tmm7, {{%zmm[0-9]+}}
+// intel: tcvtrowps2bf16l {{zmm[0-9]+}}, tmm7, [[ROW]]
+// att: tcvtrowps2bf16l $15, %tmm0, {{%zmm[0-9]+}}
+// intel: tcvtrowps2bf16l {{zmm[0-9]+}}, tmm0, 15
+#[no_mangle]
+#[target_feature(enable = "amx-avx512,avx10.2")]
+pub unsafe fn amx_rows(output: *mut u8, row: u32) {
+    output.cast::<__m512i>().write_unaligned(_tile_movrow::<7>(row));
+    output.add(64).cast::<__m512i>().write_unaligned(_tile_movrowi::<0, 15>());
+    output.add(128).cast::<__m512>().write_unaligned(_tile_cvtrowd2ps::<7>(row));
+    output.add(192).cast::<__m512>().write_unaligned(_tile_cvtrowd2psi::<0, 15>());
+    output.add(256).cast::<__m512h>().write_unaligned(_tile_cvtrowps2phh::<7>(row));
+    output.add(320).cast::<__m512h>().write_unaligned(_tile_cvtrowps2phhi::<0, 15>());
+    output.add(384).cast::<__m512h>().write_unaligned(_tile_cvtrowps2phl::<7>(row));
+    output.add(448).cast::<__m512h>().write_unaligned(_tile_cvtrowps2phli::<0, 15>());
+    output.add(512).cast::<__m512bh>().write_unaligned(_tile_cvtrowps2bf16h::<7>(row));
+    output.add(576).cast::<__m512bh>().write_unaligned(_tile_cvtrowps2bf16hi::<0, 15>());
+    output.add(640).cast::<__m512bh>().write_unaligned(_tile_cvtrowps2bf16l::<7>(row));
+    output.add(704).cast::<__m512bh>().write_unaligned(_tile_cvtrowps2bf16li::<0, 15>());
+}