diff --git a/Makefile b/Makefile
index c66ac327781f..7d4543a44300 100644
--- a/Makefile
+++ b/Makefile
@@ -545,6 +545,7 @@ SOURCE_FILES = \
   Lower.cpp \
   LowerParallelTasks.cpp \
   LowerSMEStreamingTasks.cpp \
+  LowerStructTypes.cpp \
   LowerWarpShuffles.cpp \
   Memoization.cpp \
   Module.cpp \
@@ -753,6 +754,7 @@ HEADER_FILES = \
   Lower.h \
   LowerParallelTasks.h \
   LowerSMEStreamingTasks.h \
+  LowerStructTypes.h \
   LowerWarpShuffles.h \
   Memoization.h \
   Module.h \
diff --git a/python_bindings/halide/src/halide_/PyEnums.cpp b/python_bindings/halide/src/halide_/PyEnums.cpp
index d6327abab3ba..36353e601f57 100644
--- a/python_bindings/halide/src/halide_/PyEnums.cpp
+++ b/python_bindings/halide/src/halide_/PyEnums.cpp
@@ -250,7 +250,8 @@ void define_enums(py::module &m) {
         .value("Int", Type::Int)
         .value("UInt", Type::UInt)
         .value("Float", Type::Float)
-        .value("Handle", Type::Handle);
+        .value("Handle", Type::Handle)
+        .value("Struct", Type::StructKind);
 
     py::enum_<OutputFileType>(m, "OutputFileType")
         .value("assembly", OutputFileType::assembly)
diff --git a/python_bindings/halide/src/halide_/PyExpr.cpp b/python_bindings/halide/src/halide_/PyExpr.cpp
index 473af42a2d8c..36cc646d5e80 100644
--- a/python_bindings/halide/src/halide_/PyExpr.cpp
+++ b/python_bindings/halide/src/halide_/PyExpr.cpp
@@ -39,6 +39,7 @@ void define_expr(py::module &m) {
             // for implicitly_convertible
             .def(py::init([](const FuncRef &f) -> Expr { return f; }))
             .def(py::init([](const FuncTupleElementRef &f) -> Expr { return f; }))
+            .def(py::init([](const FieldRef &f) -> Expr { return f; }))
             .def(py::init([](const Param<> &p) -> Expr { return p; }))
             .def(py::init([](const RDom &r) -> Expr { return r; }))
             .def(py::init([](const RVar &r) -> Expr { return r; }))
@@ -77,6 +78,7 @@ void define_expr(py::module &m) {
     // There must be an Expr() ctor available for each of these
     py::implicitly_convertible<FuncRef, Expr>();
     py::implicitly_convertible<FuncTupleElementRef, Expr>();
+    py::implicitly_convertible<FieldRef, Expr>();
     py::implicitly_convertible<Param<>, Expr>();
     py::implicitly_convertible<RDom, Expr>();
     py::implicitly_convertible<RVar, Expr>();
diff --git a/python_bindings/halide/src/halide_/PyIROperator.cpp b/python_bindings/halide/src/halide_/PyIROperator.cpp
index 55e19786b7cf..d7861432de99 100644
--- a/python_bindings/halide/src/halide_/PyIROperator.cpp
+++ b/python_bindings/halide/src/halide_/PyIROperator.cpp
@@ -1,7 +1,9 @@
 #include "PyIROperator.h"
 
+#include <pybind11/functional.h>
 #include <utility>
 
+#include "PyBinaryOperators.h"
 #include "PyTuple.h"
 
 namespace Halide {
@@ -72,6 +74,39 @@ py::object py_select(const py::args &args) {
 }  // namespace
 
 void define_operators(py::module &m) {
+    auto field_ref_class =
+        py::class_<FieldRef>(m, "FieldRef")
+            .def("__getitem__", &FieldRef::operator[])
+            .def("size", &FieldRef::size)
+            .def("__len__", &FieldRef::size);
+    add_binary_operators(field_ref_class);
+
+    m.def("field", static_cast<FieldRef (*)(const Expr &, const std::string &)>(&field),
+          py::arg("struct_value"), py::arg("name"));
+    m.def("field", static_cast<FieldRef (*)(const Expr &, int)>(&field),
+          py::arg("struct_value"), py::arg("index"));
+    // Per-field pack_struct: one initializer per struct field, in declaration
+    // order. Each item is an Expr (a scalar field's value; a gather() packet or
+    // a single expression with one `_` placeholder swept over an array field's
+    // extent) or a FieldRef (to copy a whole same-typed field). Mirrors the C++
+    // per-field pack_struct overload.
+    m.def(
+        "pack_struct",
+        [](const Type &type, const py::iterable &fields) -> Expr {
+            std::vector<StructFieldInit> inits;
+            for (py::handle item : fields) {
+                // Check FieldRef first: a scalar FieldRef also converts to Expr,
+                // but here it means "copy this field", not "use its value".
+                if (py::isinstance<FieldRef>(item)) {
+                    inits.emplace_back(item.cast<FieldRef>());
+                } else {
+                    inits.emplace_back(item.cast<Expr>());
+                }
+            }
+            return pack_struct(type, inits);
+        },
+        py::arg("type"), py::arg("fields"));
+
     m.def("max", [](const py::args &args) -> Expr {
         if (args.size() < 2) {
             throw py::value_error("max() must have at least 2 arguments");
@@ -189,6 +224,11 @@ void define_operators(py::module &m) {
     m.def("strict_float", &strict_float);
     m.def("scatter", static_cast<Expr (*)(const std::vector<Expr> &)>(&scatter));
     m.def("gather", static_cast<Expr (*)(const std::vector<Expr> &)>(&gather));
+    // Generator form: gather(extent, gen) builds a packet of `extent` elements
+    // by calling gen(k) for k in [0, extent), e.g. to fill a struct array field.
+    m.def("gather",
+          static_cast<Expr (*)(int, const std::function<Expr(Expr)> &)>(&gather),
+          py::arg("extent"), py::arg("gen"));
     m.def("extract_bits", static_cast<Expr (*)(Type, const Expr &, const Expr &)>(&extract_bits));
     m.def("concat_bits", &concat_bits);
     m.def("widen_right_add", &widen_right_add);
diff --git a/python_bindings/halide/src/halide_/PyType.cpp b/python_bindings/halide/src/halide_/PyType.cpp
index 31404826be05..de962e74110b 100644
--- a/python_bindings/halide/src/halide_/PyType.cpp
+++ b/python_bindings/halide/src/halide_/PyType.cpp
@@ -38,6 +38,14 @@ std::string halide_type_to_string(const Type &type) {
                 stream << "handle";
             }
             break;
+        case halide_type_struct:
+            // A struct's size is a byte count, not a bit width; report it and
+            // return early so the trailing bits() isn't appended.
+            stream << "struct" << std::to_string(type.bytes());
+            if (type.lanes() > 1) {
+                stream << "x" + std::to_string(type.lanes());
+            }
+            return stream.str();
         default:
             stream << "#unknown";
             break;
@@ -51,6 +59,48 @@ std::string halide_type_to_string(const Type &type) {
 }
 
 void define_type(py::module &m) {
+    py::class_<StructField>(m, "StructField")
+        .def(py::init([](const std::string &name, const Type &type, const py::object &array_extent) -> StructField {
+                 std::optional<int> extent;
+                 if (!array_extent.is_none()) {
+                     extent = array_extent.cast<int>();
+                 }
+                 return StructField{name, type, extent};
+             }),
+             py::arg("name"), py::arg("type"), py::arg("array_extent") = py::none())
+        .def(py::init([](const py::tuple &t) -> StructField {
+            if (t.size() < 2 || t.size() > 3) {
+                throw py::value_error("StructField requires (name, type) or (name, type, array_extent)");
+            }
+            std::optional<int> extent;
+            if (t.size() == 3 && !t[2].is_none()) {
+                extent = t[2].cast<int>();
+            }
+            return StructField{t[0].cast<std::string>(), t[1].cast<Type>(), extent};
+        }))
+        .def_readwrite("name", &StructField::name)
+        .def_readwrite("type", &StructField::type)
+        .def_property(
+            "array_extent",
+            [](const StructField &f) -> py::object {
+                return f.array_extent ? py::cast(*f.array_extent) : py::none();
+            },
+            [](StructField &f, const py::object &v) {
+                f.array_extent = v.is_none() ? std::nullopt : std::make_optional(v.cast<int>());
+            })
+        .def("__eq__", [](const StructField &a, const StructField &b) -> bool { return a == b; })
+        .def("__ne__", [](const StructField &a, const StructField &b) -> bool { return !(a == b); });
+
+    // Allow e.g. Type.Struct([("d", Float(16)), ("qs", UInt(8), 16)]),
+    // mirroring the tuple -> Range convenience conversion above.
+    py::implicitly_convertible<py::tuple, StructField>();
+
+    py::class_<StructTypeInfo>(m, "StructTypeInfo")
+        .def_readonly("fields", &StructTypeInfo::fields)
+        .def_readonly("offsets", &StructTypeInfo::offsets)
+        .def_readonly("total_bytes", &StructTypeInfo::total_bytes)
+        .def("find_field", &StructTypeInfo::find_field, py::arg("name"));
+
     py::class_<Type>(m, "Type")
         .def(py::init<>())
         .def(py::init<halide_type_code_t, int, int>(), py::arg("code"), py::arg("bits"), py::arg("lanes"))
@@ -73,6 +123,13 @@ void define_type(py::module &m) {
         .def("is_handle", &Type::is_handle)
         .def("same_handle_type", &Type::same_handle_type, py::arg("other"))
 
+        .def_static("Struct", &Type::Struct, py::arg("fields"))
+        .def("is_struct", &Type::is_struct)
+        .def("same_struct_type", &Type::same_struct_type, py::arg("other"))
+        .def_property_readonly("struct_type", [](const Type &t) -> py::object {
+            return t.struct_type() ? py::cast(StructTypeInfo(*t.struct_type())) : py::none();
+        })
+
         .def("__eq__", [](const Type &value, Type *value2) -> bool { return value2 && value == *value2; })
         .def("__ne__", [](const Type &value, Type *value2) -> bool { return !value2 || value != *value2; })
 
diff --git a/python_bindings/halide/test/correctness/CMakeLists.txt b/python_bindings/halide/test/correctness/CMakeLists.txt
index b9231af7e5c3..284a2952feb5 100644
--- a/python_bindings/halide/test/correctness/CMakeLists.txt
+++ b/python_bindings/halide/test/correctness/CMakeLists.txt
@@ -29,6 +29,7 @@ set(tests
     realize_warnings.py
     runtime_prefixes.py
     serialization.py
+    struct_type.py
     target.py
     tuple_select.py
     type.py
diff --git a/python_bindings/halide/test/correctness/struct_type.py b/python_bindings/halide/test/correctness/struct_type.py
new file mode 100644
index 000000000000..90b1a23f236c
--- /dev/null
+++ b/python_bindings/halide/test/correctness/struct_type.py
@@ -0,0 +1,163 @@
+import halide as hl
+
+
+# A block_q4_0-shaped struct, per doc/StructTypeDesign.md: one float32 scale
+# followed by 8 packed byte-wide "codes". This mirrors the C++
+# test/correctness/struct_type.cpp fixture (simplified to float32, since the
+# Python bindings have no float16 literal support).
+def make_block_type():
+    return hl.Type.Struct([("d", hl.Float(32)), ("qs", hl.UInt(8), 8)])
+
+
+def test_type_struct_basics():
+    block_a = make_block_type()
+    block_b = make_block_type()  # independently constructed, same layout
+    reordered = hl.Type.Struct([("qs", hl.UInt(8), 8), ("d", hl.Float(32))])
+    different_field = hl.Type.Struct([("d", hl.Float(32)), ("qs", hl.UInt(8), 4)])
+
+    assert block_a.is_struct()
+    assert block_b.is_struct()
+    assert block_a.bytes() == 4 + 8
+    assert block_a.code() == hl.TypeCode.Struct
+    assert not block_a.is_uint()
+    assert block_a.lanes() == 1
+
+    assert block_a == block_b
+    assert block_a.same_struct_type(block_b)
+    assert block_a != reordered
+    assert block_a != different_field
+    assert block_a != hl.UInt(8)
+
+    # A struct's type code carries its struct-ness, so with_bits/with_lanes
+    # (which don't change the code) always preserve it; only with_code to a
+    # different code drops it.
+    assert block_a.with_bits(8).is_struct()
+    assert block_a.with_bits(16).is_struct()
+    assert block_a.with_lanes(1).is_struct()
+    assert not block_a.with_code(hl.TypeCode.UInt).is_struct()
+
+    info = block_a.struct_type
+    assert info is not None
+    assert info.total_bytes == 12
+    assert [f.name for f in info.fields] == ["d", "qs"]
+    assert info.fields[0].type == hl.Float(32)
+    assert info.fields[1].type == hl.UInt(8)
+    assert info.fields[1].array_extent == 8
+    assert info.fields[0].array_extent is None
+    assert list(info.offsets) == [0, 4]
+    assert info.find_field("qs") == 1
+    assert info.find_field("nope") == -1
+
+    assert hl.Int(32).struct_type is None
+    assert not hl.Int(32).is_struct()
+
+
+# Struct types are treated as an ordinary opaque element type throughout the
+# compiler: a Func whose value is pack_struct(...), consumed via field() by
+# another Func, must produce identical results whether the producer is left
+# at its default (inlined) schedule or explicitly compute_root()'d (which
+# forces a real struct-typed buffer to be materialized).
+def _build_pipeline(schedule):
+    block_t = make_block_type()
+    blk = hl.Var("blk")
+    k = hl.Var("k")
+
+    producer = hl.Func("producer")
+    # Per-field: a scalar `d`, and the array field `qs` filled with a gather()
+    # generator (extent inferred from the struct field).
+    producer[blk] = hl.pack_struct(
+        block_t, [hl.f32(blk) + 0.5, hl.gather(8, lambda kk: hl.u8(blk * 3 + kk))]
+    )
+
+    schedule(producer)
+
+    delta = hl.Func("delta")
+    delta[blk] = hl.f32(hl.field(producer[blk], "d"))
+
+    codes = hl.Func("codes")
+    codes[k, blk] = hl.i32(hl.field(producer[blk], "qs")[k])
+
+    return delta, codes
+
+
+def _check_results(delta, codes, num_blocks):
+    delta_buf = delta.realize([num_blocks])
+    codes_buf = codes.realize([8, num_blocks])
+    for b in range(num_blocks):
+        assert abs(delta_buf[b] - (b + 0.5)) < 1e-6
+        for k in range(8):
+            assert codes_buf[k, b] == (b * 3 + k) % 256
+
+
+def test_struct_pack_and_field_inlined():
+    delta, codes = _build_pipeline(lambda f: None)
+    _check_results(delta, codes, num_blocks=6)
+
+
+def test_struct_pack_and_field_compute_root():
+    delta, codes = _build_pipeline(lambda f: f.compute_root())
+    _check_results(delta, codes, num_blocks=6)
+
+
+# FieldRef itself: scalar fields implicitly convert to Expr, array fields
+# support __getitem__/size()/__len__.
+def test_field_ref_array_indexing():
+    block_t = make_block_type()
+    blk = hl.Var("blk")
+
+    producer = hl.Func("producer")
+    # Array field filled by sweeping the `_` placeholder over its extent, with
+    # index arithmetic (blk + _).
+    producer[blk] = hl.pack_struct(block_t, [hl.f32(blk), hl.u8(blk + hl._)])
+
+    qs_field = hl.field(producer[blk], "qs")
+    assert qs_field.size() == 8
+    assert len(qs_field) == 8
+
+    out = hl.Func("out")
+    k = hl.Var("k")
+    out[k, blk] = hl.i32(qs_field[k])
+
+    out_buf = out.realize([8, 3])
+    for b in range(3):
+        for k_ in range(8):
+            assert out_buf[k_, b] == b + k_
+
+
+# pack_struct can copy a whole array field from another struct via field(),
+# and fill a scalar field from an explicit gather() over an existing list.
+def test_pack_struct_field_copy():
+    block_t = make_block_type()
+    blk = hl.Var("blk")
+    k = hl.Var("k")
+
+    src = hl.Func("src")
+    qs = [hl.u8(blk * 2 + kk) for kk in range(8)]
+    src[blk] = hl.pack_struct(block_t, [hl.f32(blk), hl.gather(qs)])
+    src.compute_root()
+
+    # Re-pack: replace `d`, keep `qs` by copying the whole array field.
+    s = src[blk]
+    repacked = hl.Func("repacked")
+    repacked[blk] = hl.pack_struct(block_t, [hl.f32(blk) + 100.0, hl.field(s, "qs")])
+
+    out_d = hl.Func("out_d")
+    out_q = hl.Func("out_q")
+    out_d[blk] = hl.f32(hl.field(repacked[blk], "d"))
+    out_q[k, blk] = hl.i32(hl.field(repacked[blk], "qs")[k])
+
+    num_blocks = 5
+    d_buf = out_d.realize([num_blocks])
+    q_buf = out_q.realize([8, num_blocks])
+    for b in range(num_blocks):
+        assert abs(d_buf[b] - (b + 100.0)) < 1e-6
+        for k_ in range(8):
+            assert q_buf[k_, b] == (b * 2 + k_) % 256
+
+
+if __name__ == "__main__":
+    test_type_struct_basics()
+    test_struct_pack_and_field_inlined()
+    test_struct_pack_and_field_compute_root()
+    test_field_ref_array_indexing()
+    test_pack_struct_field_copy()
diff --git a/src/Bounds.cpp b/src/Bounds.cpp
index 7aeff2f87bad..d2f82b8a8ab5 100644
--- a/src/Bounds.cpp
+++ b/src/Bounds.cpp
@@ -256,6 +256,8 @@ class Bounds : public IRVisitor {
         if ((t.is_uint() || t.is_int()) && t.bits() <= 16) {
             interval = Interval(t.min(), t.max());
         } else {
+            // Handle and struct types (and wider integers/floats) have no
+            // meaningful bounded numeric range.
             interval = Interval::everything();
         }
     }
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 02215600ef8e..511917c9f6b7 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -159,6 +159,7 @@ target_sources(
     Lower.h
     LowerParallelTasks.h
     LowerSMEStreamingTasks.h
+    LowerStructTypes.h
     LowerWarpShuffles.h
     Memoization.h
     Module.h
@@ -340,6 +341,7 @@ target_sources(
     Lower.cpp
     LowerParallelTasks.cpp
     LowerSMEStreamingTasks.cpp
+    LowerStructTypes.cpp
     LowerWarpShuffles.cpp
     Memoization.cpp
     Module.cpp
diff --git a/src/CodeGen_C.cpp b/src/CodeGen_C.cpp
index cd537062185f..3e201d82d63a 100644
--- a/src/CodeGen_C.cpp
+++ b/src/CodeGen_C.cpp
@@ -2383,6 +2383,11 @@ void CodeGen_C::visit(const Allocate *op) {
 
     string op_name = print_name(op->name);
     string op_type = print_type(op->type, AppendSpace);
+    string op_bytes = std::to_string(op->type.bytes());
+
+    // A struct's C storage element is a raw uint8_t (see type_to_c_type), so
+    // its per-element size must come from the true byte count, not sizeof.
+    string elem_size_expr = op->type.is_struct() ? op_bytes : ("sizeof(" + op_type + ")");
 
     // For sizes less than 8k, do a stack allocation
     bool on_stack = false;
@@ -2437,8 +2442,8 @@ void CodeGen_C::visit(const Allocate *op) {
             }
             stream << get_indent() << "if (("
                    << size_id << " > ((int64_t(1) << 31) - 1)) || (("
-                   << size_id << " * sizeof("
-                   << op_type << ")) > ((int64_t(1) << 31) - 1)))\n";
+                   << size_id << " * "
+                   << elem_size_expr << ") > ((int64_t(1) << 31) - 1)))\n";
             open_scope();
             stream << get_indent();
             // TODO: call halide_error_buffer_allocation_too_large() here instead
@@ -2468,8 +2473,10 @@ void CodeGen_C::visit(const Allocate *op) {
         stream << get_indent() << op_type;
 
         if (on_stack) {
+            // For a struct, op_type is "uint8_t", so size_id is really an element count.
+            string declared_size = op->type.is_struct() ? (size_id + " * " + op_bytes) : size_id;
             stream << op_name
-                   << "[" << size_id << "];\n";
+                   << "[" << declared_size << "];\n";
         } else {
             // Shouldn't ever currently be possible to have !on_stack && size_id.empty(),
             // but reality-check in case things change in the future.
@@ -2478,9 +2485,9 @@ void CodeGen_C::visit(const Allocate *op) {
                    << op_name
                    << " = ("
                    << op_type
-                   << " *)halide_malloc(_ucon, sizeof("
-                   << op_type
-                   << ")*" << size_id << ");\n";
+                   << " *)halide_malloc(_ucon, "
+                   << elem_size_expr
+                   << "*" << size_id << ");\n";
             heap_allocations.push(op->name);
         }
     }
diff --git a/src/CodeGen_D3D12Compute_Dev.cpp b/src/CodeGen_D3D12Compute_Dev.cpp
index f58b43452271..9c4c0f488a04 100644
--- a/src/CodeGen_D3D12Compute_Dev.cpp
+++ b/src/CodeGen_D3D12Compute_Dev.cpp
@@ -1431,6 +1431,12 @@ void CodeGen_D3D12Compute_Dev::CodeGen_D3D12Compute_C::visit(const Allocate *op)
             << "Only fixed-size allocations are supported on the gpu. "
             << "Try storing into shared memory instead.";
 
+        // A struct is stored as a raw byte array (print_storage_type emits a byte
+        // element type), so size is an element count and must be scaled to bytes.
+        if (op->type.is_struct()) {
+            size *= op->type.bytes();
+        }
+
         stream << get_indent() << print_storage_type(op->type) << " "
                << print_name(op->name) << "[" << size << "];\n";
         stream << get_indent();
diff --git a/src/CodeGen_LLVM.cpp b/src/CodeGen_LLVM.cpp
index 67b91e9b1775..c2a3da1ab7fd 100644
--- a/src/CodeGen_LLVM.cpp
+++ b/src/CodeGen_LLVM.cpp
@@ -1078,7 +1078,7 @@ void CodeGen_LLVM::compile_buffer(const Buffer<> &buf) {
     Constant *type_fields[] = {
         ConstantInt::get(i8_t, buf.type().code()),
         ConstantInt::get(i8_t, buf.type().bits()),
-        ConstantInt::get(i16_t, 0)};
+        ConstantInt::get(i16_t, buf.type().to_abi().reserved)};  // struct byte size, else 0
 
     Constant *shape = nullptr;
     if (buf.dimensions()) {
@@ -1273,7 +1273,7 @@ llvm::Function *CodeGen_LLVM::embed_metadata_getter(const std::string &metadata_
         Constant *type_fields[] = {
             ConstantInt::get(i8_t, args[arg].type.code()),
             ConstantInt::get(i8_t, args[arg].type.bits()),
-            ConstantInt::get(i16_t, 0)};  // reserved (formerly lanes); must be 0
+            ConstantInt::get(i16_t, args[arg].type.to_abi().reserved)};  // struct byte size, else 0
         Constant *type = ConstantStruct::get(type_t_type, type_fields);
 
         auto argument_estimates = args[arg].argument_estimates;
@@ -6031,6 +6031,9 @@ int CodeGen_LLVM::get_vector_num_elements(const llvm::Value *v) {
 llvm::Type *CodeGen_LLVM::llvm_type_of(LLVMContext *c, Halide::Type t,
                                        int effective_vscale) const {
     if (t.lanes() == 1) {
+        if (t.is_struct()) {
+            return llvm::ArrayType::get(llvm::Type::getInt8Ty(*c), t.bytes());
+        }
         if (t.is_float() && !t.is_bfloat()) {
             switch (t.bits()) {
             case 16:
diff --git a/src/CodeGen_Metal_Dev.cpp b/src/CodeGen_Metal_Dev.cpp
index 08145925aa10..b25b342b2280 100644
--- a/src/CodeGen_Metal_Dev.cpp
+++ b/src/CodeGen_Metal_Dev.cpp
@@ -557,8 +557,15 @@ void CodeGen_Metal_Dev::CodeGen_Metal_C::visit(const Allocate *op) {
             << "Only fixed-size allocations are supported on the gpu. "
             << "Try storing into shared memory instead.";
 
+        // A struct is stored as a raw byte array (print_storage_type emits a byte
+        // element type), so size is an element count and must be scaled to bytes.
+        if (op->type.is_struct()) {
+            size *= op->type.bytes();
+        }
+
         stream << get_indent() << print_storage_type(op->type) << " "
                << print_name(op->name) << "[" << size << "];\n";
+
         stream << get_indent() << "#define " << get_memory_space(op->name) << " thread\n";
 
         Allocation alloc;
diff --git a/src/CodeGen_OpenCL_Dev.cpp b/src/CodeGen_OpenCL_Dev.cpp
index 88e858355f1e..41b40bdfa238 100644
--- a/src/CodeGen_OpenCL_Dev.cpp
+++ b/src/CodeGen_OpenCL_Dev.cpp
@@ -798,8 +798,15 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Allocate *op) {
             << "Only fixed-size allocations are supported on the gpu. "
             << "Try storing into shared memory instead.";
 
+        // A struct is stored as a raw byte array (print_storage_type emits a byte
+        // element type), so size is an element count and must be scaled to bytes.
+        if (op->type.is_struct()) {
+            size *= op->type.bytes();
+        }
+
         stream << get_indent() << print_type(op->type) << " "
                << print_name(op->name) << "[" << size << "];\n";
+
         stream << get_indent() << "#define " << get_memory_space(op->name) << " __private\n";
 
         Allocation alloc;
diff --git a/src/CodeGen_Vulkan_Dev.cpp b/src/CodeGen_Vulkan_Dev.cpp
index b7d8dd844667..c23a8b062196 100644
--- a/src/CodeGen_Vulkan_Dev.cpp
+++ b/src/CodeGen_Vulkan_Dev.cpp
@@ -1888,6 +1888,11 @@ void CodeGen_Vulkan_Dev::SPIRV_Emitter::visit(const Allocate *op) {
         // static fixed size allocation
         if (op->extents.size() == 1 && is_const(op->extents[0])) {
             array_size = op->constant_allocation_size();
+            // A struct is stored as a raw byte array (print_storage_type emits a byte
+            // element type), so size is an element count and must be scaled to bytes.
+            if (op->type.is_struct()) {
+                array_size *= op->type.bytes();
+            }
             array_type_id = builder.declare_type(op->type, array_size);
             builder.add_symbol(variable_name + "_array_type", array_type_id, builder.current_module().id());
             debug(2) << "Vulkan: Allocate (fixed-size) " << op->name << " type=" << op->type << " array_size=" << array_size << " in shared memory on device in global scope\n";
@@ -1895,6 +1900,8 @@ void CodeGen_Vulkan_Dev::SPIRV_Emitter::visit(const Allocate *op) {
         } else {
             // dynamic allocation with unknown size at compile time ...
 
+            // TODO: what to do for struct types here? It doesn't seem to depend on op->type.
+
             // declare the array size as a specialization constant (which will get overridden at runtime)
             Type array_size_type = UInt(32);
             array_size = std::max(workgroup_size[0], uint32_t(1));  // use one item per workgroup as an initial guess
@@ -1934,6 +1941,12 @@ void CodeGen_Vulkan_Dev::SPIRV_Emitter::visit(const Allocate *op) {
             << "Allocation " << op->name << " has a dynamic size. "
             << "Only fixed-size local allocations are supported with Vulkan.";
 
+        // A struct is stored as a raw byte array (print_storage_type emits a byte
+        // element type), so size is an element count and must be scaled to bytes.
+        if (op->type.is_struct()) {
+            array_size *= op->type.bytes();
+        }
+
         debug(2) << "Vulkan: Allocate " << op->name << " type=" << op->type << " size=" << array_size << " on device in function scope\n";
 
         array_type_id = builder.declare_type(op->type, array_size);
diff --git a/src/CodeGen_WebGPU_Dev.cpp b/src/CodeGen_WebGPU_Dev.cpp
index 23db34fa02ad..089cf36d7e8a 100644
--- a/src/CodeGen_WebGPU_Dev.cpp
+++ b/src/CodeGen_WebGPU_Dev.cpp
@@ -428,6 +428,12 @@ void CodeGen_WebGPU_Dev::CodeGen_WGSL::visit(const Allocate *op) {
             << "Only fixed-size allocations are supported on the gpu. "
             << "Try storing into shared memory instead.";
 
+        // A struct is stored as a raw byte array (print_storage_type emits a byte
+        // element type), so size is an element count and must be scaled to bytes.
+        if (op->type.is_struct()) {
+            size *= op->type.bytes();
+        }
+
         stream << get_indent() << "var " << print_name(op->name)
                << " : array<" << print_type(op->type) << ", " << size << ">;\n";
 
diff --git a/src/FuseGPUThreadLoops.cpp b/src/FuseGPUThreadLoops.cpp
index a444273c9a46..fa3a48b5ea1a 100644
--- a/src/FuseGPUThreadLoops.cpp
+++ b/src/FuseGPUThreadLoops.cpp
@@ -997,13 +997,15 @@ class ExtractSharedAndHeapAllocations : public IRMutator {
                 internal_assert(ratio != 0)
                     << "alloc_type should have been at most as wide as the widest type in group\n";
                 // Sizes here are counted in units of one type or another, and
-                // are converted between them by dividing, so the types have to
-                // be whole multiples of each other.
-                internal_assert(is_power_of_two(alloc_type.bytes()) &&
-                                is_power_of_two(alloc.widest_type.bytes()))
-                    << "Allocation types must be a power of two bytes wide, but these "
-                    << "are " << alloc_type.bytes() << " and "
-                    << alloc.widest_type.bytes() << "\n";
+                // are converted between them by dividing, so the widest type
+                // has to be a whole multiple of alloc_type -- not necessarily
+                // a power of two, e.g. when alloc_type is UInt(8) (always the
+                // case when may_merge_allocs_of_different_type is set) any
+                // byte count, including a struct's, divides evenly.
+                internal_assert(alloc.widest_type.bytes() % alloc_type.bytes() == 0)
+                    << "Allocation type " << alloc.widest_type << " (" << alloc.widest_type.bytes()
+                    << " bytes) is not a whole multiple of the cluster's allocation type "
+                    << alloc_type << " (" << alloc_type.bytes() << " bytes)\n";
                 total_size += align_up(alloc.max_size * ratio,
                                        async_copy_alignment / alloc_type.bytes());
             }
diff --git a/src/IR.cpp b/src/IR.cpp
index a5ff626ed7e6..d1d3f8ccd567 100644
--- a/src/IR.cpp
+++ b/src/IR.cpp
@@ -912,6 +912,8 @@ constexpr const char *intrinsic_op_names[] = {
     "strict_mul",
     "strict_sub",
     "stringify",
+    "struct_field_read",
+    "struct_pack",
     "target_arch_is",
     "target_bits",
     "target_has_feature",
diff --git a/src/IR.h b/src/IR.h
index 113b1876a830..2204515f381c 100644
--- a/src/IR.h
+++ b/src/IR.h
@@ -890,6 +890,15 @@ struct Call : public ExprNode<Call> {
         strict_sub,
         // Convert a list of Exprs to a string
         stringify,
+        // Read one element of one field out of a struct-typed value.
+        // args = {struct value, field index as an IntImm, element index}.
+        // The field's name is resolved to its index eagerly; the element index
+        // is 0 for a scalar field.
+        struct_field_read,
+        // Construct a value of a struct type from its field elements, in
+        // declaration order, flattened element-by-element (an array field
+        // contributes array_extent args).
+        struct_pack,
         // Query properties of the compiled-for target (resolved at compile-time)
         target_arch_is,
         target_bits,
diff --git a/src/IROperator.cpp b/src/IROperator.cpp
index 710883026201..abf2398bf560 100644
--- a/src/IROperator.cpp
+++ b/src/IROperator.cpp
@@ -18,6 +18,7 @@
 #include "IRVisitor.h"
 #include "Interval.h"
 #include "StrictifyFloat.h"
+#include "Substitute.h"
 #include "Util.h"
 #include "Var.h"
 
@@ -1041,7 +1042,7 @@ void split_into_ands(const Expr &cond, std::vector<Expr> &result) {
 }
 
 Expr BufferBuilder::build() const {
-    std::vector<Expr> args(10);
+    std::vector<Expr> args(11);
     if (buffer_memory.defined()) {
         args[0] = buffer_memory;
     } else {
@@ -1079,7 +1080,11 @@ Expr BufferBuilder::build() const {
 
     args[5] = (int)type.code();
     args[6] = type.bits();
-    args[7] = dimensions;
+    // A struct element type carries its packed byte size in the ABI's reserved
+    // field; thread it through so the runtime buffer's type is faithful (and
+    // its element stride/allocation size correct). Zero for ordinary types.
+    args[7] = (int)type.to_abi().reserved;
+    args[8] = dimensions;
 
     std::vector<Expr> shape;
     for (size_t i = 0; i < (size_t)dimensions; i++) {
@@ -1107,11 +1112,11 @@ Expr BufferBuilder::build() const {
     }
     Expr shape_arg = Call::make(type_of<halide_dimension_t *>(), Call::make_struct, shape, Call::Intrinsic);
     if (shape_memory.defined()) {
-        args[8] = shape_arg;
+        args[9] = shape_arg;
     } else if (dimensions == 0) {
-        args[8] = make_zero(type_of<halide_dimension_t *>());
+        args[9] = make_zero(type_of<halide_dimension_t *>());
     } else {
-        args[8] = shape_var;
+        args[9] = shape_var;
     }
 
     Expr flags = make_zero(UInt(64));
@@ -1125,7 +1130,7 @@ Expr BufferBuilder::build() const {
                                make_const(UInt(64), halide_buffer_flag_device_dirty),
                                make_zero(UInt(64)));
     }
-    args[9] = flags;
+    args[10] = flags;
 
     Expr e = Call::make(type_of<struct halide_buffer_t *>(), Call::buffer_init, args, Call::Extern);
 
@@ -2956,6 +2961,16 @@ Expr gather(const std::vector<Expr> &args) {
     return make_scatter_gather(args);
 }
 
+Expr gather(int extent, const std::function<Expr(Expr)> &gen) {
+    user_assert(extent >= 1) << "gather() extent must be at least 1, got " << extent << ".\n";
+    std::vector<Expr> elems;
+    elems.reserve(extent);
+    for (int k = 0; k < extent; k++) {
+        elems.push_back(gen(make_const(Int(32), k)));
+    }
+    return make_scatter_gather(elems);
+}
+
 Expr extract_bits(Type t, const Expr &e, const Expr &lsb) {
     return Call::make(t, Call::extract_bits, {e, lsb}, Call::Intrinsic);
 }
@@ -2970,6 +2985,189 @@ Expr concat_bits(const std::vector<Expr> &e) {
     return Call::make(t.with_bits(t.bits() * (int)e.size()), Call::concat_bits, e, Call::Intrinsic);
 }
 
+Expr FieldRef::read(const Expr &elem_index) const {
+    return Call::make(elem_type, Call::struct_field_read,
+                      {struct_value, make_const(Int(32), field_index), elem_index}, Call::PureIntrinsic);
+}
+
+FieldRef::FieldRef(Expr struct_value, int field_index, Type elem_type, std::optional<int> array_extent)
+    : struct_value(std::move(struct_value)), field_index(field_index), elem_type(elem_type), array_extent(array_extent) {
+}
+
+FieldRef::operator Expr() const {
+    user_assert(!array_extent.has_value())
+        << "This struct field is an array field; use operator[] to access an element of it, "
+        << "not an implicit conversion to Expr.\n";
+    return read(make_const(Int(32), 0));
+}
+
+Expr FieldRef::operator[](const Expr &i) const {
+    user_assert(array_extent.has_value())
+        << "This struct field is a scalar field; it can't be indexed with operator[].\n";
+    Expr idx = cast<int32_t>(i);
+    if (auto ci = as_const_int(idx)) {
+        user_assert(*ci >= 0 && *ci < *array_extent)
+            << "Struct array field index " << *ci << " is out of range [0, " << *array_extent << ").\n";
+    }
+    return read(idx);
+}
+
+namespace {
+FieldRef make_field_ref(const Expr &struct_value, int index) {
+    user_assert(struct_value.defined() && struct_value.type().is_struct())
+        << "field() requires an Expr of a struct type.\n";
+    const StructTypeInfo *info = struct_value.type().struct_type();
+    user_assert(index >= 0 && index < (int)info->fields.size())
+        << "Struct field index " << index << " is out of range; this struct has "
+        << info->fields.size() << " fields.\n";
+    const StructField &f = info->fields[index];
+    return FieldRef(struct_value, index, f.type, f.array_extent);
+}
+}  // namespace
+
+FieldRef field(const Expr &struct_value, const std::string &name) {
+    user_assert(struct_value.defined() && struct_value.type().is_struct())
+        << "field() requires an Expr of a struct type (see Type::Struct).\n";
+    int index = struct_value.type().struct_type()->find_field(name);
+    user_assert(index >= 0) << "Struct type has no field named \"" << name << "\".\n";
+    return make_field_ref(struct_value, index);
+}
+
+FieldRef field(const Expr &struct_value, int index) {
+    return make_field_ref(struct_value, index);
+}
+
+Expr pack_struct(const Type &t, const std::vector<Expr> &field_values) {
+    user_assert(t.is_struct()) << "pack_struct() requires a struct type.\n";
+    const StructTypeInfo &info = *t.struct_type();
+
+    size_t expected = 0;
+    for (const auto &f : info.fields) {
+        expected += (size_t)f.array_extent.value_or(1);
+    }
+    user_assert(field_values.size() == expected)
+        << "pack_struct() for a " << expected << "-element struct (counting array fields "
+        << "element-by-element) was given " << field_values.size() << " values.\n";
+
+    size_t idx = 0;
+    for (const auto &f : info.fields) {
+        int n = f.array_extent.value_or(1);
+        for (int i = 0; i < n; i++) {
+            user_assert(field_values[idx].defined() && field_values[idx].type() == f.type)
+                << "pack_struct(): value " << idx << " has type " << field_values[idx].type()
+                << " but struct field \"" << f.name << "\" has type " << f.type << ".\n";
+            idx++;
+        }
+    }
+
+    return Call::make(t, Call::struct_pack, field_values, Call::PureIntrinsic);
+}
+
+namespace {
+// Collect the distinct implicit-var names (`_`, `_0`, ...) appearing in an
+// Expr, in first-encountered order. These are the placeholder(s) a pack_struct
+// array-field initializer sweeps over its extent.
+class FindImplicitVars : public IRGraphVisitor {
+    using IRGraphVisitor::visit;
+    std::set<std::string> seen;
+
+    void visit(const Variable *op) override {
+        if (Var::is_implicit(op->name) && seen.insert(op->name).second) {
+            names.push_back(op->name);
+        }
+    }
+
+public:
+    std::vector<std::string> names;
+};
+
+std::vector<std::string> find_implicit_vars(const Expr &e) {
+    FindImplicitVars finder;
+    e.accept(&finder);
+    return finder.names;
+}
+}  // namespace
+
+Expr pack_struct(const Type &t, const std::vector<StructFieldInit> &fields) {
+    user_assert(t.is_struct()) << "pack_struct() requires a struct type.\n";
+    const StructTypeInfo &info = *t.struct_type();
+
+    user_assert(fields.size() == info.fields.size())
+        << "pack_struct() for a " << info.fields.size() << "-field struct type was given "
+        << fields.size() << " field initializers. Supply exactly one initializer per field "
+        << "(an array field's elements go in a single gather() or a swept `_` expression).\n";
+
+    std::vector<Expr> flat;
+    for (size_t fi = 0; fi < info.fields.size(); fi++) {
+        const StructField &f = info.fields[fi];
+        const int extent = f.array_extent.value_or(1);
+        const bool is_array = f.array_extent.has_value();
+        const StructFieldInit &init = fields[fi];
+
+        if (init.source) {
+            // Copy a whole same-typed field out of another struct.
+            const FieldRef &r = *init.source;
+            user_assert(r.element_type() == f.type)
+                << "pack_struct(): field \"" << f.name << "\" of type " << f.type
+                << " is initialized from a field of type " << r.element_type() << ".\n";
+            user_assert(r.size() == extent)
+                << "pack_struct(): field \"" << f.name << "\" has " << extent
+                << " element(s), but is initialized from a field with " << r.size() << ".\n";
+            if (is_array) {
+                for (int k = 0; k < extent; k++) {
+                    flat.push_back(r[k]);
+                }
+            } else {
+                flat.push_back(Expr(r));
+            }
+            continue;
+        }
+
+        const Expr &e = *init.value;
+        if (const Call *g = e.as<Call>(); g != nullptr && g->is_intrinsic(Call::scatter_gather)) {
+            // An explicit gather() packet fills an array field's elements.
+            user_assert(is_array)
+                << "pack_struct(): scalar field \"" << f.name << "\" was given a gather() packet; "
+                << "a gather fills an array field.\n";
+            user_assert((int)g->args.size() == extent)
+                << "pack_struct(): array field \"" << f.name << "\" has extent " << extent
+                << ", but its gather() supplies " << g->args.size() << " element(s).\n";
+            flat.insert(flat.end(), g->args.begin(), g->args.end());
+            continue;
+        }
+
+        std::vector<std::string> implicit = find_implicit_vars(e);
+        user_assert(implicit.size() <= 1)
+            << "pack_struct(): the initializer for field \"" << f.name << "\" contains "
+            << implicit.size() << " placeholders (" << (implicit.empty() ? "" : implicit[0])
+            << " ...); an array field may sweep at most one.\n";
+
+        if (!implicit.empty()) {
+            // A single `_` placeholder swept over the field's extent.
+            user_assert(is_array)
+                << "pack_struct(): scalar field \"" << f.name << "\" was given a swept `_` "
+                << "expression; a placeholder sweep fills an array field.\n";
+            for (int k = 0; k < extent; k++) {
+                flat.push_back(substitute(implicit[0], make_const(Int(32), k), e));
+            }
+        } else {
+            // A plain value: only a scalar field.
+            user_assert(!is_array)
+                << "pack_struct(): array field \"" << f.name << "\" needs its " << extent
+                << " elements supplied via gather(), a swept `_` expression, or a field() copy, "
+                << "not a single value.\n";
+            flat.push_back(e);
+        }
+    }
+
+    // The flattened form does the final per-element type check and builds the node.
+    return pack_struct(t, flat);
+}
+
+Expr pack_struct(const Type &t, std::initializer_list<StructFieldInit> fields) {
+    return pack_struct(t, std::vector<StructFieldInit>(fields));
+}
+
 Expr target_arch_is(Target::Arch arch) {
     return Call::make(Bool(), Call::target_arch_is, {Expr((int)arch)}, Call::PureIntrinsic);
 }
diff --git a/src/IROperator.h b/src/IROperator.h
index 6489a580e106..4f65d39b45ba 100644
--- a/src/IROperator.h
+++ b/src/IROperator.h
@@ -8,8 +8,10 @@
  */
 
 #include <cmath>
+#include <functional>
 #include <map>
 #include <optional>
+#include <type_traits>
 #include <utility>
 
 #include "Bounds.h"
@@ -1610,22 +1612,34 @@ f(select(p, scatter(3, 5, 5), scatter(1, 2, 3))) = f(select(p, gather(5, 3, 3),
 *
 * Note that in the p == true case, we redundantly load from 3 and write
 * to 5 twice.
+*
+* A gather is also the way to supply the elements of a struct's array field to
+* pack_struct(): the packet's values become that field's elements, in order.
+* See pack_struct and \ref gather(int, const std::function<Expr(Expr)> &).
 */
 //@{
 Expr scatter(const std::vector<Expr> &args);
 Expr gather(const std::vector<Expr> &args);
 
-template<typename... Args>
+template<typename... Args,
+         typename = std::enable_if_t<(std::is_convertible_v<Args, Expr> && ...)>>
 Expr scatter(const Expr &e, Args &&...args) {
     return scatter({e, std::forward<Args>(args)...});
 }
 
-template<typename... Args>
+template<typename... Args,
+         typename = std::enable_if_t<(std::is_convertible_v<Args, Expr> && ...)>>
 Expr gather(const Expr &e, Args &&...args) {
     return gather({e, std::forward<Args>(args)...});
 }
 // @}
 
+/** Build a gather packet of `extent` elements by evaluating `gen(k)` for each
+ * `k` in `[0, extent)` (passed as an int32 constant). This is the general way
+ * to fill a struct array field with a computed value per element when the fill
+ * can't be written as a single placeholder sweep; see pack_struct. */
+Expr gather(int extent, const std::function<Expr(Expr)> &gen);
+
 /** Extract a contiguous subsequence of the bits of 'e', starting at the bit
  * index given by 'lsb', where zero is the least-significant bit, returning a
  * value of type 't'. Any out-of-range bits requested are filled with zeros.
@@ -1674,6 +1688,93 @@ f32.vectorize(x, 8);
  */
 Expr concat_bits(const std::vector<Expr> &e);
 
+/** A reference to one field of a struct-typed Expr, returned by field().
+ * Implicitly converts to Expr for a scalar field; supports operator[] for
+ * an array field. Using the wrong one of these for how the field was actually
+ * declared is a user_error.
+ */
+class FieldRef {
+    Expr struct_value;
+    int field_index;
+    Type elem_type;
+    std::optional<int> array_extent;
+
+    Expr read(const Expr &elem_index) const;
+
+public:
+    FieldRef(Expr struct_value, int field_index, Type elem_type, std::optional<int> array_extent);
+
+    /** Valid only for a scalar field. */
+    operator Expr() const;
+
+    /** Valid only for an array field. i may be a runtime Expr. */
+    Expr operator[](const Expr &i) const;
+
+    /** The number of elements, for an array field. 1 for a scalar field. */
+    int size() const {
+        return array_extent.value_or(1);
+    }
+
+    /** Whether this refers to an array field (vs. a scalar field). */
+    bool is_array() const {
+        return array_extent.has_value();
+    }
+
+    /** The element type of the field. */
+    Type element_type() const {
+        return elem_type;
+    }
+};
+
+/** Extract field `name` (or `index`) from a struct-typed Expr. The field's
+ * offset and type are resolved eagerly.
+ */
+// @{
+FieldRef field(const Expr &struct_value, const std::string &name);
+FieldRef field(const Expr &struct_value, int index);
+// @}
+
+/** Construct a value of struct type `t` (see Type::Struct) from one Expr
+ * per field, in declaration order. All fields must be supplied; array
+ * fields take `array_extent` Exprs, flattened into the same list. This is the
+ * low-level form; prefer the per-field overload below. */
+Expr pack_struct(const Type &t, const std::vector<Expr> &field_values);
+
+/** One field's worth of initializer for the per-field pack_struct() overload.
+ * Implicitly constructible from anything convertible to Expr -- a scalar
+ * field's value; a gather() packet whose elements fill an array field; or a
+ * single expression containing one implicit-var placeholder `_`, swept over the
+ * array field's extent (index arithmetic on `_` is allowed) -- and from a
+ * FieldRef, to copy a whole same-typed field out of another struct. */
+class StructFieldInit {
+public:
+    template<typename T,
+             typename = std::enable_if_t<std::is_convertible_v<T, Expr> &&
+                                         !std::is_same_v<std::decay_t<T>, FieldRef>>>
+    StructFieldInit(T &&e)
+        : value(Expr(std::forward<T>(e))) {
+    }
+    StructFieldInit(const FieldRef &f)
+        : source(f) {
+    }
+
+private:
+    friend Expr pack_struct(const Type &t, const std::vector<StructFieldInit> &fields);
+    // Exactly one of these is engaged.
+    std::optional<Expr> value;
+    std::optional<FieldRef> source;
+};
+
+/** Construct a value of struct type `t` with one initializer per field, in
+ * declaration order (contrast the flattened form above). A scalar field takes a
+ * scalar Expr; an array field takes a gather() packet, a single expression with
+ * one `_` placeholder swept over the field extent, or a FieldRef to a
+ * same-typed array field to copy element-by-element. */
+// @{
+Expr pack_struct(const Type &t, const std::vector<StructFieldInit> &fields);
+Expr pack_struct(const Type &t, std::initializer_list<StructFieldInit> fields);
+// @}
+
 /** Below is a collection of intrinsics for fixed-point programming. Most of
  * them can be expressed via other means, but this is more natural for some, as
  * it avoids ghost widened intermediates that don't (or shouldn't) actually show
diff --git a/src/IRPrinter.cpp b/src/IRPrinter.cpp
index b4ca72d54e3a..5e9de62b5a8a 100644
--- a/src/IRPrinter.cpp
+++ b/src/IRPrinter.cpp
@@ -46,8 +46,22 @@ ostream &operator<<(ostream &out, const Type &type) {
     case Type::BFloat:
         out << "bfloat";
         break;
+    case Type::StructKind:
+        out << "struct{";
+        if (const StructTypeInfo *info = type.struct_type()) {
+            const char *sep = "";
+            for (const auto &f : info->fields) {
+                out << sep << f.name << ": " << f.type;
+                if (f.array_extent) {
+                    out << "[" << *f.array_extent << "]";
+                }
+                sep = ", ";
+            }
+        }
+        out << "}";
+        break;
     }
-    if (!type.is_handle()) {
+    if (!type.is_handle() && !type.is_struct()) {
         out << type.bits();
     }
     if (type.lanes() > 1) {
diff --git a/src/ImageParam.cpp b/src/ImageParam.cpp
index 615ae1e6013b..a61fff5bfeb3 100644
--- a/src/ImageParam.cpp
+++ b/src/ImageParam.cpp
@@ -41,7 +41,7 @@ Func ImageParam::create_func() const {
 
 void ImageParam::set(const Buffer<> &b) {
     if (b.defined()) {
-        user_assert(b.type() == type())
+        user_assert(type().is_compatible_for_buffer_bind(b.type()))
             << "Can't bind ImageParam " << name()
             << " of type " << type()
             << " to Buffer " << b.name()
diff --git a/src/Lower.cpp b/src/Lower.cpp
index 21acfff8df2d..d6d7da79e465 100644
--- a/src/Lower.cpp
+++ b/src/Lower.cpp
@@ -46,6 +46,7 @@
 #include "LoopCarry.h"
 #include "LowerParallelTasks.h"
 #include "LowerSMEStreamingTasks.h"
+#include "LowerStructTypes.h"
 #include "LowerWarpShuffles.h"
 #include "Memoization.h"
 #include "OffloadGPULoops.h"
@@ -305,6 +306,10 @@ void lower_impl(const vector<Function> &output_funcs,
     s = storage_flattening(s, outputs, env, t);
     log("Lowering after storage flattening:", s);
 
+    debug(1) << "Lowering struct type field access...\n";
+    s = lower_struct_types(s);
+    log("Lowering after lowering struct type field access:", s);
+
     debug(1) << "Adding atomic mutex allocation...\n";
     s = add_atomic_mutex(s, outputs);
     log("Lowering after adding atomic mutex allocation:", s);
diff --git a/src/LowerStructTypes.cpp b/src/LowerStructTypes.cpp
new file mode 100644
index 000000000000..f3d6c1de2cfa
--- /dev/null
+++ b/src/LowerStructTypes.cpp
@@ -0,0 +1,330 @@
+#include "LowerStructTypes.h"
+#include "IRMutator.h"
+#include "IROperator.h"
+#include "IRVisitor.h"
+#include "Scope.h"
+
+namespace Halide {
+namespace Internal {
+
+using std::vector;
+
+namespace {
+
+// Number of pack_struct() args occupied by all fields before field_index.
+int flat_start_of_field(const StructTypeInfo &info, int field_index) {
+    int start = 0;
+    for (int i = 0; i < field_index; i++) {
+        start += info.fields[i].array_extent.value_or(1);
+    }
+    return start;
+}
+
+class LowerStructTypesMutator : public IRMutator {
+    using IRMutator::visit;
+
+    // Struct-typed Let bindings are tracked here because project_field() needs
+    // direct syntactic access to the underlying struct_pack()/Select/Load. Because
+    // struct-typed values are never materialized (only individual fields are),
+    // inlining it back in at each field() use site costs nothing at runtime.
+    Scope<Expr> struct_lets;
+
+    Expr resolve_struct_value(Expr e) const {
+        while (const Variable *v = e.as<Variable>()) {
+            internal_assert(v->type.is_struct());
+            user_assert(struct_lets.contains(v->name))
+                << "Struct-typed variable \"" << v->name << "\" is unbound; a struct-typed Expr "
+                << "may only be the value of a struct-typed Func/Store, or a direct field() "
+                << "argument, in v1.\n";
+            e = struct_lets.get(v->name);
+        }
+        return e;
+    }
+
+    static Expr fold_field_of_pack(const Call *pack, const StructTypeInfo &info, int field_index, const Expr &elem_index) {
+        const StructField &f = info.fields[field_index];
+        int flat_start = flat_start_of_field(info, field_index);
+        int extent = f.array_extent.value_or(1);
+
+        if (auto ci = as_const_int(elem_index)) {
+            int flat = flat_start + (int)*ci;
+            internal_assert(flat >= 0 && flat < (int)pack->args.size());
+            return pack->args[flat];
+        }
+
+        internal_assert(extent >= 1);
+        Expr result = pack->args[flat_start + extent - 1];
+        for (int i = extent - 2; i >= 0; i--) {
+            result = select(elem_index == i, pack->args[flat_start + i], result);
+        }
+        return result;
+    }
+
+    // Rewrite `field(e, field_index)[elem_index]`, where e is a real (already
+    // flattened) struct-typed Load into byte-addressed Loads plus a Reinterpret.
+    Expr lower_field_read_from_load(const Load *load, const Type &field_type, int field_base_offset, const Expr &elem_index_in) {
+        // A field whose own type is itself Type::Struct (a sub-struct field)
+        // can't be read this way: the byte-combine-and-Reinterpret trick
+        // below fundamentally assumes field_type's *bits* describe its
+        // *bytes* (true for every ordinary scalar type), but a struct's bits()
+        // is not its (packed, possibly not power-of-two) byte size, so
+        // Reinterpret would be asked to reinterpret between mismatched bit
+        // widths. Nested struct fields read from a literal pack_struct() go
+        // through fold_field_of_pack instead and don't hit this at all; only
+        // reading a nested field from a (non-inlined) struct-typed buffer is
+        // unsupported.
+        user_assert(!field_type.is_struct())
+            << "Reading a nested struct field (a field whose own type is Type::Struct) directly "
+            << "from a struct-typed buffer/Func that isn't fully inlined is not supported. "
+            << "Field \"" << field_type << "\" at byte offset " << field_base_offset
+            << " within " << load->name << " was requested this way.\n";
+
+        Expr elem_index = mutate(elem_index_in);
+        Expr flat_index = mutate(load->index);
+        Type index_type = flat_index.type();
+        int elem_bytes = field_type.bytes();
+
+        Expr byte_offset = flat_index * make_const(index_type, load->type.bytes()) +
+                           make_const(index_type, field_base_offset) +
+                           cast(index_type, elem_index) * make_const(index_type, elem_bytes);
+
+        auto byte_load = [&](int b) {
+            Expr idx = b == 0 ? byte_offset : byte_offset + make_const(index_type, b);
+            return Load::make(UInt(8), load->name, idx, load->image, load->param,
+                              const_true(), ModulusRemainder(), load->is_streaming);
+        };
+
+        if (elem_bytes == 1) {
+            Expr bl = byte_load(0);
+            return field_type == UInt(8) ? bl : Reinterpret::make(field_type, bl);
+        }
+
+        vector<Expr> bytes;
+        bytes.reserve(elem_bytes);
+        for (int b = 0; b < elem_bytes; b++) {
+            bytes.push_back(byte_load(b));
+        }
+
+        if (field_type.is_float()) {
+            // A float field (e.g. a packed fp16 delta) is built as a genuine
+            // shift/or chain of the bytes rather than through concat_bits's
+            // vector-shuffle-then-reinterpret lowering. Some AArch64 backends
+            // (LLVM < 23) cannot legalize a bitcast straight from a vector to
+            // a scalar half, and unlike the integer case below, there's no way
+            // to route around that once the value is vector-shaped: any chain
+            // of pure bitcasts collapses back to the direct (illegal) one
+            // during optimization. Building the packed bits with real ALU ops
+            // keeps the final Reinterpret's input a plain scalar integer, so
+            // the last step is an always-legal scalar-to-scalar bitcast; LLVM's
+            // load-combiner still typically recovers a single wide load.
+            Type uint_t = UInt(8 * elem_bytes);
+            Expr combined = cast(uint_t, bytes[0]);
+            for (int b = 1; b < elem_bytes; b++) {
+                combined = combined | (cast(uint_t, bytes[b]) << (8 * b));
+            }
+            return Reinterpret::make(field_type, combined);
+        }
+
+        // Keep a packed scalar field as a bit-concatenation of adjacent bytes
+        // rather than immediately expanding it to a shift/or tree. The
+        // concat_bits lowering turns this into a dense byte shuffle followed
+        // by a vector reinterpret, which gives LLVM enough structure to issue
+        // one (possibly unaligned) wide load. Expanding here obscures that the
+        // byte loads are adjacent once individual extracts of the field have
+        // simplified, and commonly leaves one scalar load per byte.
+        Expr combined = concat_bits(bytes);
+        return field_type == combined.type() ? combined : Reinterpret::make(field_type, combined);
+    }
+
+    // Project field field_index's element elem_index (0 for a scalar
+    // field) out of a struct-typed Expr. Recurses through the only shapes a
+    // struct-typed Expr can legally have at this point in lowering: a
+    // literal struct_pack(), a Select between two struct-typed branches, a
+    // struct-typed Let (via struct_lets), or a genuine flattened Load.
+    Expr project_field(const Expr &struct_expr, int field_index, const Expr &elem_index) {
+        Expr e = resolve_struct_value(struct_expr);
+        const StructTypeInfo *info = e.type().struct_type();
+        internal_assert(info != nullptr) << "project_field applied to a non-struct-typed Expr.\n";
+        internal_assert(field_index >= 0 && field_index < (int)info->fields.size());
+
+        if (const Call *pack = e.as<Call>(); pack != nullptr && pack->is_intrinsic(Call::struct_pack)) {
+            return mutate(fold_field_of_pack(pack, *info, field_index, elem_index));
+        }
+
+        // e is itself an unresolved nested field() access (e.g. field(field(outer, "inner"), "a"),
+        // written as two separate front-end field() calls). We resolve the inner one first, then
+        // continue projecting the requested field out of *that* result.
+        // This is what makes nested struct fields work.
+        if (const Call *nested = e.as<Call>(); nested != nullptr && nested->is_intrinsic(Call::struct_field_read)) {
+            auto nested_field_index = as_const_int(nested->args[1]);
+            internal_assert(nested_field_index) << "struct_field_read's field index must be a constant.\n";
+            Expr resolved_inner = project_field(nested->args[0], (int)*nested_field_index, nested->args[2]);
+            return project_field(resolved_inner, field_index, elem_index);
+        }
+
+        if (const Select *sel = e.as<Select>()) {
+            return Select::make(mutate(sel->condition),
+                                project_field(sel->true_value, field_index, elem_index),
+                                project_field(sel->false_value, field_index, elem_index));
+        }
+
+        if (const Let *let = e.as<Let>()) {
+            ScopedBinding<Expr> bind(struct_lets, let->name, let->value);
+            return project_field(let->body, field_index, elem_index);
+        }
+
+        const Load *load = e.as<Load>();
+        user_assert(load != nullptr && load->type.is_struct())
+            << "This struct-typed Expr cannot be read with field(): a struct-typed Expr may only "
+            << "be a literal pack_struct(), a Select between two struct-typed Exprs, or the direct "
+            << "result of calling a struct-typed ImageParam/Buffer/Func, not: " << e << "\n";
+        return lower_field_read_from_load(load, info->fields[field_index].type, info->offsets[field_index], elem_index);
+    }
+
+protected:
+    Expr visit(const Call *op) override {
+        if (op->is_intrinsic(Call::struct_field_read)) {
+            // Deliberately don't generically mutate() op->args[0] first: see
+            // project_field/resolve_struct_value, which handle its only
+            // legal shapes directly. Running it through the generic Call
+            // visitor would (correctly, for any *other* context) flag it as
+            // a leaked struct-typed Expr, but here it's about to be
+            // legitimately consumed, not leaked.
+            auto field_index_opt = as_const_int(op->args[1]);
+            internal_assert(field_index_opt) << "struct_field_read's field index must be a constant.\n";
+            return project_field(op->args[0], (int)*field_index_opt, op->args[2]);
+        }
+
+        if (op->is_intrinsic(Call::struct_pack)) {
+            // Just mutate the pack elements. The actual packing is handled at
+            // consumption time (either by another field or by a Store).
+            return IRMutator::visit(op);
+        }
+
+        if (op->type.is_struct()) {
+            user_error << "Struct-typed Expr \"" << Expr(op) << "\" is not the direct argument of "
+                       << "field(), and is not a struct-typed Store's own value. A struct-typed Expr "
+                       << "may only appear in one of those two positions -- it can't be passed to an "
+                       << "arbitrary intrinsic or stored in a Let generically.\n";
+        }
+
+        return IRMutator::visit(op);
+    }
+
+    Stmt visit(const Store *op) override {
+        if (!op->value.type().is_struct()) {
+            return IRMutator::visit(op);
+        }
+
+        Type struct_t = op->value.type();
+        const StructTypeInfo &info = *struct_t.struct_type();
+        Expr dest_index = mutate(op->index);
+        Type index_type = dest_index.type();
+        Expr dest_byte_base = dest_index * make_const(index_type, struct_t.bytes());
+        Expr predicate = mutate(op->predicate);
+
+        vector<Stmt> stores;
+        for (int field_index = 0; field_index < (int)info.fields.size(); field_index++) {
+            const StructField &f = info.fields[field_index];
+            // A field whose own type is itself Type::Struct (a nested struct field) can't be
+            // split into byte-stores this way: the Reinterpret below fundamentally assumes
+            // f.type's *bits* describe its *bytes* (true for every ordinary scalar type), but a
+            // struct's bits() is not its (packed, possibly not power-of-two) byte size --
+            // Reinterpret would be asked to reinterpret between mismatched bit widths. Only
+            // supported for a fully-inlined struct producer (see lower_field_read_from_load's
+            // matching restriction on the read side).
+            user_assert(!f.type.is_struct())
+                << "Storing a nested struct field (a field whose own type is Type::Struct) as "
+                << "part of a materialized (non-inlined) struct-typed Store is not supported. "
+                << "Field \"" << f.name << "\" of type " << f.type << " in a Store to "
+                << op->name << " was requested this way.\n";
+            int extent = f.array_extent.value_or(1);
+            int elem_bytes = f.type.bytes();
+            for (int e = 0; e < extent; e++) {
+                Expr value = project_field(op->value, field_index, make_const(Int(32), e));
+                Expr bits = elem_bytes == 1 ? value : Reinterpret::make(UInt(8 * elem_bytes), value);
+                Expr elem_byte_base = dest_byte_base +
+                                      make_const(index_type, info.offsets[field_index] + e * elem_bytes);
+                for (int b = 0; b < elem_bytes; b++) {
+                    Expr byte_val = elem_bytes == 1 ? bits : extract_bits(UInt(8), bits, make_const(Int(32), 8 * b));
+                    Expr byte_idx = b == 0 ? elem_byte_base : elem_byte_base + make_const(index_type, b);
+                    stores.push_back(Store::make(op->name, byte_val, byte_idx, op->param, predicate, ModulusRemainder(), op->is_streaming));
+                }
+            }
+        }
+        return Block::make(stores);
+    }
+
+    Expr visit(const Let *op) override {
+        if (op->value.type().is_struct()) {
+            ScopedBinding bind(struct_lets, op->name, op->value);
+            return mutate(op->body);
+        }
+        return IRMutator::visit(op);
+    }
+
+    Expr visit(const Variable *op) override {
+        if (op->type.is_struct()) {
+            user_assert(struct_lets.contains(op->name))
+                << "Struct-typed variable \"" << op->name << "\" is unbound; a struct-typed Expr "
+                << "may only be the value of a struct-typed Func/Store, or a direct field() "
+                << "argument, in v1.\n";
+            return mutate(struct_lets.get(op->name));
+        }
+        return op;
+    }
+};
+
+// After LowerStructTypes, no struct-typed node of any kind
+// should survive in the Stmt. This is purely defensive.
+class CheckNoStructTypesRemain : public IRGraphVisitor {
+    using IRGraphVisitor::visit;
+
+    static void fail(const Expr &e, const char *what) {
+        internal_error << "Struct-typed " << what << " survived struct type lowering: " << e << "\n";
+    }
+
+protected:
+    void visit(const Call *op) override {
+        if (op->type.is_struct() || op->is_intrinsic(Call::struct_field_read) || op->is_intrinsic(Call::struct_pack)) {
+            fail(op, "Call");
+        }
+        IRGraphVisitor::visit(op);
+    }
+    void visit(const Load *op) override {
+        if (op->type.is_struct()) {
+            fail(op, "Load");
+        }
+        IRGraphVisitor::visit(op);
+    }
+    void visit(const Variable *op) override {
+        if (op->type.is_struct()) {
+            fail(op, "Variable");
+        }
+        IRGraphVisitor::visit(op);
+    }
+    void visit(const Let *op) override {
+        if (op->value.type().is_struct()) {
+            fail(op->value, "Let value");
+        }
+        IRGraphVisitor::visit(op);
+    }
+    void visit(const Store *op) override {
+        if (op->value.type().is_struct()) {
+            fail(op->value, "Store value");
+        }
+        IRGraphVisitor::visit(op);
+    }
+};
+
+}  // namespace
+
+Stmt lower_struct_types(const Stmt &s) {
+    Stmt result = LowerStructTypesMutator()(s);
+    CheckNoStructTypesRemain checker;
+    result.accept(&checker);
+    return result;
+}
+
+}  // namespace Internal
+}  // namespace Halide
diff --git a/src/LowerStructTypes.h b/src/LowerStructTypes.h
new file mode 100644
index 000000000000..febfdf7f9b39
--- /dev/null
+++ b/src/LowerStructTypes.h
@@ -0,0 +1,27 @@
+#ifndef HALIDE_LOWER_STRUCT_TYPES_H
+#define HALIDE_LOWER_STRUCT_TYPES_H
+
+/** \file
+ * Defines the lowering pass that rewrites field()/pack_struct() intrinsics
+ * into ordinary loads and Reinterpret nodes.
+ *
+ * Struct types are treated as an ordinary, opaque, N-byte element type
+ * throughout scheduling, bounds inference, and storage flattening. This
+ * pass must run *after* storage_flattening, once every struct-typed
+ * Provide/Call has become a concrete Store/Load with a real flat byte
+ * index. It splits a struct-typed Store into per-field byte stores and
+ * rewrites field() reads into byte-shifted loads, using that flat index.
+ */
+
+#include "IR.h"
+
+namespace Halide {
+namespace Internal {
+
+/** Lower all field()/pack_struct() usage in a Stmt. */
+Stmt lower_struct_types(const Stmt &s);
+
+}  // namespace Internal
+}  // namespace Halide
+
+#endif
diff --git a/src/Parameter.cpp b/src/Parameter.cpp
index 76f43177f495..9cc5e8b26879 100644
--- a/src/Parameter.cpp
+++ b/src/Parameter.cpp
@@ -211,7 +211,7 @@ const halide_buffer_t *Parameter::raw_buffer() const {
 void Parameter::set_buffer(const Buffer<> &b) {
     check_is_buffer();
     if (b.defined()) {
-        user_assert(contents->type == b.type())
+        user_assert(contents->type.is_compatible_for_buffer_bind(b.type()))
             << "Can't bind Parameter " << name()
             << " of type " << contents->type
             << " to Buffer " << b.name()
diff --git a/src/Serialization.cpp b/src/Serialization.cpp
index 2dd7bf4f33aa..c4aeadcb528f 100644
--- a/src/Serialization.cpp
+++ b/src/Serialization.cpp
@@ -401,6 +401,11 @@ Offset<String> Serializer::serialize_string(FlatBufferBuilder &builder, const st
 }
 
 Offset<Serialize::Type> Serializer::serialize_type(FlatBufferBuilder &builder, const Type &type) {
+    // The serialization schema's TypeCode has no struct variant (a struct also
+    // carries a field layout the flat Type table can't hold), so reject struct
+    // types outright rather than emitting an out-of-range code.
+    user_assert(!type.is_struct())
+        << "Serialization of struct types is not supported.\n";
     const int bits = type.bits();
     const int lanes = type.lanes();
     halide_type_code_t code = type.code();
diff --git a/src/Type.cpp b/src/Type.cpp
index 824e105a40e4..680b2d29ec80 100644
--- a/src/Type.cpp
+++ b/src/Type.cpp
@@ -1,8 +1,10 @@
 #include "ConstantBounds.h"
 #include "IR.h"
 #include <cfloat>
+#include <map>
 #include <mutex>
 #include <sstream>
+#include <tuple>
 
 namespace Halide {
 
@@ -25,8 +27,110 @@ int64_t min_int(int bits) {
 
 }  // namespace
 
+bool StructField::operator==(const StructField &other) const {
+    return name == other.name && type == other.type && array_extent == other.array_extent;
+}
+
+bool StructField::operator<(const StructField &other) const {
+    if (name != other.name) {
+        return name < other.name;
+    }
+    if (type != other.type) {
+        return type < other.type;
+    }
+    return array_extent < other.array_extent;
+}
+
+int StructTypeInfo::find_field(const std::string &name) const {
+    for (size_t i = 0; i < fields.size(); i++) {
+        if (fields[i].name == name) {
+            return (int)i;
+        }
+    }
+    return -1;
+}
+
+int Type::bytes() const {
+    if (is_struct()) {
+        const StructTypeInfo *info = struct_type();
+        return info ? info->total_bytes : 0;
+    }
+    return (bits() + 7) / 8;
+}
+
+Type Type::Struct(const std::vector<StructField> &fields) {
+    user_assert(!fields.empty()) << "Type::Struct requires at least one field.\n";
+
+    // Deliberately leaked: like halide_handle_cplusplus_type, this info table
+    // must remain valid for the lifetime of the program, since the intern
+    // table holds a (non-owning) pointer to it and Types get freely copied
+    // and compared throughout compilation.
+    auto *info = new StructTypeInfo();
+    info->fields = fields;
+    info->offsets.reserve(fields.size());
+
+    int offset = 0;
+    for (const auto &f : fields) {
+        user_assert(!f.type.is_struct() || f.type.struct_type() != nullptr)
+            << "Struct field \"" << f.name << "\" has an invalid nested struct type.\n";
+        user_assert(f.array_extent.value_or(1) > 0)
+            << "Struct field \"" << f.name << "\" has a non-positive array extent.\n";
+        info->offsets.push_back(offset);
+        offset += f.type.bytes() * f.array_extent.value_or(1);
+    }
+    info->total_bytes = offset;
+
+    // A struct has its own honest type code; its byte size lives in the interned
+    // StructTypeInfo (and is carried in the ABI's reserved field, see to_abi()).
+    // type_bits is not meaningful for a struct -- it's set to a byte's worth so
+    // the erased ABI tag is well-formed, but bytes() is the real size.
+    Type t(StructKind, 8, 1);
+    t.metadata_index_ = Internal::intern_struct_type(info);
+    return t;
+}
+
+bool Type::same_struct_type(const Type &other) const {
+    const StructTypeInfo *a = struct_type();
+    const StructTypeInfo *b = other.struct_type();
+
+    if (a == b) {
+        return true;
+    }
+    if (a == nullptr || b == nullptr) {
+        return false;
+    }
+    return a->fields == b->fields;
+}
+
+bool Type::operator<(const Type &other) const {
+    if (std::tie(type_code, type_bits, type_lanes) <
+        std::tie(other.type_code, other.type_bits, other.type_lanes)) {
+        return true;
+    }
+    if (std::tie(other.type_code, other.type_bits, other.type_lanes) <
+        std::tie(type_code, type_bits, type_lanes)) {
+        return false;
+    }
+    if (code() == Handle) {
+        return handle_type() < other.handle_type();
+    }
+    if (is_struct()) {
+        // Equal (type_code, type_bits, type_lanes) above already implies both
+        // are structs here, so we only order by field layout.
+        if (same_struct_type(other)) {
+            // Consistent with operator==: two struct types with identical
+            // field lists always compare as neither-less-than-the-other, even
+            // if they're backed by different interned StructTypeInfo tables.
+            return false;
+        }
+        return struct_type()->fields < other.struct_type()->fields;
+    }
+    return false;
+}
+
 /** Return an expression which is the maximum value of this type */
 Halide::Expr Type::max() const {
+    user_assert(!is_struct()) << "Type::max() is not defined for a struct type: " << *this << "\n";
     if (is_vector()) {
         return Internal::Broadcast::make(element_of().max(), lanes());
     } else if (is_int()) {
@@ -34,7 +138,7 @@ Halide::Expr Type::max() const {
     } else if (is_uint()) {
         return Internal::UIntImm::make(*this, max_uint(bits()));
     } else {
-        internal_assert(is_float());
+        internal_assert(is_float()) << "Type::max() is not defined for " << *this << "\n";
         if (bits() == 16) {
             return Internal::FloatImm::make(*this, (double)float16_t::make_infinity());
         } else if (bits() == 32) {
@@ -51,6 +155,7 @@ Halide::Expr Type::max() const {
 
 /** Return an expression which is the minimum value of this type */
 Halide::Expr Type::min() const {
+    user_assert(!is_struct()) << "Type::min() is not defined for a struct type: " << *this << "\n";
     if (is_vector()) {
         return Internal::Broadcast::make(element_of().min(), lanes());
     } else if (is_int()) {
@@ -58,7 +163,7 @@ Halide::Expr Type::min() const {
     } else if (is_uint()) {
         return Internal::UIntImm::make(*this, 0);
     } else {
-        internal_assert(is_float());
+        internal_assert(is_float()) << "Type::min() is not defined for " << *this << "\n";
         if (bits() == 16) {
             return Internal::FloatImm::make(*this, (double)float16_t::make_negative_infinity());
         } else if (bits() == 32) {
@@ -225,46 +330,99 @@ bool Type::can_represent(double x) const {
 
 namespace Internal {
 namespace {
-std::mutex &handle_type_intern_mutex() {
-    static std::mutex m;
-    return m;
+
+/** A process-wide table that maps externally-owned, program-lifetime metadata
+ * pointers (handle-type descriptors, struct-type layouts) to small 1-based
+ * indices, so a `Type` can reference them in 4 bytes instead of an 8-byte
+ * inline pointer. Index 0 is reserved to mean "none". Pointers are deduped by
+ * identity (the pointees are stable, and distinct pointers with equal content
+ * are reconciled by the deep comparisons in Type). Entries are never removed. */
+template<typename T>
+class InternTable {
+    std::mutex mutex;
+    std::vector<const T *> table;
+
+public:
+    uint32_t intern(const T *ptr) {
+        if (ptr == nullptr) {
+            return 0;
+        }
+        std::scoped_lock lock(mutex);
+        for (size_t i = 0; i < table.size(); i++) {
+            if (table[i] == ptr) {
+                return (uint32_t)(i + 1);  // +1: index 0 is reserved for null
+            }
+        }
+        table.push_back(ptr);
+        internal_assert(table.size() < (size_t)0xffffffffu) << "intern table overflow";
+        return (uint32_t)table.size();
+    }
+
+    const T *get(uint32_t index) {
+        if (index == 0) {
+            return nullptr;
+        }
+        std::scoped_lock lock(mutex);
+        internal_assert(index <= table.size()) << "invalid intern index";
+        return table[index - 1];
+    }
+};
+
+InternTable<halide_handle_cplusplus_type> &handle_type_intern_table() {
+    static InternTable<halide_handle_cplusplus_type> t;
+    return t;
 }
-std::vector<const halide_handle_cplusplus_type *> &handle_type_intern_table() {
-    static std::vector<const halide_handle_cplusplus_type *> t;
+
+InternTable<StructTypeInfo> &struct_type_intern_table() {
+    static InternTable<StructTypeInfo> t;
     return t;
 }
+
 }  // namespace
 
 uint32_t intern_handle_type(const halide_handle_cplusplus_type *handle_type) {
-    // The overwhelmingly common case: a non-handle type. Costs a null check,
-    // no lock.
-    if (handle_type == nullptr) {
-        return 0;
-    }
-    std::scoped_lock lock(handle_type_intern_mutex());
-    auto &table = handle_type_intern_table();
-    // Dedup by pointer identity. Handle types are rare and few, so a linear
-    // scan is cheap; the pointees are externally owned and stable, so pointer
-    // identity is a safe key. (Distinct pointers with equal content still
-    // compare equal via Type::same_handle_type's deep comparison.)
-    for (size_t i = 0; i < table.size(); i++) {
-        if (table[i] == handle_type) {
-            return (uint32_t)(i + 1);  // +1: index 0 is reserved for null
-        }
-    }
-    table.push_back(handle_type);
-    internal_assert(table.size() < (size_t)0xffffffffu) << "handle-type intern table overflow";
-    return (uint32_t)table.size();
+    // The overwhelmingly common case (a non-handle type) is the nullptr fast
+    // path inside intern(), which costs just a null check and no lock.
+    return handle_type_intern_table().intern(handle_type);
 }
 
 const halide_handle_cplusplus_type *get_interned_handle_type(uint32_t index) {
-    if (index == 0) {
-        return nullptr;
+    return handle_type_intern_table().get(index);
+}
+
+uint32_t intern_struct_type(const StructTypeInfo *struct_type) {
+    internal_assert(struct_type != nullptr) << "intern_struct_type(nullptr)";
+    return struct_type_intern_table().intern(struct_type);
+}
+
+const StructTypeInfo *get_interned_struct_type(uint32_t index) {
+    return struct_type_intern_table().get(index);
+}
+
+uint32_t intern_opaque_struct_type(int total_bytes) {
+    // Reconstructing a struct Type from its ABI form recovers only the size,
+    // not the field layout. Memoize one field-opaque StructTypeInfo per size so
+    // repeated round-trips (e.g. Buffer::type()) don't leak a table entry each.
+    static std::mutex m;
+    static std::map<int, const StructTypeInfo *> by_size{};
+    const StructTypeInfo *info = nullptr;
+    {
+        std::scoped_lock lock(m);
+        auto it = by_size.find(total_bytes);
+        if (it == by_size.end()) {
+            // A single opaque byte-blob field of the right length; deliberately
+            // leaked, like every other interned StructTypeInfo.
+            auto *fresh = new StructTypeInfo();
+            fresh->fields = {StructField{"", UInt(8), total_bytes > 0 ? std::optional<int>(total_bytes) : std::nullopt}};
+            fresh->offsets = {0};
+            fresh->total_bytes = total_bytes;
+            by_size[total_bytes] = fresh;
+            info = fresh;
+        } else {
+            info = it->second;
+        }
     }
-    std::scoped_lock lock(handle_type_intern_mutex());
-    auto &table = handle_type_intern_table();
-    internal_assert(index <= table.size()) << "invalid handle-type intern index";
-    return table[index - 1];
+    return intern_struct_type(info);
 }
 }  // namespace Internal
 
@@ -294,7 +452,12 @@ std::string type_to_c_type(Type type, bool include_space, bool c_plus_plus) {
     bool needs_space = true;
     ostringstream oss;
 
-    if (type.is_bfloat()) {
+    if (type.is_struct()) {
+        // The C backend byte-addresses struct storage (field access is lowered
+        // to byte loads/stores), so a struct's C element type is a raw byte;
+        // callers scale allocations by Type::bytes() for the true size.
+        oss << "uint8_t";
+    } else if (type.is_bfloat()) {
         oss << "bfloat" << type.bits() << "_t";
     } else if (type.is_float()) {
         if (type.bits() == 32) {
diff --git a/src/Type.h b/src/Type.h
index fae4db772562..ecc8979ade7c 100644
--- a/src/Type.h
+++ b/src/Type.h
@@ -7,6 +7,8 @@
 #include "runtime/HalideRuntime.h"
 #include <cstdint>
 #include <optional>
+#include <string>
+#include <vector>
 
 /** \file
  * Defines halide types
@@ -268,6 +270,9 @@ struct halide_handle_traits {
 
 namespace Halide {
 
+struct StructField;
+struct StructTypeInfo;
+
 namespace Internal {
 struct ConstantInterval;
 
@@ -280,6 +285,19 @@ struct ConstantInterval;
  * 0 without locking, so non-handle type construction pays nothing. */
 uint32_t intern_handle_type(const halide_handle_cplusplus_type *handle_type);
 const halide_handle_cplusplus_type *get_interned_handle_type(uint32_t index);
+
+/** Struct types carry their field layout the same way handles carry their C++
+ * pointee metadata: by a 4-byte index into a process-wide intern table of
+ * (leaked, program-lifetime) StructTypeInfo pointers, rather than an inline
+ * pointer, so a struct-typed `Type` stays 8 bytes. Index 0 means "no struct
+ * metadata"; a well-formed struct Type always has a nonzero index. */
+uint32_t intern_struct_type(const StructTypeInfo *struct_type);
+const StructTypeInfo *get_interned_struct_type(uint32_t index);
+
+/** Intern (memoized by size) a field-opaque struct of the given total byte
+ * size, used to reconstruct a struct Type from its ABI form, which records
+ * only the size. Returns a nonzero struct-table index. */
+uint32_t intern_opaque_struct_type(int total_bytes);
 }  // namespace Internal
 
 struct Expr;
@@ -294,9 +312,12 @@ struct Type {
     halide_type_code_t type_code;
     uint8_t type_bits;
     uint16_t type_lanes;
-    // Index into the process-wide handle-type intern table (0 == no handle
-    // metadata).
-    uint32_t handle_index_ = 0;
+    // Index into a process-wide intern table. Its meaning is selected by
+    // type_code: for a Handle it indexes the handle-type table (0 == void *),
+    // for a Struct it indexes the struct-type table (0 == malformed struct).
+    // Unused (0) for numeric types. A Type is never both a handle and a struct,
+    // so one slot serves both and keeps sizeof(Type) == 8.
+    uint32_t metadata_index_ = 0;
 
 public:
     /** Aliases for halide_type_code_t values for legacy compatibility
@@ -307,6 +328,7 @@ struct Type {
     static constexpr halide_type_code_t Float = halide_type_float;
     static constexpr halide_type_code_t BFloat = halide_type_bfloat;
     static constexpr halide_type_code_t Handle = halide_type_handle;
+    static constexpr halide_type_code_t StructKind = halide_type_struct;
     // @}
 
     /** Exposed so code that needs to reason about the maximum representable
@@ -314,11 +336,35 @@ struct Type {
      * it from Type itself instead of hardcoding a width. */
     static constexpr int kLanesBits = 8 * sizeof(type_lanes);
 
-    /** The number of bytes required to store a single scalar value of this type. Ignores vector lanes. */
-    int bytes() const {
-        return (bits() + 7) / 8;
+    /** The number of bytes required to store a single scalar value of this
+     * type. Ignores vector lanes. For a struct type, this is the packed
+     * size of the whole struct. */
+    int bytes() const;
+
+    /** Construct a packed, named-field aggregate type, matching a C struct
+     * with no alignment padding. The resulting Type has its own honest type
+     * code (Type::Struct); its packed byte size is recorded so it survives
+     * erasure to the ABI (see to_abi()). */
+    static Type Struct(const std::vector<StructField> &fields);
+
+    /** Is this type a struct type, as constructed by `Type::Struct`? */
+    HALIDE_ALWAYS_INLINE
+    bool is_struct() const {
+        return code() == StructKind;
+    }
+
+    /** If this is a struct type, its field list/offsets. Null otherwise.
+     * Backed by an intern index rather than an inline pointer (like
+     * handle_type()), so this is an accessor rather than a public field. */
+    const StructTypeInfo *struct_type() const {
+        return is_struct() ? Internal::get_interned_struct_type(metadata_index_) : nullptr;
     }
 
+    /** Check that the struct layout of two struct types matches structurally
+     * (same field names, types, order, and array extents), not just by
+     * pointer identity of the interned `StructTypeInfo`. */
+    bool same_struct_type(const Type &other) const;
+
     // Default ctor initializes everything to predictable-but-unlikely values
     constexpr Type()
         : type_code(Handle), type_bits(0), type_lanes(0) {
@@ -331,7 +377,7 @@ struct Type {
     HALIDE_ALWAYS_INLINE
     Type(halide_type_code_t code, int bits, int lanes, const halide_handle_cplusplus_type *handle_type = nullptr)
         : type_code(code), type_bits((uint8_t)bits), type_lanes((uint16_t)lanes),
-          handle_index_(handle_type ? Internal::intern_handle_type(handle_type) : 0) {
+          metadata_index_(handle_type ? Internal::intern_handle_type(handle_type) : 0) {
         internal_assert(lanes == (int)type_lanes)
             << "Halide only supports vector types with up to 65535 lanes. " << lanes << " lanes requested.";
         internal_assert(bits == (int)type_bits)
@@ -346,23 +392,41 @@ struct Type {
     /** Trivial copy assignment operator. */
     Type &operator=(const Type &that) = default;
 
-    /** Construct a (scalar) language Type from an ABI element type. */
+    /** Construct a (scalar) language Type from an ABI element type. A struct
+     * ABI type carries only its packed byte size, not its field layout (just
+     * as a buffer of handles forgets its C++ pointee type), so it rebuilds a
+     * faithful-sized but field-opaque struct -- enough for bytes()/to_abi() to
+     * round-trip and for buffer-bind compatibility, but not for field access. */
     HALIDE_ALWAYS_INLINE
     Type(halide_type_t that, const halide_handle_cplusplus_type *handle_type = nullptr)
         : Type(that.code, that.bits, 1, handle_type) {
+        if (that.code == halide_type_struct) {
+            metadata_index_ = Internal::intern_opaque_struct_type(that.reserved);
+        }
     }
 
     /** Erase this language Type to the ABI/runtime halide_type_t for use in
      * runtime calls, buffer/argument metadata, etc. This is a *checked*
      * erasure: the ABI element type has no lanes, so a vector type
      * (lanes >= 2) has no image here and asserts rather than silently
-     * dropping its lanes. */
+     * dropping its lanes. A struct type erases to code halide_type_struct
+     * carrying its packed byte size in the reserved field, so the aggregate's
+     * true size survives the boundary (an ordinary numeric type erases with
+     * reserved == 0). */
     HALIDE_ALWAYS_INLINE
     halide_type_t to_abi() const {
         internal_assert(type_lanes < 2)
             << "Cannot erase a vector type with " << type_lanes
             << " lanes to a scalar ABI halide_type_t.\n";
-        return halide_type_t(code(), type_bits);
+        halide_type_t t(code(), type_bits);
+        if (is_struct()) {
+            int n = bytes();
+            internal_assert(n <= 0xffff)
+                << "Struct type of " << n << " bytes is too large to erase to "
+                << "the ABI; the halide_type_t size field is 16 bits.\n";
+            t.reserved = (uint16_t)n;
+        }
+        return t;
     }
 
     /** Return the underlying data type of an element as an enum value. */
@@ -371,7 +435,10 @@ struct Type {
         return type_code;
     }
 
-    /** Return the bit size of a single element of this type. */
+    /** Return the bit size of a single element of this type. Not meaningful
+     * for a struct type (whose real size is a byte count that need not be a
+     * whole number of machine words, and can exceed what this field holds) --
+     * use bytes() for a struct's size. */
     HALIDE_ALWAYS_INLINE
     int bits() const {
         return type_bits;
@@ -388,7 +455,13 @@ struct Type {
     Type with_code(halide_type_code_t new_code) const {
         Type t = *this;
         t.type_code = new_code;
-        t.handle_index_ = new_code != code() ? 0 : handle_index_;  // Changing the type code invalidates any handle metadata.
+        if (new_code != code()) {
+            // The kind of type is changing, so any handle or struct metadata
+            // (which described the old kind) no longer applies. This keeps the
+            // code-based is_struct()/is_handle() consistent with metadata: a
+            // struct only stays a struct while its code stays halide_type_struct.
+            t.metadata_index_ = 0;
+        }
         return t;
     }
 
@@ -397,7 +470,10 @@ struct Type {
     Type with_bits(int new_bits) const {
         Type t = *this;
         t.type_bits = (uint8_t)new_bits;
-        t.handle_index_ = new_bits != bits() ? 0 : handle_index_;  // Changing the bit width invalidates any handle metadata.
+        // The type code (and thus its kind) is unchanged, so handle/struct
+        // metadata survives -- bits() is not meaningful for either anyway, and
+        // a struct must not silently degrade into a plain integer partway
+        // through lowering (e.g. StorageFlattening rounding bits up to a byte).
         return t;
     }
 
@@ -407,6 +483,9 @@ struct Type {
     Type with_lanes(int new_lanes) const {
         Type t = *this;
         t.type_lanes = (uint16_t)new_lanes;
+        // Handle and struct metadata describe the *element* type, so they
+        // survive a change in lane count (a vector of handles/structs keeps
+        // the same element metadata); no need to touch metadata_index_.
         return t;
     }
 
@@ -442,9 +521,11 @@ struct Type {
      * index rather than stored inline, so this is now an accessor rather than a
      * public field. */
     const halide_handle_cplusplus_type *handle_type() const {
-        // Inline the common (non-handle) case so it needs no call or table
-        // lookup; only a real handle index reaches the intern table.
-        return handle_index_ == 0 ? nullptr : Internal::get_interned_handle_type(handle_index_);
+        // Inline the common case so it needs no call or table lookup. The
+        // shared metadata_index_ only names the handle table when this is a
+        // handle; for a struct it indexes the struct table instead, so guard
+        // on the code to avoid a cross-table lookup.
+        return (code() != Handle || metadata_index_ == 0) ? nullptr : Internal::get_interned_handle_type(metadata_index_);
     }
 
     /** Is this type boolean (represented as UInt(1))? */
@@ -485,7 +566,10 @@ struct Type {
         return code() == Int;
     }
 
-    /** Is this type an unsigned integer type? */
+    /** Is this type an unsigned integer type? A struct type has its own honest
+     * code (halide_type_struct), so this reports false for it -- struct types
+     * no longer masquerade as UInt(8), and numeric reasoning is naturally
+     * excluded without a separate is_struct() guard. */
     HALIDE_ALWAYS_INLINE
     bool is_uint() const {
         return code() == UInt;
@@ -521,8 +605,12 @@ struct Type {
     /** Compare two types for equality */
     HALIDE_ALWAYS_INLINE
     bool operator==(const Type &other) const {
+        // Equal type_code already implies both are (or aren't) handles/structs,
+        // so the deep metadata comparison only needs to run for that code.
         return type_code == other.type_code && type_bits == other.type_bits &&
-               type_lanes == other.type_lanes && (code() != Handle || same_handle_type(other));
+               type_lanes == other.type_lanes &&
+               (code() != Handle || same_handle_type(other)) &&
+               (code() != StructKind || same_struct_type(other));
     }
 
     /** Compare two types for inequality */
@@ -530,6 +618,20 @@ struct Type {
         return !(*this == other);
     }
 
+    /** Is `buffer_type` an acceptable type for a `halide_buffer_t` bound to a
+     * Parameter/ImageParam declared with this Type? This mirrors the runtime
+     * type check emitted by AddImageChecks: the buffer's ABI element type must
+     * equal this type's. For a struct that means the same struct code and
+     * packed byte size -- the buffer only carries the ABI halide_type_t, which
+     * forgets a struct's field layout (like a buffer of handles forgets its
+     * pointee type), so two structs of equal size are ABI-compatible. */
+    bool is_compatible_for_buffer_bind(const Type &buffer_type) const {
+        if (*this == buffer_type) {
+            return true;
+        }
+        return is_struct() && buffer_type.to_abi() == to_abi();
+    }
+
     /** Compare a language type to an ABI element type. Equal iff this type is a
      * single element (not a vector) with the same code and bits. */
     bool operator==(const halide_type_t &other) const {
@@ -541,17 +643,12 @@ struct Type {
         return !(*this == other);
     }
 
-    /** Compare ordering of two types so they can be used in certain containers and algorithms */
-    bool operator<(const Type &other) const {
-        if (std::tie(type_code, type_bits, type_lanes) <
-            std::tie(other.type_code, other.type_bits, other.type_lanes)) {
-            return true;
-        }
-        if (code() == Handle) {
-            return handle_type() < other.handle_type();
-        }
-        return false;
-    }
+    /** Compare ordering of two types so they can be used in certain
+     * containers and algorithms. Consistent with operator==: two struct
+     * types that compare equal (same field names/types/order/extents, per
+     * same_struct_type) always compare as neither-less-than-the-other,
+     * regardless of whether they share a StructTypeInfo pointer. */
+    bool operator<(const Type &other) const;
 
     /** Produce the scalar type (that of a single element) of this vector type */
     HALIDE_ALWAYS_INLINE
@@ -591,9 +688,35 @@ struct Type {
     Expr min() const;
 };
 
-static_assert(sizeof(Type) == 8, "Halide::Type is a code+bits+lanes triple plus a 4-byte handle-type index");
+// Type stays 8 bytes: struct field layout is referenced by the same 4-byte
+// intern index that carries handle metadata (a Type is never both), not by an
+// inline side-table pointer. It must also stay trivially copyable, since Types
+// are copied all over the compiler.
+static_assert(sizeof(Type) == 8, "Type should stay 8 bytes");
 static_assert(std::is_trivially_copyable_v<Type>, "Type must stay trivially copyable");
 
+/** One field of a struct: a name, a type, and (for array fields)
+ * a number of elements. */
+struct StructField {
+    std::string name;
+    Type type;
+    std::optional<int> array_extent = std::nullopt;
+
+    bool operator==(const StructField &other) const;
+    bool operator<(const StructField &other) const;
+};
+
+/** The compile-time information a struct carries: its field list
+ * and the byte offset of each field */
+struct StructTypeInfo {
+    std::vector<StructField> fields;
+    std::vector<int> offsets;  // one per field, in bytes
+    int total_bytes = 0;
+
+    /** Find a field by name. Returns -1 if not found. */
+    int find_field(const std::string &name) const;
+};
+
 /** Constructing a signed integer type */
 HALIDE_ALWAYS_INLINE Type Int(int bits, int lanes = 1) {
     return Type(Type::Int, bits, lanes);
diff --git a/src/runtime/HalideRuntime.h b/src/runtime/HalideRuntime.h
index e1610b9028c8..e61c5a38b906 100644
--- a/src/runtime/HalideRuntime.h
+++ b/src/runtime/HalideRuntime.h
@@ -499,6 +499,8 @@ typedef enum halide_type_code_t
     halide_type_float = 2,   ///< IEEE floating point numbers
     halide_type_handle = 3,  ///< opaque pointer type (void *)
     halide_type_bfloat = 4,  ///< floating point numbers in the bfloat format
+    halide_type_struct = 5,  ///< a packed, by-value aggregate; its total byte
+                             ///< size is carried in halide_type_t::reserved
 } halide_type_code_t;
 
 // Note that while __attribute__ can go before or after the declaration,
@@ -532,7 +534,9 @@ struct halide_type_t {
     HALIDE_ATTRIBUTE_ALIGN(1)
     uint8_t bits;
 
-    /** Reserved for future element-kind payloads. */
+    /** Element-kind payload. For a struct type (code == halide_type_struct)
+     * this holds the total packed byte size of the aggregate; the ordinary
+     * `bits` field is not meaningful for a struct. Zero for all other types. */
     HALIDE_ATTRIBUTE_ALIGN(2)
     uint16_t reserved;
 
@@ -563,9 +567,11 @@ struct halide_type_t {
         return static_cast<uint32_t>(*this) < static_cast<uint32_t>(other);
     }
 
-    /** Size in bytes for a single element of this type. */
+    /** Size in bytes for a single element of this type. For a struct type the
+     * packed aggregate size is carried in `reserved` rather than derived from
+     * `bits`. */
     HALIDE_ALWAYS_INLINE constexpr int bytes() const {
-        return (bits + 7) / 8;
+        return code == halide_type_struct ? reserved : (bits + 7) / 8;
     }
 
     HALIDE_ALWAYS_INLINE constexpr operator uint32_t() const {
@@ -2502,6 +2508,8 @@ inline std::ostream &operator<<(std::ostream &os, const halide_type_t &type) {
         return os << "(void*)";
     case halide_type_bfloat:
         return os << "bfloat" << (int)type.bits;
+    case halide_type_struct:
+        return os << "struct(" << (int)type.reserved << " bytes)";
     }
     return os;
 }
diff --git a/src/runtime/halide_buffer_t.cpp b/src/runtime/halide_buffer_t.cpp
index a3bf1e68f9b2..5b47b9c4d373 100644
--- a/src/runtime/halide_buffer_t.cpp
+++ b/src/runtime/halide_buffer_t.cpp
@@ -97,14 +97,17 @@ halide_buffer_t *_halide_buffer_init(halide_buffer_t *dst,
                                      void *host,
                                      uint64_t device,
                                      const halide_device_interface_t *device_interface,
-                                     int type_code, int type_bits,
+                                     int type_code, int type_bits, int type_reserved,
                                      int dimensions,
                                      halide_dimension_t *shape,
                                      uint64_t flags) {
     dst->host = (uint8_t *)host;
     dst->device = device;
     dst->device_interface = device_interface;
+    // type_reserved carries a struct element type's packed byte size (0 for
+    // ordinary types); see halide_type_t::reserved and Type::to_abi().
     dst->type = {(halide_type_code_t)type_code, (uint8_t)type_bits};
+    dst->type.reserved = (uint16_t)type_reserved;
     dst->dimensions = dimensions;
     dst->dim = dst_shape;
     if (shape != dst->dim) {
diff --git a/src/runtime/to_string.cpp b/src/runtime/to_string.cpp
index 7bb36d76d4cd..05f90cae556f 100644
--- a/src/runtime/to_string.cpp
+++ b/src/runtime/to_string.cpp
@@ -281,6 +281,10 @@ WEAK char *halide_type_to_string(char *dst, char *end, const halide_type_t *t) {
     case halide_type_bfloat:
         code_name = "bfloat";
         break;
+    case halide_type_struct:
+        // A struct's size lives in `reserved`, not `bits`; print the byte count.
+        dst = halide_string_to_string(dst, end, "struct");
+        return halide_uint64_to_string(dst, end, t->reserved, 1);
     default:
         code_name = "bad_type_code";
         break;
diff --git a/test/correctness/CMakeLists.txt b/test/correctness/CMakeLists.txt
index 5d102a63c202..6794e958e5e1 100644
--- a/test/correctness/CMakeLists.txt
+++ b/test/correctness/CMakeLists.txt
@@ -182,6 +182,7 @@ tests(
     gpu_reuse_shared_memory.cpp
     gpu_specialize.cpp
     gpu_store_in_register_with_no_lanes_loop.cpp
+    gpu_struct_type.cpp
     gpu_sum_scan.cpp
     gpu_texture.cpp
     gpu_thread_barrier.cpp
@@ -361,6 +362,8 @@ tests(
     strict_float_bounds.cpp
     strict_fma.cpp
     strided_load.cpp
+    struct_type.cpp
+    struct_type_dot_product.cpp
     sve_codegen_predicated.cpp
     sve_codegen_reinterpret.cpp
     target.cpp
diff --git a/test/correctness/gpu_struct_type.cpp b/test/correctness/gpu_struct_type.cpp
new file mode 100644
index 000000000000..f7c3934d7580
--- /dev/null
+++ b/test/correctness/gpu_struct_type.cpp
@@ -0,0 +1,77 @@
+#include "Halide.h"
+#include <cmath>
+#include <cstdio>
+
+using namespace Halide;
+
+// GPU torture test for struct types (see Type::Struct, test/correctness/struct_type.cpp):
+// a struct-valued Func consumed on the GPU, once compute_at'd inside a
+// gpu_threads() loop (thread-local/register scratch on the device) and once
+// compute_at'd at the gpu_blocks() level and stored in GPUShared memory.
+//
+// Device-side codegen is a completely separate code generator per backend,
+// with its own type-to-size mapping -- this is what caught two bugs beyond
+// the CPU LLVM backend's own (CodeGen_LLVM::llvm_type_of mapping a struct to
+// i8 instead of an array of its true byte size): CodeGen_Metal_Dev's
+// non-shared Allocate path under-sized a struct-typed stack scratch buffer
+// the same way, and FindIntrinsics.cpp's generic lower_intrinsic fallback
+// (used by every CodeGen_C-derived backend, i.e. every GPU target here) had
+// no case for extract_bits/concat_bits at all -- those were only ever
+// handled by CodeGen_LLVM calling the lowering helpers directly. Both fixed
+// in CodeGen_Metal_Dev.cpp / FindIntrinsics.cpp. Run with
+// HL_JIT_TARGET=host-metal or HL_JIT_TARGET=host-opencl (etc.) to select a
+// backend; skipped if the environment doesn't select a GPU target at all.
+void run_variant(const char *name, float d_offset, int code_mul, MemoryType producer_storage) {
+    Type block_t = Type::Struct({{"d", Float(32)}, {"qs", Int(8), 8}});
+    const int num_blocks = 6;
+
+    Var i("i"), blk("blk"), bo("bo"), bi("bi");
+
+    Func producer("producer");
+    std::vector<Expr> vals;
+    vals.push_back(cast<float>(blk) + d_offset);
+    for (int k = 0; k < 8; k++) {
+        vals.push_back(cast<int8_t>(blk * code_mul + k));
+    }
+    producer(blk) = pack_struct(block_t, vals);
+
+    Func result("result");
+    result(i, blk) = cast<float>(field(producer(blk), "qs")[i]) * cast<float>(field(producer(blk), "d"));
+
+    result.gpu_tile(blk, bo, bi, 1);
+    if (producer_storage == MemoryType::GPUShared) {
+        producer.compute_at(result, bo);
+        producer.store_in(MemoryType::GPUShared);
+    } else {
+        producer.compute_at(result, bi);
+    }
+
+    Buffer<float> out = result.realize({8, num_blocks});
+    out.copy_to_host();
+
+    for (int b = 0; b < num_blocks; b++) {
+        float d = (float)b + d_offset;
+        for (int k = 0; k < 8; k++) {
+            float expected = (float)(int8_t)(b * code_mul + k) * d;
+            if (std::abs(out(k, b) - expected) > 1e-4f) {
+                printf("gpu_struct_type (%s): result(%d, %d) = %f instead of %f\n",
+                       name, k, b, out(k, b), expected);
+                exit(1);
+            }
+        }
+    }
+}
+
+int main(int argc, char **argv) {
+    Target target = get_jit_target_from_environment();
+    if (!target.has_gpu_feature()) {
+        printf("[SKIP] No GPU target enabled.\n");
+        return 0;
+    }
+
+    run_variant("gpu_threads compute_at", 0.5f, 3, MemoryType::Register);
+    run_variant("gpu shared compute_at", 1.5f, 5, MemoryType::GPUShared);
+
+    printf("Success!\n");
+    return 0;
+}
diff --git a/test/correctness/struct_type.cpp b/test/correctness/struct_type.cpp
new file mode 100644
index 000000000000..341b9f51ad5f
--- /dev/null
+++ b/test/correctness/struct_type.cpp
@@ -0,0 +1,763 @@
+#include "Halide.h"
+#include <cmath>
+#include <cstring>
+#include <stdio.h>
+
+using namespace Halide;
+using namespace Halide::Internal;
+
+namespace {
+
+// A block_q4_0-shaped struct: one float16 scale followed by 16 packed byte-wide "codes".
+Type make_block_type() {
+    return Type::Struct({{"d", Float(16)}, {"qs", UInt(8), 16}});
+}
+
+void test_type_struct_basics() {
+    Type block_a = make_block_type();
+    Type block_b = make_block_type();  // independently constructed, same layout
+    Type reordered = Type::Struct({{"qs", UInt(8), 16}, {"d", Float(16)}});
+    Type different_field = Type::Struct({{"d", Float(16)}, {"qs", UInt(8), 8}});
+
+    if (!block_a.is_struct() || !block_b.is_struct()) {
+        printf("Expected Type::Struct() results to report is_struct()\n");
+        exit(1);
+    }
+    if (block_a.bytes() != 18) {
+        printf("Expected block_a.bytes() == 18, got %d\n", block_a.bytes());
+        exit(1);
+    }
+    if (block_a.code() != halide_type_struct || block_a.lanes() != 1) {
+        printf("Expected a struct type to have code halide_type_struct and one lane\n");
+        exit(1);
+    }
+    if (block_a.is_uint() || block_a.is_int() || block_a.is_handle()) {
+        printf("A struct type should not report as a numeric or handle type\n");
+        exit(1);
+    }
+    if (!(block_a == block_b)) {
+        printf("Two independently-declared-but-identical struct types should compare equal\n");
+        exit(1);
+    }
+    if (block_a == reordered) {
+        printf("Struct types with differently-ordered fields should not compare equal\n");
+        exit(1);
+    }
+    if (block_a == different_field) {
+        printf("Struct types with different field shapes should not compare equal\n");
+        exit(1);
+    }
+    if (block_a == UInt(8)) {
+        printf("A struct type should not compare equal to a plain UInt(8)\n");
+        exit(1);
+    }
+    // A struct's type code carries its struct-ness, so with_bits/with_lanes
+    // (which never change the code) always preserve it -- bits() and lanes()
+    // aren't meaningful for a struct, and a struct-typed Load/Store must not
+    // silently degrade into a plain integer partway through lowering (e.g.
+    // StorageFlattening rounding bits up to a byte multiple). Only with_code
+    // to a different code drops struct-ness.
+    if (!block_a.with_bits(8).is_struct() || !block_a.with_bits(16).is_struct()) {
+        printf("Type::with_bits should preserve struct-ness\n");
+        exit(1);
+    }
+    if (block_a.with_bits(16).bytes() != 18) {
+        printf("with_bits should not disturb a struct's real byte size\n");
+        exit(1);
+    }
+    if (!block_a.with_lanes(1).is_struct()) {
+        printf("Type::with_lanes(same lanes) should preserve struct-ness\n");
+        exit(1);
+    }
+    if (block_a.with_code(Type::UInt).is_struct()) {
+        printf("Type::with_code(non-struct code) should drop struct-ness\n");
+        exit(1);
+    }
+}
+
+// operator< must agree with operator==: two independently-allocated but
+// structurally-identical struct types (block_a/block_b, from two separate
+// Type::Struct() calls) carry different StructTypeInfo pointers, so a
+// pointer-based operator< would report them as ordered relative to each
+// other despite comparing equal, breaking the usual invariant that
+// !(a < b) && !(b < a) iff a == b (which std::map/std::set/std::sort all
+// rely on).
+void test_struct_type_ordering_consistent_with_equality() {
+    Type block_a = make_block_type();
+    Type block_b = make_block_type();  // independently constructed, same layout
+    Type reordered = Type::Struct({{"qs", UInt(8), 16}, {"d", Float(16)}});
+
+    if (block_a < block_b || block_b < block_a) {
+        printf("Equal struct types must not be ordered relative to each other\n");
+        exit(1);
+    }
+    if (!(block_a < reordered) && !(reordered < block_a)) {
+        printf("Unequal struct types must be ordered one way or the other\n");
+        exit(1);
+    }
+    if (block_a < reordered && reordered < block_a) {
+        printf("Struct type ordering must be antisymmetric\n");
+        exit(1);
+    }
+    if (UInt(8) < block_a && block_a < UInt(8)) {
+        printf("Struct-vs-non-struct ordering must be antisymmetric\n");
+        exit(1);
+    }
+}
+
+// constant_integer_bounds() (used e.g. for overflow checks ahead of a cast)
+// must treat a struct-typed Expr as having no known integer bounds -- it
+// isn't a number, so there's nothing to bound. bounds_helper's generic
+// fallback (for a node shape it doesn't specifically recognize, e.g. a
+// struct-typed Select) computes ConstantInterval::bounds_of_type(e.type())
+// and asserts the result is representable by e.type(); can_represent()
+// rejects a struct outright, so bounds_of_type() must agree and report "no
+// known bounds" too. A Select (rather than the struct_pack Expr itself) is
+// used here because bounds_helper has no case for Select at all, so it's
+// guaranteed to hit that generic fallback rather than Call's own (more
+// specific) dispatch. Built via a raw struct_pack Call rather than the
+// pack_struct() front-end helper, since that helper validates one initializer
+// per declared field (with arrays swept via gather()/`_`) -- more machinery
+// than this Type/Expr-level test needs.
+void test_constant_integer_bounds_of_struct() {
+    Type block_t = Type::Struct({{"d", Float(32)}, {"qs", UInt(8), 4}});
+    Expr a = Call::make(block_t, Call::struct_pack,
+                        {make_zero(Float(32)), make_zero(UInt(8)), make_zero(UInt(8)), make_zero(UInt(8)), make_zero(UInt(8))},
+                        Call::PureIntrinsic);
+    Expr cond = Variable::make(Bool(), "c");
+    Expr sel = Select::make(cond, a, a);
+
+    ConstantInterval ci = constant_integer_bounds(sel);
+    if (ci.min_defined || ci.max_defined) {
+        printf("constant_integer_bounds() of a struct-typed Expr should be unbounded, got "
+               "min_defined=%d max_defined=%d\n",
+               ci.min_defined, ci.max_defined);
+        exit(1);
+    }
+
+    if (block_t.can_represent((int64_t)0)) {
+        printf("A struct type should not claim to represent any integer constant\n");
+        exit(1);
+    }
+}
+
+// Reading struct-typed data through an ImageParam/hand-built Buffer. A
+// struct-typed ImageParam/Func keeps *exactly* the dimensionality the user
+// declared -- no hidden extra dimension anywhere -- so its buffer shape is
+// built exactly like any other typed buffer's would be, with strides in
+// units of "one whole struct" (Type::bytes() is what scales that up to an
+// actual byte address, entirely inside field()'s lowering).
+void test_read_from_buffer() {
+    Type block_t = make_block_type();
+    const int num_blocks = 5, num_rows = 3;
+
+    int block_bytes = block_t.bytes();
+    std::vector<uint8_t> data((size_t)block_bytes * num_blocks * num_rows);
+    auto at = [&](int byte, int blk, int row) -> uint8_t & {
+        return data[(size_t)byte + blk * block_bytes + row * block_bytes * num_blocks];
+    };
+    for (int row = 0; row < num_rows; row++) {
+        for (int blk = 0; blk < num_blocks; blk++) {
+            float16_t d_val(1.5f + 0.25f * blk + 0.1f * row);
+            uint16_t bits = d_val.to_bits();
+            memcpy(&at(0, blk, row), &bits, 2);
+            for (int k = 0; k < 16; k++) {
+                at(2 + k, blk, row) = (uint8_t)((blk * 3 + k + row) % 251);
+            }
+        }
+    }
+    // Exactly 2 dims (blk, row), matching Wt's own declared dimensionality --
+    // strides are in units of "one struct" (dense AoS: blk's stride is 1,
+    // row's stride is num_blocks), the same convention any other typed buffer
+    // uses. The buffer's element type is the struct itself, so its element size
+    // is block_bytes and the strides scale up to the byte layout above.
+    halide_dimension_t shape[2] = {
+        {0, num_blocks, 1, 0},
+        {0, num_rows, num_blocks, 0},
+    };
+    Buffer<> Wt_buf(block_t, data.data(), 2, shape);
+
+    ImageParam Wt(block_t, 2, "Wt");
+    Wt.set(Wt_buf);
+
+    Var k("k"), blk("blk"), row("row");
+
+    Func delta("delta");
+    delta(blk, row) = cast<float>(field(Wt(blk, row), "d"));
+
+    Func code("code");
+    code(k, blk, row) = cast<int32_t>(field(Wt(blk, row), "qs")[k]);
+
+    Buffer<float> delta_result = delta.realize({num_blocks, num_rows});
+    Buffer<int32_t> code_result = code.realize({16, num_blocks, num_rows});
+
+    for (int row = 0; row < num_rows; row++) {
+        for (int blk = 0; blk < num_blocks; blk++) {
+            float16_t expected_d(1.5f + 0.25f * blk + 0.1f * row);
+            float expected = (float)expected_d;
+            if (delta_result(blk, row) != expected) {
+                printf("delta(%d, %d) = %f instead of %f\n", blk, row, delta_result(blk, row), expected);
+                exit(1);
+            }
+            for (int k = 0; k < 16; k++) {
+                int32_t expected_code = (blk * 3 + k + row) % 251;
+                if (code_result(k, blk, row) != expected_code) {
+                    printf("code(%d, %d, %d) = %d instead of %d\n",
+                           k, blk, row, code_result(k, blk, row), expected_code);
+                    exit(1);
+                }
+            }
+        }
+    }
+}
+
+// A struct-valued Func that IS the pipeline's own output, realized directly.
+// Struct types are treated as an ordinary opaque element type all the way
+// through scheduling/bounds_inference/storage_flattening -- producer keeps
+// its declared dimensionality (1, just `blk`) throughout, and Pipeline's own
+// realize() needs no struct-specific handling: it allocates a struct-typed,
+// 1-dimensional output buffer exactly like it would for any other type.
+void test_realized_writer_as_output() {
+    Type block_t = make_block_type();
+    Var blk("blk");
+    const int num_blocks = 6;
+
+    Func producer("producer");
+    // Explicit gather() generator form for the array field (extent inferred
+    // from the struct field), alongside a scalar field initializer.
+    producer(blk) = pack_struct(
+        block_t,
+        {
+            cast<float16_t>(0.5f + cast<float>(blk)),                        // d
+            gather(16, [&](Expr k) { return cast<uint8_t>(blk * 5 + k); }),  // qs
+        });
+
+    Buffer<> raw = producer.realize({num_blocks});
+    // The realized output buffer reports a struct element type of the right byte
+    // size, but not its field layout -- Halide::Buffer<>::type() round-trips
+    // through the ABI-level halide_type_t, which carries a struct's packed size
+    // but forgets its fields (the same limitation Handle-typed buffers have for
+    // their pointee type). Dimensionality is exactly what producer declared
+    // (1, just `blk`) -- no hidden extra dimension.
+    if (!raw.type().is_struct() || raw.type().bytes() != block_t.bytes()) {
+        printf("Realized struct output should report a struct element type of %d bytes, got %d-byte %s\n",
+               block_t.bytes(), raw.type().bytes(), raw.type().is_struct() ? "struct" : "non-struct");
+        exit(1);
+    }
+    if (raw.dimensions() != 1 || raw.dim(0).extent() != num_blocks) {
+        printf("Unexpected realized struct output shape: %d dims, dim0 extent %d\n",
+               raw.dimensions(), raw.dim(0).extent());
+        exit(1);
+    }
+
+    // Read it back through a fresh struct-typed ImageParam wrapping the same
+    // buffer directly -- no coercion needed, both sides agree on both type
+    // and dimensionality.
+    ImageParam Wt(block_t, 1, "Wt2");
+    Wt.set(raw);
+    Var blk2("blk2");
+    Func delta("delta");
+    delta(blk2) = cast<float>(field(Wt(blk2), "d"));
+    Func code0("code0");
+    code0(blk2) = cast<int32_t>(field(Wt(blk2), "qs")[0]);
+
+    Buffer<float> delta_result = delta.realize({num_blocks});
+    Buffer<int32_t> code0_result = code0.realize({num_blocks});
+    for (int blk_i = 0; blk_i < num_blocks; blk_i++) {
+        float16_t expected_d(0.5f + (float)blk_i);
+        if (delta_result(blk_i) != (float)expected_d) {
+            printf("realized writer: delta(%d) = %f instead of %f\n",
+                   blk_i, delta_result(blk_i), (float)expected_d);
+            exit(1);
+        }
+        int32_t expected_code0 = (uint8_t)(blk_i * 5);
+        if (code0_result(blk_i) != expected_code0) {
+            printf("realized writer: code0(%d) = %d instead of %d\n",
+                   blk_i, code0_result(blk_i), expected_code0);
+            exit(1);
+        }
+    }
+}
+
+// One initializer per field of make_block_type(): a scalar `d`, and the array
+// field `qs` filled by sweeping the `_` placeholder over its extent -- note the
+// index arithmetic (blk * code_mul + _), which the bare placeholder form
+// supports because the extent is a compile-time constant.
+std::vector<StructFieldInit> block_field_values(const Var &blk, float d_offset, int code_mul) {
+    return {
+        cast<float16_t>(d_offset + cast<float>(blk)),  // d
+        cast<uint8_t>(blk * code_mul + _),             // qs[16], swept over [0, 16)
+    };
+}
+
+float expected_d(float d_offset, int blk) {
+    return (float)float16_t(d_offset + (float)blk);
+}
+
+// Torture test: a struct-valued Func consumed *directly* by another Func in
+// the same pipeline (not realized standalone, not read back via ImageParam).
+// This is exactly the case that broke bounds_inference under the old
+// "desugar into a hidden extra dimension" architecture -- producer here
+// keeps its true declared dimensionality (1) throughout, so there's nothing
+// unusual for bounds_inference to trip on.
+void test_struct_consumed_by_another_func() {
+    Type block_t = make_block_type();
+    Var blk("blk");
+    const int num_blocks = 6;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 0.5f, 5));
+    producer.compute_root();
+
+    Func consumer("consumer");
+    consumer(blk) = cast<float>(field(producer(blk), "d"));
+
+    Buffer<float> result = consumer.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = expected_d(0.5f, i);
+        if (result(i) != expected) {
+            printf("struct_consumed_by_another_func: consumer(%d) = %f instead of %f\n", i, result(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: compute_at a struct-valued producer inside a consumer's
+// loop, rather than compute_root or fully inlining it.
+void test_compute_at() {
+    Type block_t = make_block_type();
+    Var blk("blk"), blk_o("blk_o"), blk_i("blk_i");
+    const int num_blocks = 8;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 1.0f, 1));
+
+    Func consumer("consumer");
+    consumer(blk) = cast<float>(field(producer(blk), "d")) +
+                    cast<float>(field(producer(blk), "qs")[0]);
+    consumer.split(blk, blk_o, blk_i, 4);
+    producer.compute_at(consumer, blk_o);
+
+    Buffer<float> result = consumer.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = expected_d(1.0f, i) + (float)(uint8_t)(i * 1);
+        if (std::abs(result(i) - expected) > 1e-3f) {
+            printf("compute_at: consumer(%d) = %f instead of %f\n", i, result(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: one struct-valued intermediate consumed by two different
+// downstream Funcs, reading different fields with different schedules
+// (scalar vs. vectorized). Confirms both share the same producer bounds
+// without either schedule interfering with the other's view of the struct.
+void test_two_consumers_different_schedules() {
+    Type block_t = make_block_type();
+    Var blk("blk"), k("k");
+    const int num_blocks = 8;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 2.0f, 7));
+    producer.compute_root();
+
+    Func delta("delta");
+    delta(blk) = cast<float>(field(producer(blk), "d"));
+
+    Func codes("codes");
+    codes(k, blk) = cast<int32_t>(field(producer(blk), "qs")[k]);
+    codes.bound(k, 0, 16).vectorize(k, 16);
+
+    Buffer<float> delta_result = delta.realize({num_blocks});
+    Buffer<int32_t> codes_result = codes.realize({16, num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = expected_d(2.0f, i);
+        if (delta_result(i) != expected) {
+            printf("two_consumers: delta(%d) = %f instead of %f\n", i, delta_result(i), expected);
+            exit(1);
+        }
+        for (int k_i = 0; k_i < 16; k_i++) {
+            int32_t expected_code = (uint8_t)(i * 7 + k_i);
+            if (codes_result(k_i, i) != expected_code) {
+                printf("two_consumers: codes(%d, %d) = %d instead of %d\n",
+                       k_i, i, codes_result(k_i, i), expected_code);
+                exit(1);
+            }
+        }
+    }
+}
+
+// Torture test: a struct-valued Func's own value is a bare passthrough Call
+// to another struct-valued Func (e.g. via Func::in()), not a literal
+// pack_struct() -- this is what a scheduling wrapper looks like in
+// practice, and it must be treated as a valid struct-typed Store value (a
+// whole-struct copy split field-by-field) rather than only recognizing a
+// literal pack_struct().
+void test_passthrough_wrapper() {
+    Type block_t = make_block_type();
+    Var blk("blk");
+    const int num_blocks = 5;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 3.0f, 3));
+    producer.compute_root();
+
+    Func wrapper = producer.in();
+    wrapper.compute_root();
+
+    Func consumer("consumer");
+    consumer(blk) = cast<float>(field(wrapper(blk), "d")) +
+                    cast<float>(field(wrapper(blk), "qs")[2]);
+
+    Buffer<float> result = consumer.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = expected_d(3.0f, i) + (float)(uint8_t)(i * 3 + 2);
+        if (std::abs(result(i) - expected) > 1e-3f) {
+            printf("passthrough_wrapper: consumer(%d) = %f instead of %f\n", i, result(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: a struct-valued Func's value is a select() between two
+// struct-typed Exprs, and the Func is compute_root'd -- so the Store
+// splitter must push field projection through the Select, not just
+// recognize a literal pack_struct() directly as the whole value.
+void test_select_between_structs() {
+    Type block_t = make_block_type();
+    Var blk("blk");
+    const int num_blocks = 6;
+
+    Func producer("producer");
+    Expr cond = (blk % 2) == 0;
+    producer(blk) = select(cond,
+                           pack_struct(block_t, block_field_values(blk, 10.0f, 1)),
+                           pack_struct(block_t, block_field_values(blk, 20.0f, 1)));
+    producer.compute_root();
+
+    Func consumer("consumer");
+    consumer(blk) = cast<float>(field(producer(blk), "d"));
+
+    Buffer<float> result = consumer.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = (i % 2 == 0) ? expected_d(10.0f, i) : expected_d(20.0f, i);
+        if (result(i) != expected) {
+            printf("select_between_structs: consumer(%d) = %f instead of %f\n", i, result(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: a struct-typed sub-expression used twice within one Func's
+// value gets hoisted into a Let by common_subexpression_elimination.
+// field()'s lowering must resolve a struct-typed Let-bound Variable back to
+// the real (buffer-backed, in this case) value rather than erroring out.
+void test_cse_shared_struct_let() {
+    Type block_t = make_block_type();
+    Var blk("blk");
+    const int num_blocks = 6;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 1.0f, 7));
+    producer.compute_root();
+
+    Func result("result");
+    Expr blk_struct = producer(blk);
+    result(blk) = cast<float>(field(blk_struct, "d")) + cast<float>(field(blk_struct, "qs")[0]);
+
+    Buffer<float> r = result.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = expected_d(1.0f, i) + (float)(uint8_t)(i * 7);
+        if (std::abs(r(i) - expected) > 1e-3f) {
+            printf("cse_shared_struct_let: result(%d) = %f instead of %f\n", i, r(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: an array field read with a *runtime*, vectorized index --
+// must lower to ordinary byte-addressed loads that the normal vectorizer
+// (which runs well after struct lowering) can still turn into wide loads,
+// not a per-lane select-chain or scalar loop.
+void test_vectorized_array_index() {
+    Type block_t = make_block_type();
+    Var blk("blk"), k("k");
+    const int num_blocks = 4;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 4.0f, 9));
+    producer.compute_root();
+
+    Func codes("codes");
+    codes(k, blk) = cast<int32_t>(field(producer(blk), "qs")[k]);
+    codes.bound(k, 0, 16).vectorize(k, 16);
+
+    Buffer<int32_t> result = codes.realize({16, num_blocks});
+    for (int blk_i = 0; blk_i < num_blocks; blk_i++) {
+        for (int k_i = 0; k_i < 16; k_i++) {
+            int32_t expected = (uint8_t)(blk_i * 9 + k_i);
+            if (result(k_i, blk_i) != expected) {
+                printf("vectorized_array_index: codes(%d, %d) = %d instead of %d\n",
+                       k_i, blk_i, result(k_i, blk_i), expected);
+                exit(1);
+            }
+        }
+    }
+}
+
+// Torture test: a struct field whose own type is itself Type::Struct (a
+// sub-struct field), read through two chained front-end field() calls. Only
+// supported for the fully-inlined case (see test/error/struct_nested_field_store.cpp
+// and test/error/struct_nested_field_load.cpp for the unsupported,
+// buffer-backed cases) -- this test exercises the supported, inlined path.
+void test_nested_struct() {
+    Type inner_t = Type::Struct({{"a", Int(32)}, {"b", UInt(8)}});
+    Type outer_t = Type::Struct({{"inner", inner_t}, {"c", Float(32)}});
+    Var i("i");
+    const int n = 4;
+
+    Func producer("producer");
+    producer(i) = pack_struct(outer_t, {pack_struct(inner_t, {cast<int32_t>(i * 10), cast<uint8_t>(i)}), cast<float>(i) + 0.5f});
+    // Left at the default (inlined) schedule: this is the supported path.
+
+    Func consumer("consumer");
+    Expr outer_val = producer(i);
+    Expr inner_val = field(outer_val, "inner");
+    consumer(i) = field(inner_val, "a") + cast<int32_t>(field(inner_val, "b")) + cast<int32_t>(field(outer_val, "c"));
+
+    Buffer<int32_t> result = consumer.realize({n});
+    for (int idx = 0; idx < n; idx++) {
+        int32_t expected = idx * 10 + idx + (int32_t)(idx + 0.5f);
+        if (result(idx) != expected) {
+            printf("nested_struct: consumer(%d) = %d instead of %d\n", idx, result(idx), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: a struct-typed sub-expression that is itself a nested struct
+// field (not a top-level struct value), used twice within its enclosing
+// pack_struct's own definition -- forces common_subexpression_elimination to
+// hoist it into a struct-typed Let whose *body* refers to it as a bare
+// Variable directly inside another pack_struct's argument list (not through
+// field()). This exercises LowerStructTypesMutator's generic struct-typed
+// Variable resolution path (visit(Variable)), not just project_field's own
+// internal Let/Variable handling.
+void test_nested_struct_cse_shared_inner() {
+    Type inner_t = Type::Struct({{"a", Int(32)}, {"b", UInt(8)}});
+    Type outer_t = Type::Struct({{"inner1", inner_t}, {"inner2", inner_t}});
+    Var i("i");
+    const int n = 5;
+
+    Func producer("producer");
+    Expr shared_inner = pack_struct(inner_t, {cast<int32_t>(i * 10), cast<uint8_t>(i)});
+    producer(i) = pack_struct(outer_t, {shared_inner, shared_inner});
+
+    Func consumer("consumer");
+    Expr outer_val = producer(i);
+    Expr inner1 = field(outer_val, "inner1");
+    Expr inner2 = field(outer_val, "inner2");
+    consumer(i) = field(inner1, "a") + cast<int32_t>(field(inner1, "b")) +
+                  field(inner2, "a") + cast<int32_t>(field(inner2, "b"));
+
+    Buffer<int32_t> result = consumer.realize({n});
+    for (int idx = 0; idx < n; idx++) {
+        int32_t expected = 2 * (idx * 10 + idx);
+        if (result(idx) != expected) {
+            printf("nested_struct_cse_shared_inner: consumer(%d) = %d instead of %d\n", idx, result(idx), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: an ImageParam over a non-dense (padded/strided) layout --
+// struct field addressing must compose with an arbitrary declared stride,
+// not just the dense case exercised by test_read_from_buffer.
+void test_non_dense_stride() {
+    Type block_t = make_block_type();
+    const int num_blocks = 4, row_stride_in_blocks = 6;  // 2 extra padding blocks per row
+    int block_bytes = block_t.bytes();
+
+    std::vector<uint8_t> data((size_t)block_bytes * row_stride_in_blocks);
+    auto at = [&](int byte, int blk) -> uint8_t & { return data[(size_t)byte + blk * block_bytes]; };
+    for (int blk = 0; blk < num_blocks; blk++) {
+        float16_t d_val(7.0f + blk);
+        uint16_t bits = d_val.to_bits();
+        memcpy(&at(0, blk), &bits, 2);
+    }
+    halide_dimension_t shape[1] = {
+        {0, num_blocks, 1, 0},
+    };
+    Buffer<> Wt_buf(block_t, data.data(), 1, shape);
+
+    ImageParam Wt(block_t, 1, "WtPadded");
+    Wt.set(Wt_buf);
+
+    Var blk("blk");
+    Func delta("delta");
+    delta(blk) = cast<float>(field(Wt(blk), "d"));
+
+    Buffer<float> result = delta.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = (float)float16_t(7.0f + i);
+        if (result(i) != expected) {
+            printf("non_dense_stride: delta(%d) = %f instead of %f\n", i, result(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: a struct-valued Func left at its default (compute_inline)
+// schedule and consumed by another Func -- must still fold away to nothing
+// materialized, matching the doc's "no bytes ever touched for an inlined
+// struct producer" claim.
+void test_inlined_default_schedule() {
+    Type block_t = make_block_type();
+    Var blk("blk");
+    const int num_blocks = 6;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 6.0f, 2));
+    // No explicit schedule -- defaults to compute_inline().
+
+    Func consumer("consumer");
+    consumer(blk) = cast<float>(field(producer(blk), "d"));
+
+    Buffer<float> result = consumer.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = expected_d(6.0f, i);
+        if (result(i) != expected) {
+            printf("inlined_default_schedule: consumer(%d) = %f instead of %f\n", i, result(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: .parallel() on a compute_root'd struct producer's logical
+// dim. Each parallel iteration splits its own point's struct value into N
+// field-stores at disjoint byte ranges, so there's no cross-iteration
+// hazard from the splitting itself.
+void test_parallel_schedule() {
+    Type block_t = make_block_type();
+    Var blk("blk");
+    const int num_blocks = 32;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 8.0f, 11));
+    producer.compute_root().parallel(blk);
+
+    Func consumer("consumer");
+    consumer(blk) = cast<float>(field(producer(blk), "d")) +
+                    cast<float>(field(producer(blk), "qs")[15]);
+
+    Buffer<float> result = consumer.realize({num_blocks});
+    for (int i = 0; i < num_blocks; i++) {
+        float expected = expected_d(8.0f, i) + (float)(uint8_t)(i * 11 + 15);
+        if (std::abs(result(i) - expected) > 1e-3f) {
+            printf("parallel_schedule: consumer(%d) = %f instead of %f\n", i, result(i), expected);
+            exit(1);
+        }
+    }
+}
+
+// Torture test: an array field read with a *runtime* index directly off of
+// an inlined (never materialized) struct producer, so field()'s lowering
+// sees a literal pack_struct() rather than a Load -- this exercises
+// fold_field_of_pack's runtime-index select-chain fallback (as opposed to
+// its constant-index fast path, which every other array-field test here
+// takes because their producer is compute_root'd/compute_at'd first).
+void test_runtime_index_into_inlined_pack() {
+    Type block_t = make_block_type();
+    Var blk("blk"), k("k");
+    const int num_blocks = 3;
+
+    Func producer("producer");
+    producer(blk) = pack_struct(block_t, block_field_values(blk, 0.0f, 4));
+    // No explicit schedule -- stays compute_inline(), so field() below sees
+    // the literal pack_struct() directly instead of a buffer Load.
+
+    Func codes("codes");
+    codes(k, blk) = cast<int32_t>(field(producer(blk), "qs")[k]);
+    // k is a runtime (non-constant) index into the array field.
+
+    Buffer<int32_t> result = codes.realize({16, num_blocks});
+    for (int blk_i = 0; blk_i < num_blocks; blk_i++) {
+        for (int k_i = 0; k_i < 16; k_i++) {
+            int32_t expected = (uint8_t)(blk_i * 4 + k_i);
+            if (result(k_i, blk_i) != expected) {
+                printf("runtime_index_into_inlined_pack: codes(%d, %d) = %d instead of %d\n",
+                       k_i, blk_i, result(k_i, blk_i), expected);
+                exit(1);
+            }
+        }
+    }
+}
+
+// pack_struct can initialize an array field by copying a whole same-typed
+// array field out of another struct value, via field(other, "name"). Here we
+// re-pack a struct keeping its `qs` array verbatim while replacing `d`.
+void test_pack_struct_field_copy() {
+    Type block_t = make_block_type();
+    Var blk("blk"), k("k");
+    const int num_blocks = 5;
+
+    Func src("src");
+    src(blk) = pack_struct(block_t, block_field_values(blk, 2.0f, 3));
+    src.compute_root();  // force a materialized struct buffer, so the copy reads a real field
+
+    Func repacked("repacked");
+    Expr s = src(blk);
+    repacked(blk) = pack_struct(
+        block_t,
+        {
+            cast<float16_t>(100.0f + cast<float>(blk)),  // new d
+            field(s, "qs"),                              // copy qs[16] from src
+        });
+
+    Func out_d("out_d"), out_q("out_q");
+    out_d(blk) = cast<float>(field(repacked(blk), "d"));
+    out_q(k, blk) = cast<int32_t>(field(repacked(blk), "qs")[k]);
+
+    Buffer<float> d_res = out_d.realize({num_blocks});
+    Buffer<int32_t> q_res = out_q.realize({16, num_blocks});
+    for (int blk_i = 0; blk_i < num_blocks; blk_i++) {
+        float exp_d = (float)float16_t(100.0f + (float)blk_i);
+        if (d_res(blk_i) != exp_d) {
+            printf("field_copy: d(%d) = %f instead of %f\n", blk_i, d_res(blk_i), exp_d);
+            exit(1);
+        }
+        for (int k_i = 0; k_i < 16; k_i++) {
+            int32_t exp_q = (uint8_t)(blk_i * 3 + k_i);  // matches block_field_values(blk, 2.0f, 3)
+            if (q_res(k_i, blk_i) != exp_q) {
+                printf("field_copy: qs(%d, %d) = %d instead of %d\n", k_i, blk_i, q_res(k_i, blk_i), exp_q);
+                exit(1);
+            }
+        }
+    }
+}
+
+}  // namespace
+
+int main(int argc, char **argv) {
+    test_type_struct_basics();
+    test_struct_type_ordering_consistent_with_equality();
+    test_constant_integer_bounds_of_struct();
+    test_read_from_buffer();
+    test_realized_writer_as_output();
+    test_struct_consumed_by_another_func();
+    test_compute_at();
+    test_two_consumers_different_schedules();
+    test_passthrough_wrapper();
+    test_select_between_structs();
+    test_cse_shared_struct_let();
+    test_vectorized_array_index();
+    test_nested_struct();
+    test_nested_struct_cse_shared_inner();
+    test_non_dense_stride();
+    test_inlined_default_schedule();
+    test_parallel_schedule();
+    test_runtime_index_into_inlined_pack();
+    test_pack_struct_field_copy();
+    printf("Success!\n");
+    return 0;
+}
diff --git a/test/correctness/struct_type_dot_product.cpp b/test/correctness/struct_type_dot_product.cpp
new file mode 100644
index 000000000000..63067a80bb3d
--- /dev/null
+++ b/test/correctness/struct_type_dot_product.cpp
@@ -0,0 +1,295 @@
+#include "Halide.h"
+#include "halide_test_dirs.h"
+
+#include <algorithm>
+#include <cmath>
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <fstream>
+#include <new>
+#include <random>
+#include <sstream>
+
+using namespace Halide;
+
+namespace {
+
+// The block layouts of ggml's q4_0 (4-bit weights) and q8_0 (8-bit
+// activations) quantization formats, expressed as Halide struct types.
+//
+//   block_q4_0 { fp16 d; uint8 qs[16]; }   // 32 quants, 2 nibbles per byte
+//   block_q8_0 { fp16 d; int8  qs[32]; }   // 32 quants
+//
+// This mirrors ggml_vec_dot_q4_0_q8_0 (ggml-cpu/arch/arm/quants.c): for each
+// pair of blocks, the 16 q4_0 bytes hold 32 signed 4-bit weights (low nibbles
+// map to quants 0..15, high nibbles to 16..31, each biased by -8), dotted with
+// the 32 int8 q8_0 activations, scaled by the product of the two fp16 deltas,
+// and summed across all blocks.
+//
+// The *algorithm* is written as the obvious thing: dequantize each side to
+// float on its own -- x_wt for the 4-bit weights, y_wt for the int8
+// activations, each scaling by its block's fp16 delta -- and sum the products
+// over all quants. result's update is then just
+//   result() += x_wt(k, b) * y_wt(k, b)
+// a plain float dot product over a single quant index k in [0, 32), with none
+// of the packed-integer kernel structure spelled out.
+//
+// The efficient kernel is then *derived by the schedule*, composing the four
+// quantization directives plus a split of the quant reduction:
+//
+//   * split(r_quant, ko, ki, 16) cuts the 32 quants into the two nibble halves
+//     ko and the packed byte ki. Because x_wt indexes qs by k % 16 and picks
+//     its nibble by k < 16, unrolling ko folds those to a constant nibble and a
+//     dense byte index, so both qs arrays load as whole vectors over ki rather
+//     than a strided gather.
+//   * rfactor(r_block, u) splits the (block, quant) reduction so each block's
+//     quants reduce into their own partial sum, partial(u).
+//   * eager_inline(x_wt, y_wt) substitutes the two dequantizers into partial's
+//     update, surfacing the per-block deltas xd(u), yd(u) and the integer
+//     weight/activation terms as explicit leaves hoist_invariants() can see.
+//   * hoist_invariants() lifts xd(u) and yd(u) -- invariant across the quants
+//     of a block -- out of the reduction, leaving a scale-free integer dot
+//     dot(u) and a single scale multiply in the write-back.
+//   * change_type(Int(32)) proves the per-block term can't overflow Int(32) and
+//     retypes dot's accumulation accordingly, so the quants reduce with
+//     smull/addv instead of a float round-trip.
+//
+// x_wt and y_wt are left unscheduled: eager_inline() has already folded them
+// into the reduction, so nothing else computes them.
+Type q4_0_type() {
+    return Type::Struct({{"d", Float(16)}, {"qs", UInt(8), 16}});
+}
+Type q8_0_type() {
+    return Type::Struct({{"d", Float(16)}, {"qs", Int(8), 32}});
+}
+Type q5_0_type() {
+    // Deliberately packed: qh is a four-byte scalar at the unaligned offset 2.
+    return Type::Struct({{"d", Float(16)}, {"qh", UInt(32)}, {"qs", UInt(8), 16}});
+}
+
+struct Pipeline {
+    ImageParam x{q4_0_type(), 1, "x"};  // nb q4_0 weight blocks
+    ImageParam y{q8_0_type(), 1, "y"};  // nb q8_0 activation blocks
+    Func result{"qdot"};
+
+    Pipeline() {
+        // Alignment hints: dense, base-aligned blocks.
+        x.dim(0).set_min(0);
+        y.dim(0).set_min(0);
+        x.set_host_alignment(16);
+        y.set_host_alignment(16);
+
+        Var b("b"), k("k"), u("u");
+
+        // Dequantize one 4-bit weight to float: quant k is a nibble of the
+        // packed byte k % 16 -- its low nibble for k < 16, its high nibble
+        // otherwise -- biased by -8 to signed [-8, 7], times the block delta.
+        Expr nib = field(x(b), "qs")[k % 16];  // uint8
+        Expr w = cast<int32_t>(select(k < 16, nib % 16, nib / 16)) - 8;
+        Func x_wt("x_wt");
+        x_wt(k, b) = cast<float>(field(x(b), "d")) * cast<float>(w);
+
+        // Dequantize one int8 activation to float: quant k is int8 k, times the
+        // block delta.
+        Expr a = cast<int32_t>(field(y(b), "qs")[k]);
+        Func y_wt("y_wt");
+        y_wt(k, b) = cast<float>(field(y(b), "d")) * cast<float>(a);
+
+        // The algorithm: dequantize each side and sum every product.
+        RDom r(0, x.dim(0).extent(), 0, 32, "r");  // block, quant
+        result() = 0.0f;
+        result() += x_wt(r[1], r[0]) * y_wt(r[1], r[0]);
+
+        // Derive the packed-integer kernel from the schedule. Split the 32
+        // quants into the two nibble halves ko and the packed byte ki first, so
+        // the derived stages carry that structure through.
+        RVar ko("ko"), ki("ki");
+        result.update().split(r[1], ko, ki, 16);
+        Func partial = result.update().rfactor(r[0], u);
+        Func dot = partial.update().eager_inline(x_wt, y_wt).hoist_invariants();
+        Func dot_i32 = dot.change_type(Int(32));
+
+        // One block's scaled contribution per outer step; its quants reduce
+        // horizontally, vectorized across the packed byte and unrolled over the
+        // two nibble halves.
+        partial.compute_at(result, r[0]);
+        dot_i32.compute_at(partial, u)
+            .update()
+            .reorder(ki, ko, u)
+            .atomic()
+            .vectorize(ki, 16)
+            .unroll(ko);
+    }
+};
+
+Target arm_target() {
+    return Target{"arm-64-linux-no_asserts-no_runtime-no_bounds_query-arm_dot_prod-arm_fp16"};
+}
+
+std::string read_file(const std::string &path) {
+    std::ifstream f(path);
+    std::ostringstream ss;
+    ss << f.rdbuf();
+    return ss.str();
+}
+
+int count_occurrences(const std::string &haystack, const std::string &needle) {
+    int n = 0;
+    for (size_t pos = haystack.find(needle); pos != std::string::npos; pos = haystack.find(needle, pos + 1)) {
+        n++;
+    }
+    return n;
+}
+
+// A heap allocation aligned to 16 bytes, since the pipeline declares a
+// 16-byte host alignment on its inputs and plain std::vector's allocator
+// does not guarantee that (only happens to on some 64-bit ABIs).
+struct AlignedBytes {
+    explicit AlignedBytes(size_t n)
+        : n(n), ptr(static_cast<uint8_t *>(::operator new(n, std::align_val_t(16)))) {
+    }
+    ~AlignedBytes() {
+        ::operator delete(ptr, std::align_val_t(16));
+    }
+    AlignedBytes(const AlignedBytes &) = delete;
+    AlignedBytes &operator=(const AlignedBytes &) = delete;
+
+    uint8_t *data() {
+        return ptr;
+    }
+    const uint8_t *data() const {
+        return ptr;
+    }
+    uint8_t &operator[](size_t i) {
+        return ptr[i];
+    }
+    size_t size() const {
+        return n;
+    }
+
+private:
+    size_t n;
+    uint8_t *ptr;
+};
+
+// ---- Reference implementation (scalar, for correctness checking) ----
+float reference_dot(const AlignedBytes &x_blocks,
+                    const AlignedBytes &y_blocks, int nb) {
+    // x block = 18 bytes {fp16 d, u8 qs[16]}; y block = 34 bytes {fp16 d, i8 qs[32]}.
+    float sumf = 0.0f;
+    for (int i = 0; i < nb; i++) {
+        const uint8_t *xb = x_blocks.data() + i * 18;
+        const uint8_t *yb = y_blocks.data() + i * 34;
+        float16_t xd = float16_t::make_from_bits(xb[0] | (xb[1] << 8));
+        float16_t yd = float16_t::make_from_bits(yb[0] | (yb[1] << 8));
+        int isum = 0;
+        for (int j = 0; j < 16; j++) {
+            int nib = xb[2 + j];
+            int x_lo = (nib % 16) - 8;
+            int x_hi = (nib / 16) - 8;
+            int y_lo = (int8_t)yb[2 + j];
+            int y_hi = (int8_t)yb[2 + j + 16];
+            isum += x_lo * y_lo + x_hi * y_hi;
+        }
+        sumf += (float)xd * (float)yd * (float)isum;
+    }
+    return sumf;
+}
+
+void test_correctness() {
+    const int nb = 7;
+    std::mt19937 rng(12345);
+    AlignedBytes x_data((size_t)nb * 18), y_data((size_t)nb * 34);
+    for (size_t i = 0; i < x_data.size(); i++)
+        x_data[i] = (uint8_t)(rng() & 0xff);
+    for (size_t i = 0; i < y_data.size(); i++)
+        y_data[i] = (uint8_t)(rng() & 0xff);
+    // Give the two deltas modest, exactly-representable fp16 values so the
+    // reference and Halide agree bit-for-bit-ish.
+    for (int i = 0; i < nb; i++) {
+        uint16_t xd = float16_t(0.5f + 0.25f * i).to_bits();
+        uint16_t yd = float16_t(1.0f + 0.5f * i).to_bits();
+        memcpy(&x_data[i * 18], &xd, 2);
+        memcpy(&y_data[i * 34], &yd, 2);
+    }
+
+    halide_dimension_t xs[1] = {{0, nb, 1, 0}};
+    halide_dimension_t ys[1] = {{0, nb, 1, 0}};
+    Buffer<> xbuf(q4_0_type(), x_data.data(), 1, xs);
+    Buffer<> ybuf(q8_0_type(), y_data.data(), 1, ys);
+
+    Pipeline p;
+    p.x.set(xbuf);
+    p.y.set(ybuf);
+    Buffer<float> out = p.result.realize();
+
+    float expected = reference_dot(x_data, y_data, nb);
+    if (std::abs(out() - expected) > 1e-2f * std::max(1.0f, std::abs(expected))) {
+        printf("q4_0/q8_0 dot product = %f, expected %f\n", out(), expected);
+        exit(1);
+    }
+}
+
+void test_arm_codegen() {
+    Pipeline p;
+    std::string s_path = Internal::get_test_tmp_dir() + "struct_type_dot_product.s";
+    p.result.compile_to_assembly(s_path, {p.x, p.y}, "qdot", arm_target());
+    std::string asm_text = read_file(s_path);
+
+    // The struct's packed qs arrays -- 16 q4_0 nibble bytes and 32 q8_0 int8s
+    // per block -- must load as whole 128-bit vectors, not element-by-element.
+    // A regression to per-field byte loads would show up as a burst of `ldrb`
+    // in the block loop; a good lowering emits `ldr`/`ldur q` (or `ld1`).
+    int byte_loads = count_occurrences(asm_text, "ldrb");
+    int vector_loads = count_occurrences(asm_text, "ldr\tq") +
+                       count_occurrences(asm_text, "ldur\tq") +
+                       count_occurrences(asm_text, "ld1\t");
+    // The fp16 deltas should convert inline (native fcvt), not via a byte load
+    // + soft-float path.
+    int fp16_converts = count_occurrences(asm_text, "fcvt\ts");
+
+    bool ok = byte_loads == 0 && vector_loads >= 3 && fp16_converts >= 2;
+    if (!ok) {
+        printf("Unexpected q4_0/q8_0 dot-product ARM codegen: byte_loads=%d (want 0), "
+               "vector_loads=%d (want >= 3), fp16_converts=%d (want >= 2).\n",
+               byte_loads, vector_loads, fp16_converts);
+        printf("---- ARM assembly ----\n%s\n", asm_text.c_str());
+        exit(1);
+    }
+}
+
+void test_packed_scalar_field_codegen() {
+    ImageParam blocks(q5_0_type(), 1, "blocks");
+    Var b("b");
+    Func qh("read_qh");
+    qh(b) = cast<uint32_t>(field(blocks(b), "qh"));
+
+    std::string s_path = Internal::get_test_tmp_dir() + "struct_type_packed_scalar.s";
+    qh.compile_to_assembly(s_path, {blocks}, "read_qh", arm_target());
+    std::string asm_text = read_file(s_path);
+
+    // The packed UInt(32) field starts at byte offset two and is therefore
+    // unaligned. It should nevertheless remain one wide load; decomposing it
+    // into four byte loads obscures common reuse by subsequent bit extracts.
+    int byte_loads = count_occurrences(asm_text, "ldrb");
+    int word_loads = count_occurrences(asm_text, "ldr\tw") +
+                     count_occurrences(asm_text, "ldur\tw");
+    if (byte_loads != 0 || word_loads < 1) {
+        printf("Unexpected packed UInt(32) struct-field ARM codegen: byte_loads=%d (want 0), "
+               "word_loads=%d (want >= 1).\n",
+               byte_loads, word_loads);
+        printf("---- ARM assembly ----\n%s\n", asm_text.c_str());
+        exit(1);
+    }
+}
+
+}  // namespace
+
+int main(int argc, char **argv) {
+    test_correctness();
+    test_arm_codegen();
+    test_packed_scalar_field_codegen();
+    printf("Success!\n");
+    return 0;
+}
diff --git a/test/error/CMakeLists.txt b/test/error/CMakeLists.txt
index 5b949b48a112..b8b1c6cc747e 100644
--- a/test/error/CMakeLists.txt
+++ b/test/error/CMakeLists.txt
@@ -135,6 +135,8 @@ tests(
     store_root_without_compute_root.cpp
     stream_loads_illegal_self.cpp
     stream_stores_illegal_rvar.cpp
+    struct_nested_field_load.cpp
+    struct_nested_field_store.cpp
     thread_id_outside_block_id.cpp
     too_many_args.cpp
     treat_rvar_as_var.cpp
diff --git a/test/error/struct_nested_field_load.cpp b/test/error/struct_nested_field_load.cpp
new file mode 100644
index 000000000000..1ca8a595c2c4
--- /dev/null
+++ b/test/error/struct_nested_field_load.cpp
@@ -0,0 +1,39 @@
+#include "Halide.h"
+#include <stdio.h>
+#include <vector>
+
+using namespace Halide;
+
+// A struct field whose own type is itself Type::Struct (a nested struct
+// field) can't be read directly out of a real (buffer-backed) struct-typed
+// Load: LowerStructTypes.cpp's lower_field_read_from_load can only
+// Reinterpret a field's bytes into a same-sized plain integer, which doesn't
+// work for a nested struct field (whose bits() is not its byte size).
+// Populating the buffer manually (rather than via a materialized Halide
+// Store) sidesteps the separate, equally-unsupported Store-side restriction
+// (see struct_nested_field_store.cpp) so that this specific, documented
+// read-side restriction is what actually fires.
+int main(int argc, char **argv) {
+    Type inner_t = Type::Struct({{"a", Int(32)}, {"b", UInt(8)}});
+    Type outer_t = Type::Struct({{"inner", inner_t}, {"c", Float(32)}});
+    const int num_blocks = 4;
+
+    std::vector<uint8_t> data((size_t)outer_t.bytes() * num_blocks, 0);
+    halide_dimension_t shape[1] = {
+        {0, num_blocks, 1, 0},
+    };
+    Buffer<> Wt_buf(outer_t, data.data(), 1, shape);
+
+    ImageParam Wt(outer_t, 1, "Wt");
+    Wt.set(Wt_buf);
+
+    Var i("i");
+    Func consumer("consumer");
+    Expr inner_val = field(field(Wt(i), "inner"), "a");
+    consumer(i) = inner_val;
+
+    Buffer<int32_t> result = consumer.realize({num_blocks});
+
+    printf("Success!\n");
+    return 0;
+}
diff --git a/test/error/struct_nested_field_store.cpp b/test/error/struct_nested_field_store.cpp
new file mode 100644
index 000000000000..722f74fc1b23
--- /dev/null
+++ b/test/error/struct_nested_field_store.cpp
@@ -0,0 +1,25 @@
+#include "Halide.h"
+#include <stdio.h>
+
+using namespace Halide;
+
+// A struct field whose own type is itself Type::Struct (a nested struct
+// field) can't be materialized into a real buffer: LowerStructTypes.cpp's
+// Store splitter can only Reinterpret a field's value into a same-sized
+// plain integer, and a struct's runtime tag is always plain UInt(8)
+// regardless of its real byte size, so splitting a nested struct field this
+// way is rejected with a clear error instead of miscompiling.
+int main(int argc, char **argv) {
+    Type inner_t = Type::Struct({{"a", Int(32)}, {"b", UInt(8)}});
+    Type outer_t = Type::Struct({{"inner", inner_t}, {"c", Float(32)}});
+    Var i("i");
+
+    Func producer("producer");
+    producer(i) = pack_struct(outer_t, {pack_struct(inner_t, {cast<int32_t>(i * 10), cast<uint8_t>(i)}), cast<float>(i) + 0.5f});
+    producer.compute_root();
+
+    Buffer<> result = producer.realize({4});
+
+    printf("Success!\n");
+    return 0;
+}