From 48eae68156168f4ec9fc6e99a85ab9ed8e3837ba Mon Sep 17 00:00:00 2001
From: Jim Blandy <jimb@red-bean.com>
Date: Sat, 14 Jun 2025 11:28:42 -0700
Subject: [PATCH] [naga spv-in] Lay out entry point output structs correctly.

In the SPIR-V front end, when generating Naga IR `Struct` types to
represent a SPIR-V entry point's `Output` variables, instead of saying
"0xFFFF, // shouldn't matter", follow the usual rules in assigning
struct member offsets and computing an overall
size (`TypeInner::Struct::span`) for the resulting struct type.
---
 naga/src/front/spv/function.rs                | 21 +++++++++++++++++--
 .../out/ir/spv-spec-constants.compact.ron     |  4 ++--
 naga/tests/out/ir/spv-spec-constants.ron      |  4 ++--
 naga/tests/out/msl/spv-quad-vert.msl          |  3 ++-
 4 files changed, 25 insertions(+), 7 deletions(-)

diff --git a/naga/src/front/spv/function.rs b/naga/src/front/spv/function.rs
index fe56ba833b..9c59d0a850 100644
--- a/naga/src/front/spv/function.rs
+++ b/naga/src/front/spv/function.rs
@@ -374,7 +374,15 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
         );
 
         // 3. copy the outputs from privates to the result
+        //
+        // It would be nice to share struct layout code here with `parse_type_struct`,
+        // but that case needs to take into account offset decorations, which makes an
+        // abstraction harder to follow than just writing out what we mean. `Layouter`
+        // and `Alignment` cover the worst parts already.
         let mut members = Vec::new();
+        self.layouter.update(module.to_ctx()).unwrap();
+        let mut next_member_offset = 0;
+        let mut struct_alignment = crate::proc::Alignment::ONE;
         let mut components = Vec::new();
         for &v_id in ep.variable_ids.iter() {
             let lvar = self.lookup_variable.lookup(v_id)?;
@@ -435,6 +443,11 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
                                 }
                             }
 
+                            let member_alignment = self.layouter[sm.ty].alignment;
+                            next_member_offset = member_alignment.round_up(next_member_offset);
+                            sm.offset = next_member_offset;
+                            struct_alignment = struct_alignment.max(member_alignment);
+                            next_member_offset += self.layouter[sm.ty].size;
                             members.push(sm);
 
                             components.push(function.expressions.append(
@@ -454,12 +467,16 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
                             }
                         }
 
+                        let member_alignment = self.layouter[result.ty].alignment;
+                        next_member_offset = member_alignment.round_up(next_member_offset);
                         members.push(crate::StructMember {
                             name: None,
                             ty: result.ty,
                             binding,
-                            offset: 0,
+                            offset: next_member_offset,
                         });
+                        struct_alignment = struct_alignment.max(member_alignment);
+                        next_member_offset += self.layouter[result.ty].size;
                         // populate just the globals first, then do `Load` in a
                         // separate step, so that we can get a range.
                         components.push(expr_handle);
@@ -545,7 +562,7 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
                         name: None,
                         inner: crate::TypeInner::Struct {
                             members,
-                            span: 0xFFFF, // shouldn't matter
+                            span: struct_alignment.round_up(next_member_offset),
                         },
                     },
                     span,
diff --git a/naga/tests/out/ir/spv-spec-constants.compact.ron b/naga/tests/out/ir/spv-spec-constants.compact.ron
index 669ddb338b..698a807258 100644
--- a/naga/tests/out/ir/spv-spec-constants.compact.ron
+++ b/naga/tests/out/ir/spv-spec-constants.compact.ron
@@ -160,10 +160,10 @@
                         binding: Some(BuiltIn(Position(
                             invariant: false,
                         ))),
-                        offset: 0,
+                        offset: 16,
                     ),
                 ],
-                span: 65535,
+                span: 32,
             ),
         ),
     ],
diff --git a/naga/tests/out/ir/spv-spec-constants.ron b/naga/tests/out/ir/spv-spec-constants.ron
index 43ed9003b6..5afd71ed65 100644
--- a/naga/tests/out/ir/spv-spec-constants.ron
+++ b/naga/tests/out/ir/spv-spec-constants.ron
@@ -251,10 +251,10 @@
                         binding: Some(BuiltIn(Position(
                             invariant: false,
                         ))),
-                        offset: 0,
+                        offset: 16,
                     ),
                 ],
-                span: 65535,
+                span: 32,
             ),
         ),
     ],
diff --git a/naga/tests/out/msl/spv-quad-vert.msl b/naga/tests/out/msl/spv-quad-vert.msl
index e96b906fdf..25bdd23df9 100644
--- a/naga/tests/out/msl/spv-quad-vert.msl
+++ b/naga/tests/out/msl/spv-quad-vert.msl
@@ -15,6 +15,7 @@ struct gl_PerVertex {
 };
 struct VertexOutput {
     metal::float2 member;
+    char _pad1[8];
     metal::float4 gl_Position;
 };
 
@@ -53,6 +54,6 @@ vertex main_Output main_(
     main_1(v_uv, a_uv_1, unnamed, a_pos_1);
     metal::float2 _e7 = v_uv;
     metal::float4 _e8 = unnamed.gl_Position;
-    const auto _tmp = VertexOutput {_e7, _e8};
+    const auto _tmp = VertexOutput {_e7, {}, _e8};
     return main_Output { _tmp.member, _tmp.gl_Position };
 }