wgsl: @align(n) must divide required-align-of, for all structs #4978

dneto0 · 2024-11-19T01:13:19Z

Builds on #4974

Reverses: #3756

- The previous phrasing was a note that was "implied" by other rules.
  It wasn't actually perfectly implied.
- Avoids a silly and confusing case with the first element of
  a struct.
- Brings the spec more in line with Naga and WebKit.
- More clearly applies the rule for *all* structs, no matter
  if it is actually instantiated by a variable.

github-actions · 2024-11-19T01:17:33Z

Previews, as seen when this build job started (f059fe6):
WebGPU _{webgpu.idl | Explainer | Correspondence Reference}
WGSL _{grammar.js | wgsl.lalr.txt}

alan-baker

My understanding is that both Naga and WebKit also check it evenly divides align-of too. Should that be an additional constraint?

alan-baker · 2024-11-19T01:18:09Z

wgsl/index.bs

@@ -10594,7 +10605,7 @@ used in address space |C|.
  <tr algorithm="alignment of an runtime-sized array">
      <td>array&lt;T&gt;
      <td>[=AlignOf=](|S|)
-      <td>[=roundUp=](16, [=AlignOf=](|S|))
+      <td>not applicable


Yes, good catch.

dneto0 · 2024-11-19T01:27:34Z

My understanding is that both Naga and WebKit also check it evenly divides align-of too. Should that be an additional constraint?

I think with this rule change, @align on a member can't be too small, therefore Align of a struct always evenly divides RequiredAlignOf for the struct.

dneto0 · 2024-11-19T15:14:43Z

Here is some sample code showing differences between Tint, Naga, and WebKit compilers.

litmus.wgsl.txt

// Preliminaries: If compilation fails here, then struct attributes are
// checked independently of how they are used in variables.
struct S     { a: u32, @align(4) b: vec4u }        // pass tint, fail naga, fail wgslc
struct S2    { a: u32, @align(4) b: vec4u }        // pass tint, fail naga, fail wgslc

// SS and SU put the bad element at offset 0.
struct SS    { @align(4) a: vec4u }                // pass tint, fail naga, fail wgslc
struct SU    { @align(4) a: vec4u }                // pass tint, fail naga, fail wgslc
struct SS4   { @align(1) a: u32 }                  // pass tint, fail naga, fail wgslc
struct SU4   { @align(1) a: u32 }                  // pass tint, fail naga, fail wgslc

// Validation of bool alignment.
//struct B0 { @align(0) a: bool }   //XFAIL(positive) fail tint, fail naga, fail wgslc
struct B4 { @align(4) a: bool }                    // pass tint, pass naga, pass wgslc
struct B1 { @align(1) a: bool }                    // pass tint, pass naga, pass wgslc
// struct B3 { @align(3) a: bool }//XFAIL(power_of_2) fail tint, fail naga, fail wgslc

// Ensure compilers check the @size attribute at all
// Validation of u32 size
struct US4 { @size(4) a: u32 }                     // pass tint, pass naga, pass wgslc
//struct US1 { @size(1) a: u32 }              //XFAIL fail tint, fail naga, fail wgslc
//struct US3 { @size(3) a: u32 }              //XFAIL fail tint, fail naga, fail wgslc

// Validation of bool size
// struct BS0 { @size(0) a: bool }  //XFAIL(positive) fail tint, fail naga, fail wgslc
struct BS4 { @size(4) a: bool }                    // pass tint, pass naga, pass wgslc
//struct BS1 { @size(1) a: bool }                  // fail tint(boolsize4), pass naga, pass wgslc
//struct BS3 { @size(3) a: bool }                  // fail tint(boolsize4), pass naga, pass wgslc


// If we got past the preliminaries, then see which address spaces are checked.
// var<workgroup> workgroup_var: S;                               // fail tint, fail naga, fail wgslc
// var<private> private_var: S;                                   // fail tint, fail naga, fail wgslc

// These have the bad alignment at offset 0.
//var<workgroup> wg_ss: SS;                                       // fail tint, fail naga, fail wgslc
//var<private> private_ss: SS;                                    // fail tint, fail naga, fail wgslc
// @group(0) @binding(0) var<storage> storage_var: SS;            // fail tint, fail naga, fail wgslc
// @group(0) @binding(1) var<storage> storage_array: array<SS>;   // fail tint, fail naga, fail wgslc
// @group(0) @binding(1) var<uniform> uniform_var: SU;            // fail tint, fail naga, fail wgslc
// @group(0) @binding(1) var<uniform> uniform_array: array<SU,2>; // fail tint, fail naga, fail wgslc

@compute @workgroup_size(1)
fn main() {
 //let check_all_structs = S();                    // pass tint, fail naga, fail wgslc
 //let check_all_arrays = array<S2,2>();           // pass tint, fail naga, fail wgslc

// var private_var: S;                             // fail tint, fail naga, fail wgslc
// var private_array: array<S2,2>;                 // fail tint, fail naga, fail wgslc

// None of the following are needed for Tint to validate.
// _ = &workgroup_var;
// _ = &private_var;
// _ = &storage_var;
// _ = &storage_array;
// _ = &uniform_var;
// _ = &uniform_array;
}

jrprice · 2024-11-19T16:43:06Z

wgsl/index.bs

    then |n| [=shader-creation error|must=] satisfy:
-    |n|&nbsp;=&nbsp;|k|&nbsp;&times;&nbsp;[=RequiredAlignOf=](|T|,|C|)
+    <blockquote>
+    |n|&nbsp;=&nbsp;|k|&nbsp;&times;&nbsp;[=RequiredAlignOf=](|T|,|AS|)


Could we not just use AlignOf() here instead of RequiredAlignOf(), to make this rule address-space agnostic? Then you wouldn't need the !uniform condition.

ah yes, that works. It works because it recurses to a nested type.

Fixed this now.

I don't think so, because we should plan for a future when we do scalar block layout.
We need to keep Align(..) only about setting layouts, and use RequiredAlignOf to be the constraints the environment places on memory accesses.

kdashg · 2024-11-19T20:02:26Z

WGSL 2024-11-19 Minutes

Milestone?
DN: Had reason to look at this again. Put this rule into tint, and compared to naga/webkit, and tint was doing something different: Only when instantiating in a variable, gets corner cases for looking at the first field of a struct. Confusing, we decided we wanted to just make the spec match what naga/webkit does. Looking at the history here, we made things complicated in order to leave options open for future, but now think that we should just change things if we need to in the future.
JB: So concretely, if you ask for a not strict enough alignment, it becomes an error?
DN: yes
JB,MW: +1
Resolved: Accepted.

kdashg

Approved by WG in meeting.

Reverses: gpuweb#3756 - The previous phrasing was a note that was "implied" by other rules. It wasn't actually perfectly implied. - Avoids a silly and confusing case with the first element of a struct. - Brings the spec more in line with Naga and WebKit. - More clearly applies the rule for *all* structs, no matter if it is actually instantiated by a variable.

dneto0 · 2024-11-20T15:00:34Z

wgsl/index.bs

    then |n| [=shader-creation error|must=] satisfy:
    <blockquote>
-    |n|&nbsp;=&nbsp;|k|&nbsp;&times;&nbsp;[=RequiredAlignOf=](|T|,|AS|)
+    |n|&nbsp;=&nbsp;|k|&nbsp;&times;&nbsp;[=AlignOf=](|T|)


Thinking about this again. Not sure this is the right change.
Does this still leave open the ability to do scalar block layout, by lowering the alignment of vec4u to that of u32?

ok, talked this through with @dj2

I will revert this change because:

right now the two phrasings are the same.

but in future when we want to introduce something like scalar block layout, we need a path to doing that.

if we only have Align as a knob, then we have to change Align(vec4u) to be 4.

But then that is also the default alignment used to set the layout. And then that changes the layout of structs that *don't have annotations: e.g. { a:u32, b:vec4u} would then place b at offset 8 instead of 16.

dneto0 · 2024-11-20T16:14:55Z

Here's some separating cases, looking forward to scalar-block-layout #4040

// SBL = scalar block layout.  Changes RequiredAlignOf(vec4<T>) to sizeof(T)
// 'align' after a struct is alignment of the struct type.

struct Natural     { a: u32,            b: vec4u };  // offsets 0 16, align 16  // should pass
struct Force4      { a: u32, @align(4)  b: vec4u };  // offsets 0 4,  align 4.  // should fail today, pass with SBL
struct Force64     { a: u32, @align(64) b: vec4u };  // offsets 0 64, align 64  // should pass

// nesting
struct Natural_Natural  { nestA: u32,             nestB: Natural } // offsets 0 16,  align 16  // should pass
struct Force4__Natural  { nestA: u32, @align(4)   nestB: Natural } // offsets 0 4,   align 4   // should fail today, pass with SBL
struct Force64_Natural  { nestA: u32, @align(64)  nestB: Natural } // offsets 0 64,  align 64  // should pass

struct Natural_Force4   { nestA: u32,             nestB: Force4  } // offsets 0, 4,  align 4   // should fail today, pass with SBL
struct Force4__Force4   { nestA: u32, @align(4)   nestB: Force4  } // offsets 0, 4,  align 4   // should fail today, pass with SBL
struct Force64_Force4   { nestA: u32, @align(64)  nestB: Force4  } // offsets 0, 64, align 64  // should pass

struct Natural_Force64  { nestA: u32,             nestB: Force64 } // offsets 0, 64, align 64  // should pass
struct Force4__Force64  { nestA: u32, @align(4)   nestB: Force64 } // offsets 0, 4,  align 4   // fail: 4 does not divide align(Force64)
struct Force64_Force64  { nestA: u32, @align(64)  nestB: Force64 } // offsets 0, 64, align 64  // should pass

alan-baker

Minor typo

wgsl/index.bs

Co-authored-by: alan-baker <alanbaker@google.com>

dneto0 requested a review from alan-baker November 19, 2024 01:13

dneto0 added the wgsl WebGPU Shading Language Issues label Nov 19, 2024

dneto0 force-pushed the align-normative branch from e5bb03e to 29c76b0 Compare November 19, 2024 01:15

alan-baker reviewed Nov 19, 2024

View reviewed changes

dneto0 mentioned this pull request Nov 19, 2024

Test fixes for alignment validation landing gpuweb/cts#4016

Merged

jrprice reviewed Nov 19, 2024

View reviewed changes

dneto0 mentioned this pull request Nov 19, 2024

Plainly state alignment requirement for scalars and vectors #4965

Closed

kdashg approved these changes Nov 19, 2024

View reviewed changes

dneto0 force-pushed the align-normative branch from 29c76b0 to 077b1b8 Compare November 19, 2024 21:32

dneto0 requested a review from jrprice November 19, 2024 21:32

dneto0 commented Nov 20, 2024

View reviewed changes

dneto0 force-pushed the align-normative branch from 077b1b8 to 1962a00 Compare November 20, 2024 15:18

dneto0 mentioned this pull request Nov 20, 2024

Extension: Scalar block layout #4040

Open

alan-baker requested changes Nov 20, 2024

View reviewed changes

wgsl/index.bs Outdated Show resolved Hide resolved

Add missing word

f059fe6

Co-authored-by: alan-baker <alanbaker@google.com>

dneto0 requested a review from alan-baker November 21, 2024 04:35

Merge branch 'main' into align-normative

2a4d7f9

alan-baker approved these changes Nov 21, 2024

View reviewed changes

dneto0 merged commit bd061d4 into gpuweb:main Nov 21, 2024
4 checks passed

dneto0 deleted the align-normative branch November 21, 2024 14:45

dneto0 mentioned this pull request Nov 21, 2024

enhancement: uniform buffer can use standard layout #4973

Open

wgsl: @align(n) must divide required-align-of, for all structs #4978

wgsl: @align(n) must divide required-align-of, for all structs #4978

Uh oh!

Conversation

dneto0 commented Nov 19, 2024

Uh oh!

github-actions bot commented Nov 19, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

alan-baker left a comment

Choose a reason for hiding this comment

Uh oh!

alan-baker Nov 19, 2024

Choose a reason for hiding this comment

Uh oh!

dneto0 commented Nov 19, 2024

Uh oh!

dneto0 commented Nov 19, 2024

Uh oh!

jrprice Nov 19, 2024

Choose a reason for hiding this comment

Uh oh!

dneto0 Nov 19, 2024

Choose a reason for hiding this comment

Uh oh!

dneto0 Nov 19, 2024

Choose a reason for hiding this comment

Uh oh!

dneto0 Nov 20, 2024

Choose a reason for hiding this comment

Uh oh!

kdashg commented Nov 19, 2024

Uh oh!

kdashg left a comment

Choose a reason for hiding this comment

Uh oh!

dneto0 Nov 20, 2024

Choose a reason for hiding this comment

Uh oh!

dneto0 Nov 20, 2024

Choose a reason for hiding this comment

Uh oh!

dneto0 commented Nov 20, 2024

Uh oh!

alan-baker left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

github-actions bot commented Nov 19, 2024 •

edited

Loading