diff --git a/cuda-oxide/crates/mir-lower/src/convert/types/enum_layout.rs b/cuda-oxide/crates/mir-lower/src/convert/types/enum_layout.rs
index 25eb043d10..74b80948d4 100644
--- a/cuda-oxide/crates/mir-lower/src/convert/types/enum_layout.rs
+++ b/cuda-oxide/crates/mir-lower/src/convert/types/enum_layout.rs
@@ -2011,4 +2011,195 @@ mod tests {
             "{error}"
         );
     }
+
+    #[test]
+    fn enum_slot_map_pointer_niche_with_u128_payload_preserves_align16() {
+        let mut ctx = make_ctx();
+
+        let logical = mir_uint(&mut ctx, 8);
+        let wide = mir_uint(&mut ctx, 128);
+        let u32_ty = mir_uint(&mut ctx, 32);
+        let pointer: TypeHandle = MirPtrType::get_generic(&mut ctx, u32_ty, false).into();
+
+        let payload: TypeHandle = MirTupleType::get_with_layout(
+            &mut ctx,
+            vec![wide, pointer],
+            vec![0, 1],
+            vec![0, 16],
+            32,
+            16,
+        )
+        .into();
+
+        let enum_ty: TypeHandle = MirEnumType::get_with_encoding(
+            &mut ctx,
+            "MaybeWideAndPointer".into(),
+            logical,
+            vec![0, 1],
+            vec![
+                EnumVariant::unit("None".into()),
+                EnumVariant::new_with_layout("Some".into(), vec![payload], vec![0], vec![32]),
+            ],
+            EnumEncoding {
+                tag_offset: 16,
+                total_size: 32,
+                abi_align: 16,
+                layout_kind: EnumLayoutKind::Niche,
+                carrier_kind: EnumCarrierKind::Pointer,
+                carrier_width: 64,
+                untagged_variant: 1,
+                variant_inhabited: vec![1, 1],
+                ..EnumEncoding::default()
+            },
+        )
+        .into();
+
+        let map = build_enum_slot_map(&mut ctx, enum_ty).unwrap();
+
+        let fields = struct_fields(&ctx, map.llvm_struct_ty);
+
+        // Required enum alignment is 16, while the 16-byte filler itself is
+        // deliberately represented as two i64 leaves with natural alignment 8.
+        // build_enum_slot_map supplies the zero-length alignment anchor separately.
+        assert_eq!(fields.len(), 4);
+
+        let filler_ref = fields[1].deref(&ctx);
+        let filler = filler_ref
+            .downcast_ref::<llvm_types::ArrayType>()
+            .expect("16-byte enum filler should lower to an LLVM array");
+
+        assert_eq!(
+            filler.size(),
+            2,
+            "16-byte enum filler should use two 64-bit leaves"
+        );
+
+        assert_eq!(
+            filler
+                .elem_type()
+                .deref(&ctx)
+                .downcast_ref::<IntegerType>()
+                .map(IntegerType::width),
+            Some(64),
+            "16-byte enum filler elements should be i64"
+        );
+
+        assert_eq!(
+            llvm_type_size_align(&ctx, map.llvm_struct_ty),
+            Some((32, 16)),
+            "typed 16-byte filler must preserve rustc's enum size and alignment"
+        );
+    }
+
+    /// `(element count, element width in bits)` of an array-typed filler.
+    fn filler_array(ctx: &Context, filler: TypeHandle) -> (u64, u32) {
+        let filler_ref = filler.deref(ctx);
+        let array = filler_ref
+            .downcast_ref::<llvm_types::ArrayType>()
+            .expect("enum filler should lower to an LLVM array");
+        let elem_width = array
+            .elem_type()
+            .deref(ctx)
+            .downcast_ref::<IntegerType>()
+            .map(IntegerType::width)
+            .expect("enum filler elements should be integers");
+        (array.size(), elem_width)
+    }
+
+    #[test]
+    fn enum_filler_16_bytes_at_align4_uses_four_i32s() {
+        let mut ctx = make_ctx();
+
+        let filler = make_enum_filler_type(&mut ctx, 0, 16, 4);
+
+        assert_eq!(
+            filler_array(&ctx, filler),
+            (4, 32),
+            "i64 would raise the enum alignment to 8; i32 is the widest word that fits"
+        );
+        assert_eq!(llvm_type_size_align(&ctx, filler), Some((16, 4)));
+    }
+
+    #[test]
+    fn enum_filler_16_bytes_uses_two_i64s_at_align8() {
+        let mut ctx = make_ctx();
+
+        let filler = make_enum_filler_type(&mut ctx, 0, 16, 8);
+
+        let filler_ref = filler.deref(&ctx);
+        let array = filler_ref
+            .downcast_ref::<llvm_types::ArrayType>()
+            .expect("16-byte filler should use an LLVM array");
+
+        assert_eq!(array.size(), 2);
+
+        assert_eq!(
+            array
+                .elem_type()
+                .deref(&ctx)
+                .downcast_ref::<IntegerType>()
+                .map(IntegerType::width),
+            Some(64)
+        );
+
+        assert_eq!(llvm_type_size_align(&ctx, filler), Some((16, 8)));
+    }
+
+    #[test]
+    fn enum_filler_keeps_bytes_at_an_odd_offset() {
+        let mut ctx = make_ctx();
+
+        let filler = make_enum_filler_type(&mut ctx, 1, 16, 8);
+
+        assert_eq!(
+            filler_array(&ctx, filler),
+            (16, 8),
+            "no integer width divides offset 1, so a wider word would move every later field"
+        );
+        assert_eq!(llvm_type_size_align(&ctx, filler), Some((16, 1)));
+    }
+
+    #[test]
+    fn enum_filler_keeps_bytes_in_a_byte_aligned_enum() {
+        let mut ctx = make_ctx();
+
+        let filler = make_enum_filler_type(&mut ctx, 0, 16, 1);
+
+        assert_eq!(
+            filler_array(&ctx, filler),
+            (16, 8),
+            "any wider word would raise the enum alignment above rustc's 1"
+        );
+        assert_eq!(llvm_type_size_align(&ctx, filler), Some((16, 1)));
+    }
+
+    #[test]
+    fn enum_filler_keeps_bytes_when_no_width_divides_the_size() {
+        let mut ctx = make_ctx();
+
+        let filler = make_enum_filler_type(&mut ctx, 0, 15, 8);
+
+        assert_eq!(filler_array(&ctx, filler), (15, 8));
+        assert_eq!(llvm_type_size_align(&ctx, filler), Some((15, 1)));
+    }
+
+    #[test]
+    fn enum_filler_24_bytes_uses_three_i64s_at_align8() {
+        let mut ctx = make_ctx();
+
+        let filler = make_enum_filler_type(&mut ctx, 8, 24, 8);
+
+        assert_eq!(filler_array(&ctx, filler), (3, 64));
+        assert_eq!(llvm_type_size_align(&ctx, filler), Some((24, 8)));
+    }
+
+    #[test]
+    fn enum_filler_12_bytes_uses_three_i32s_at_align4() {
+        let mut ctx = make_ctx();
+
+        let filler = make_enum_filler_type(&mut ctx, 4, 12, 4);
+
+        assert_eq!(filler_array(&ctx, filler), (3, 32));
+        assert_eq!(llvm_type_size_align(&ctx, filler), Some((12, 4)));
+    }
 }
diff --git a/cuda-oxide/crates/mir-lower/src/convert/types/layout.rs b/cuda-oxide/crates/mir-lower/src/convert/types/layout.rs
index 54083e2424..2488ed3fe1 100644
--- a/cuda-oxide/crates/mir-lower/src/convert/types/layout.rs
+++ b/cuda-oxide/crates/mir-lower/src/convert/types/layout.rs
@@ -24,44 +24,75 @@ pub(super) fn make_padding_type(ctx: &mut Context, size: u64) -> TypeHandle {
 /// Storage for `size` bytes of enum filler at byte `offset`: bytes the payload
 /// occupies that no typed slot claims.
 ///
-/// `[N x i8]` is byte-exact but costs one leaf *per byte* everywhere a payload
-/// value is built, merged, or taken apart. `Option<&[T]>` is the shape that
-/// shows it: the niche carrier claims the pointer, the length is left to an
-/// 8-byte filler, and building the value then decomposes that length into
-/// eight `i8` `insertvalue`s, every block merge carries eight `phi i8`, and
-/// reading it back is a chain of `prmt` byte permutes with the aggregate
-/// spilled to `.local` in between. Measured on sm_86, it costs the *pointer*
-/// too: it rides through the same byte soup, `InferAddressSpaces` can no
-/// longer follow it, and the access lands in the generic window (`ld.v2.b64`)
-/// instead of `ld.global.v2.b64`.
+/// The filler is a vehicle, not a claim. Nothing reads it field-wise: a
+/// payload without a typed slot round-trips through memory as a whole
+/// aggregate, and both ends of that round trip are byte-exact. What the
+/// filler's type does decide is the leaf count of the enum struct, where a
+/// leaf is one scalar field left after every array and nested struct is
+/// flattened.
 ///
-/// One integer of the same width is the same bytes in one leaf. It is only
-/// legal where it moves nothing:
+/// The leaf count matters because of `scalarize_block_args`. pliron's
+/// `mem2reg` turns an enum local into a block argument of the whole struct,
+/// and that pass splits such an argument into scalar block arguments only
+/// while the struct has at most `MAX_LEAVES` (16) leaves: `collect_leaves`
+/// refuses the 17th leaf and any array over 16 elements. A `[N x i8]` filler
+/// costs one leaf per byte. `Option<&[T]>` leaves its length to an 8-byte
+/// filler, and as `[8 x i8]` that length was eight `i8` `insertvalue`s, eight
+/// `phi i8` at every merge, and a chain of `prmt` byte permutes with the
+/// aggregate spilled to `.local` in between. `Option<([u64; 2], &[T])>`
+/// stored as `{ [16 x i8], ptr, i64 }` is 18 leaves, so its `Some`/`None`
+/// join stays one aggregate phi; `opt` cannot take that apart,
+/// `InferAddressSpaces` cannot follow the pointer through the `extractvalue`,
+/// and the access lands in the generic window (`ld.v2.b64`) instead of
+/// `ld.global.v2.b64`. Stored as `{ [2 x i64], ptr, i64 }` it is 4 leaves,
+/// the join becomes scalar phis, and the pointer keeps its address space.
 ///
-/// - the width is 2, 4 or 8 bytes, so the integer is byte-faithful (a multiple
-///   of 8 bits, hence no padding of its own). 16 is deliberately excluded:
-///   `i128` is legal LLVM but lowers to register pairs on NVPTX, so widening
-///   that far trades one win for another cost and wants its own measurement;
-/// - `offset` is a multiple of the width, so LLVM inserts no gap ahead of the
-///   field and every later field keeps its byte offset;
-/// - the width does not exceed the enum's alignment, or the struct's natural
-///   alignment would rise above rustc's.
+/// The rule: take the widest width among 8, 4 and 2 bytes that divides both
+/// `size` and `offset` and does not exceed the enum's alignment; the filler
+/// is that integer when it covers `size` alone, else an array of it. rustc
+/// types its own padding by the same rule (`type_padding_filler`: the largest
+/// of i64, i32 and i16 whose alignment fits). Each condition keeps the layout
+/// fixed:
 ///
-/// Anything else keeps the byte array. The filler is a *vehicle*, not a claim:
-/// nothing reads it field-wise, because a payload that has no typed slot
-/// round-trips through memory as a whole aggregate, and both ends of that
-/// round trip are byte-exact. `build_enum_slot_map`'s own size and alignment
-/// assertions — hard errors, not debug checks — backstop all three conditions.
+/// - the width divides `size`, so the array covers the same bytes;
+/// - `offset` is a multiple of the width, so the natural struct walk inserts
+///   no gap ahead of the field and every later field keeps its byte offset;
+/// - the width does not exceed the enum's alignment, so the struct's natural
+///   alignment never rises above rustc's.
+///
+/// `i128` is left out: it would impose 16-byte natural alignment, where
+/// `[2 x i64]` carries the same 16 bytes at alignment 8.
+///
+/// One integer word may cover payload bytes plus Rust padding. Those padding
+/// bits are read as undef, the same assumption the 8-byte rule already makes
+/// when one `i64` covers a small field and its tail padding.
+///
+/// What stays a byte array: a run no width divides, such as the 23 bytes at
+/// offset 1 behind a `bool` carrier in `Option<(bool, [u64; 2])>`, and every
+/// filler of a byte-aligned enum. What still exceeds the leaf budget even when
+/// typed: a filler above 14 i64 words (112 bytes) beside a pointer and a
+/// length, since 14 + 2 is the last count that fits. Both keep an aggregate
+/// phi and the generic access that comes with it; the `bool` carrier case
+/// also keeps its `.local` spill. The budget itself is the follow-up.
+/// `build_enum_slot_map`'s size and alignment checks are hard errors and
+/// backstop every case.
 pub(super) fn make_enum_filler_type(
     ctx: &mut Context,
     offset: u64,
     size: u64,
     abi_align: u64,
 ) -> TypeHandle {
-    if matches!(size, 2 | 4 | 8) && offset.is_multiple_of(size) && size <= abi_align.max(1) {
-        let width = u32::try_from(size * 8).expect("size is at most 8, so width is at most 64");
-        return IntegerType::get(ctx, width, Signedness::Signless).into();
+    for width in [8u64, 4, 2] {
+        if size.is_multiple_of(width) && offset.is_multiple_of(width) && width <= abi_align {
+            let bits = u32::try_from(width * 8).expect("width is at most 8, so at most 64 bits");
+            let int_ty: TypeHandle = IntegerType::get(ctx, bits, Signedness::Signless).into();
+            if size == width {
+                return int_ty;
+            }
+            return llvm_types::ArrayType::get(ctx, int_ty, size / width).into();
+        }
     }
+
     make_padding_type(ctx, size)
 }
 
diff --git a/cuda-oxide/crates/mir-lower/src/scalarize_block_args.rs b/cuda-oxide/crates/mir-lower/src/scalarize_block_args.rs
index 908bb1e0db..c26c80d59b 100644
--- a/cuda-oxide/crates/mir-lower/src/scalarize_block_args.rs
+++ b/cuda-oxide/crates/mir-lower/src/scalarize_block_args.rs
@@ -653,4 +653,79 @@ mod tests {
         let branch = op_cast::<dyn BranchOpInterface>(term_obj.as_ref()).unwrap();
         assert_eq!(branch.successor_operands(&ctx, 0).len(), 1);
     }
+
+    /// Enum storage `{ [count x iN], i64, i64 }`: one filler array beside a
+    /// pointer-sized carrier and a length (an `i64` stands in for the `ptr`
+    /// leaf; only the leaf count matters here). Builds a merge of that
+    /// aggregate fed by two predecessors, runs the pass, and returns how many
+    /// arguments the merge block carries afterwards: one when the aggregate
+    /// stays whole, one per leaf when it splits.
+    fn merge_arguments_for_filler_storage(elem_bits: u32, count: u64) -> usize {
+        let mut ctx = make_ctx();
+        let elem: TypeHandle = IntegerType::get(&ctx, elem_bits, Signedness::Signless).into();
+        let i64_ty: TypeHandle = IntegerType::get(&ctx, 64, Signedness::Signless).into();
+        let filler: TypeHandle = llvm_types::ArrayType::get(&ctx, elem, count).into();
+        let agg_ty: TypeHandle = llvm_types::StructType::get_unnamed(
+            &ctx,
+            (
+                vec![filler, i64_ty, i64_ty],
+                llvm_types::StructLayout::Unpacked,
+            ),
+        )
+        .into();
+
+        let (module_ptr, entry) = build_kernel(&mut ctx, vec![], vec![]);
+
+        let undef_op = llvm::UndefOp::new(&mut ctx, agg_ty);
+        undef_op.get_operation().insert_at_back(entry, &ctx);
+        let aggregate = undef_op.get_operation().deref(&ctx).get_result(0);
+
+        let other = append_block(&mut ctx, entry, vec![]);
+        let merge = append_block(&mut ctx, entry, vec![agg_ty]);
+        let br_op = llvm::BrOp::new(&mut ctx, merge, vec![aggregate]);
+        br_op.get_operation().insert_at_back(entry, &ctx);
+
+        let other_undef = llvm::UndefOp::new(&mut ctx, agg_ty);
+        other_undef.get_operation().insert_at_back(other, &ctx);
+        let other_agg = other_undef.get_operation().deref(&ctx).get_result(0);
+        let other_br = llvm::BrOp::new(&mut ctx, merge, vec![other_agg]);
+        other_br.get_operation().insert_at_back(other, &ctx);
+
+        let return_op = llvm::ReturnOp::new(&mut ctx, None);
+        return_op.get_operation().insert_at_back(merge, &ctx);
+
+        scalarize_aggregate_block_args(&mut ctx, module_ptr).unwrap();
+
+        merge.deref(&ctx).get_num_arguments()
+    }
+
+    /// The filler's type decides whether enum storage fits the leaf budget.
+    /// `Option<([u64; 2], &[T])>` stored as `{ [16 x i8], ptr, i64 }` is 18
+    /// leaves and stays one aggregate argument, which is how the pointer
+    /// inside it lost its address space; the same bytes typed as `[2 x i64]`
+    /// are 4 leaves and split. `[14 x i64]` beside two words is the last
+    /// filler that fits (16 leaves); `[15 x i64]` is one over.
+    #[test]
+    fn typed_fillers_decide_whether_enum_storage_splits() {
+        assert_eq!(
+            merge_arguments_for_filler_storage(8, 16),
+            1,
+            "{{ [16 x i8], i64, i64 }} is 18 leaves and must stay whole"
+        );
+        assert_eq!(
+            merge_arguments_for_filler_storage(64, 2),
+            4,
+            "{{ [2 x i64], i64, i64 }} is 4 leaves and must split"
+        );
+        assert_eq!(
+            merge_arguments_for_filler_storage(64, 14),
+            16,
+            "{{ [14 x i64], i64, i64 }} is 16 leaves, the last count that fits"
+        );
+        assert_eq!(
+            merge_arguments_for_filler_storage(64, 15),
+            1,
+            "{{ [15 x i64], i64, i64 }} is 17 leaves and must stay whole"
+        );
+    }
 }
diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs
index 6a5cf6115b..255b16ace7 100644
--- a/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs
+++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/vectorization/src/main.rs
@@ -155,6 +155,136 @@ mod view_kernels {
         empty.as_mut_ptr()
     }
 
+    /// Regression coverage for a 16-byte enum filler beside a pointer niche.
+    ///
+    /// The `[u64; 2]` payload occupies 16 bytes before the slice pointer. Enum
+    /// storage must represent that filler without obscuring the slice pointer's
+    /// global address space; otherwise the final vector load becomes generic
+    /// `ld.v4.*` instead of `ld.global.v4.*`.
+    #[kernel]
+    pub fn f32x4_view_copy_through_wide_option(
+        input: &[f32],
+        mut output: DisjointSlice<f32>,
+        lo: u64,
+        hi: u64,
+    ) {
+        let i = thread::index_1d().get();
+
+        let out_len = output.len();
+        let out_flat = unsafe { core::slice::from_raw_parts_mut(output.as_mut_ptr(), out_len) };
+
+        let Some(input_quads) = vector::as_vectors::<F32x4>(input) else {
+            return;
+        };
+
+        let Some(out_quads) = vector::as_vectors_mut::<F32x4>(out_flat) else {
+            return;
+        };
+
+        let cookie = [lo, hi];
+
+        let wrapped = if lo == u64::MAX {
+            None
+        } else {
+            Some((cookie, input_quads))
+        };
+
+        let Some((_cookie, quads)) = wrapped else {
+            return;
+        };
+
+        if i < quads.len() && i < out_quads.len() {
+            out_quads[i] = quads[i];
+        }
+    }
+
+    /// Regression coverage for a 16-byte enum filler beside a mutable pointer niche.
+    ///
+    /// This mirrors `f32x4_view_copy_through_wide_option`, but places the output
+    /// slice behind the wide enum payload. Losing its address space would turn the
+    /// vector store into generic `st.v4.*` instead of `st.global.v4.*`.
+    #[kernel]
+    pub fn f32x4_view_store_through_wide_option(
+        input: &[f32],
+        mut output: DisjointSlice<f32>,
+        lo: u64,
+        hi: u64,
+    ) {
+        let i = thread::index_1d().get();
+
+        let Some(input_quads) = vector::as_vectors::<F32x4>(input) else {
+            return;
+        };
+
+        let out_len = output.len();
+        let out_flat = unsafe { core::slice::from_raw_parts_mut(output.as_mut_ptr(), out_len) };
+
+        let Some(out_quads) = vector::as_vectors_mut::<F32x4>(out_flat) else {
+            return;
+        };
+
+        let cookie = [lo, hi];
+
+        let wrapped = if lo == u64::MAX {
+            None
+        } else {
+            Some((cookie, out_quads))
+        };
+
+        let Some((_cookie, out_quads)) = wrapped else {
+            return;
+        };
+
+        if i < input_quads.len() && i < out_quads.len() {
+            out_quads[i] = input_quads[i];
+        }
+    }
+
+    /// Regression coverage for a 24-byte enum filler beside a pointer niche.
+    ///
+    /// `[u64; 3]` leaves 24 unclaimed bytes before the slice pointer. Typed
+    /// as `[3 x i64]` the enum storage is 5 leaves, under the block-argument
+    /// scalarizer's budget of 16, so the `Some`/`None` join splits into
+    /// scalar phis. As `[24 x i8]` it was 26 leaves, the join stayed one
+    /// aggregate phi, and the vector load lost its global window the same
+    /// way the 16-byte case did.
+    #[kernel]
+    pub fn f32x4_view_copy_through_24_byte_option(
+        input: &[f32],
+        mut output: DisjointSlice<f32>,
+        lo: u64,
+        hi: u64,
+    ) {
+        let i = thread::index_1d().get();
+
+        let out_len = output.len();
+        let out_flat = unsafe { core::slice::from_raw_parts_mut(output.as_mut_ptr(), out_len) };
+
+        let Some(input_quads) = vector::as_vectors::<F32x4>(input) else {
+            return;
+        };
+
+        let Some(out_quads) = vector::as_vectors_mut::<F32x4>(out_flat) else {
+            return;
+        };
+
+        let cookie = [lo, hi, lo ^ hi];
+
+        let wrapped = if lo == u64::MAX {
+            None
+        } else {
+            Some((cookie, input_quads))
+        };
+
+        let Some((_cookie, quads)) = wrapped else {
+            return;
+        };
+
+        if i < quads.len() && i < out_quads.len() {
+            out_quads[i] = quads[i];
+        }
+    }
+
     /// Copy one `F32x4` quad per thread between flat `f32` buffers through
     /// checked [`vector::as_vectors`] views. The quad is moved whole, never
     /// decomposed into lanes, so both the load and the store fuse into
@@ -250,6 +380,53 @@ fn main() {
         .expect("f32x4_view_copy launch");
     let flat_round_trip = flat_out.to_host_vec(&stream).expect("flat readback") == flat;
 
+    // The wide-option kernels route the same copy through
+    // `Option<([u64; 2], &[F32x4])>` and `Option<([u64; 3], &[F32x4])>`: a
+    // 16- or 24-byte payload beside the slice pointer. Launch each into a
+    // fresh zeroed buffer and check the round trip, then take the `None` arm
+    // once (`lo == u64::MAX`) and check that the output stays untouched.
+    let wide_option_output = |launch: &dyn Fn(&mut DeviceBuffer<f32>)| -> Vec<f32> {
+        let mut out = DeviceBuffer::<f32>::zeroed(&stream, N * 4).expect("wide-option output");
+        launch(&mut out);
+        out.to_host_vec(&stream).expect("wide-option readback")
+    };
+    // SAFETY (each launch below): launch shape/resources match the kernel;
+    // buffers cover its accesses.
+    let wide_copy_round_trip = wide_option_output(&|out| {
+        unsafe {
+            view_module.f32x4_view_copy_through_wide_option(&stream, cfg, &flat_in, out, 123, 456)
+        }
+        .expect("f32x4_view_copy_through_wide_option launch");
+    }) == flat;
+    let wide_store_round_trip = wide_option_output(&|out| {
+        unsafe {
+            view_module.f32x4_view_store_through_wide_option(&stream, cfg, &flat_in, out, 123, 456)
+        }
+        .expect("f32x4_view_store_through_wide_option launch");
+    }) == flat;
+    let wide_none_untouched = wide_option_output(&|out| {
+        unsafe {
+            view_module.f32x4_view_copy_through_wide_option(
+                &stream,
+                cfg,
+                &flat_in,
+                out,
+                u64::MAX,
+                456,
+            )
+        }
+        .expect("f32x4_view_copy_through_wide_option launch, None arm");
+    })
+    .iter()
+    .all(|v| *v == 0.0);
+    let option_24_round_trip = wide_option_output(&|out| {
+        unsafe {
+            view_module
+                .f32x4_view_copy_through_24_byte_option(&stream, cfg, &flat_in, out, 123, 456)
+        }
+        .expect("f32x4_view_copy_through_24_byte_option launch");
+    }) == flat;
+
     // Inspect the code we just launched: read the PTX back from the artifact
     // bundle embedded in this binary. Those are the same bytes `kernels::load`
     // handed to the driver, so the shape assertion can never diverge from what
@@ -269,6 +446,22 @@ fn main() {
         errors += 1;
         println!("  !! f32x4_view_copy round-trip copy mismatch");
     }
+    if !wide_copy_round_trip {
+        errors += 1;
+        println!("  !! f32x4_view_copy_through_wide_option round-trip copy mismatch");
+    }
+    if !wide_store_round_trip {
+        errors += 1;
+        println!("  !! f32x4_view_store_through_wide_option round-trip copy mismatch");
+    }
+    if !wide_none_untouched {
+        errors += 1;
+        println!("  !! f32x4_view_copy_through_wide_option wrote output through the None arm");
+    }
+    if !option_24_round_trip {
+        errors += 1;
+        println!("  !! f32x4_view_copy_through_24_byte_option round-trip copy mismatch");
+    }
 
     if let Some(ptx) = &ptx {
         let document = ptx_parse::Document::parse(ptx).expect("parse embedded PTX");
@@ -326,6 +519,32 @@ fn main() {
                 }
             }
         }
+
+        // A wide enum payload next to the slice pointer must not hide the
+        // pointer's global address space from NVPTX address-space inference.
+        // The enum storage has to stay under the block-argument scalarizer's
+        // leaf budget: a typed filler (`[2 x i64]`, `[3 x i64]`) does, a byte
+        // array (`[16 x i8]`, `[24 x i8]`) does not, and then the merged
+        // pointer rides an aggregate phi that `InferAddressSpaces` cannot
+        // see through. The copy kernels take the input slice through the
+        // enum (load side); the store kernel takes the output slice.
+        for (kernel, dir, filler) in [
+            ("f32x4_view_copy_through_wide_option", "ld", "16-byte"),
+            ("f32x4_view_store_through_wide_option", "st", "16-byte"),
+            ("f32x4_view_copy_through_24_byte_option", "ld", "24-byte"),
+        ] {
+            let body = kernel_body(&document, kernel);
+            match find_128bit_global(body, dir) {
+                Some(line) => println!("{kernel}: {line}"),
+                None => {
+                    errors += 1;
+                    let access = if dir == "ld" { "load" } else { "store" };
+                    println!(
+                        "  !! {kernel}: {filler} enum filler lost the 128-bit global {access}"
+                    );
+                }
+            }
+        }
     } else {
         println!(
             "PTX shape check skipped: the embedded bundle has no PTX text to \