Designing a uniformity opt-out

[alan-baker]

Split from #3479

This issue is for discussion of a design for a uniformity opt out mechanism.

There have been three proposals (to my knowledge):

1. An attribute (e.g. @uniform).
2. A new block statement (e.g. unchecked_uniformity { ... }).
3. Variants of built-in functions (e.g. uniform_textureSample(...)).

CC @jimblandy

[brendan-duncan]

For Unity shaders, the important part of the opt-out design is that shader translators like Tint or Naga can generate the opt-outs, as we will not be able to author the source shaders with the extra information.

[krogovin]

Looking at WGSL 2022-10-25 Minutes it starts with trying to make the escape hatch as local as possible, and then it returns to marking blocks as unsafe-like. In contrast, just having unsafe function versions maximizes locality.

What is the motivation for marking blocks of code as unsafe-like instead of just having unsafe functions? I do not see any benefit at all from the marking of blocks of code over having unsafe function variants.

[alan-baker]

We should also consider how uniformity might evolve in the future when designing this feature. Along those lines, two future directions come to mind:

1. Additional scopes
2. Descriptor arrays

Descriptor arrays in Vulkan require dynamically uniform access by default. There is a non-uniform capability, but that leads to less efficient code so indicating that type of access is likely useful to indicate.

Additional scopes are more future looking. I believe this will be useful for future features like subgroups, but might rely on underlying language extensions.

These two features make me leans towards using an attribute over new builtin names. It would be awkward to have builtins for each scope and wouldn't work at all for descriptor array access.

I personally prefer either attributes or builtins over a new statement type. Attributes have the extra advantage of being able to attached to more things than statements. For example, we could attach an attribute to a function parameter without having the statement opt-out cover a whole function.

[dneto0]

For example, we could attach an attribute to a function parameter without having the statement opt-out cover a whole function.

I think it's worth fleshing out a critical-user-journey of how this is used. For example, the "programmer knows better, imprecise analysis at fault" case:

1. Developer writes shader code, and feeds it to CreateShaderModule
2. CreateShaderModule fails with a message. The diagnostic should point to root cause: which value is non-uniform.
3. The programmer, knowing better, annotates the value to indicate it's actually a uniform value but that the analysis is too weak to know it.
4. Developer feeds the shader module to CreateShaderModule and CreateShaderModule succeeds.

Concretely, this is the original program:

@group(0) @binding(0) var t: texture_2d<f32>;
@group(0) @binding(1) var s: sampler;

@fragment
fn main(
   @builtin(position) p: vec4<f32>,
   @location(0) f: f32
) {
  if (f<0) {
    let c = textureSample(t, s, p.xy);
  }
}

Compiling with Tint, we get:

$ tint a.wgsl
a.wgsl:10:13 error: 'textureSample' must only be called from uniform control flow
    let c = textureSample(t, s, p.xy);
            ^^^^^^^^^^^^^

a.wgsl:9:3 note: control flow depends on non-uniform value
  if (f<0) {
  ^^

a.wgsl:9:7 note: reading from user-defined input 'f' may result in a non-uniform value
  if (f<0) {
      ^

So in step 3 the programmer sanitizes the uniformity analysis, e.g. by declaring that they know that the f value is actually uniform:

@group(0) @binding(0) var t: texture_2d<f32>;
@group(0) @binding(1) var s: sampler;

@fragment
fn main(
   @builtin(position) p: vec4<f32>,
   @uniform @location(0) f: f32  // Programmer says they know the 'f' value is uniform.
) {
  if (f<0) {
    let c = textureSample(t, s, p.xy);
  }
}

Benefits:

• Cooperation/interaction between tooling and developer. The analysis acts as a guardrail that teaches but allows the bypass
• Using an attribute like this captures programmer intent: "I know better than the analysis can know, that the 'f' value is uniform".
  • For this case, the attribute is better than options 2 and 3, which "just" silence the analysis.

Another interesting case: "programmer knows better, accepts tradeoff in quality due to the incorrectly computed implicit derivative". In this scenario "just" silencing the analysis might be a better match for programmer intent. It would actually be a lie to assert in the source code that a value is uniform when the programmer knows it isn't uniform.

[dneto0]

Where could you put an attribute? We have a lot of choices in the design space:

• on a formal parameter, as above: Asserts the value is always uniform
• on a let: Asserts the value is always uniform
• on a var: Asserts the value stored in the memory locations is always uniform, regardless of what initializations and assignments occur
• could also be on an expression. I'd suggest limiting it to a singular expression
  • This could include a parenthesized expression. So if/when we get differently-scoped uniformity, we could use two attribute names to avoid the parsing ambigutity when you reach a parenthesis @uniform (without params) and a differently named thing like @uniform_for(subgroup)

I'm reminded again that SPIR-V has Uniform and UniformId which very closely mirrors this attribute option.

[tex3d]

I like the attribute approach described here by @dneto0. I also like the ability to supply the uniform scope (like group, subgroup, quad) with something like @uniform_for(subgroup). I think that scope is essential to describe the different degrees of uniformity necessary for different operations such as the screen-space gradients in fragment shaders. Of course, this scope impacts uniform analysis, so it would have to be updated to independently track the different uniform scopes.

To build on it, I think a name like @assume_uniform or @assert_uniform could be a better fit. You really are trying to promise that the value is uniform for analysis purposes. What @assert_uniform implies is that some debug-mode runtime could potentially catch a violated assumption by adding checks. The potential for debug-mode checks and the implication that you're lying if you violate the assumption or assertion based on the wording in the attribute, psychologically discourages blindly applying it anytime you get an analysis error (unlike a function/block/built-in variant approach).

I don't think we should be making it normal to call functions that require a level of uniformity that it isn't actually present, and just rely on implementations to have some predictable fallback behavior when requirements are violated. If someone is doing that, they should fix the shader code.

Regarding this debate between variable/expression attributes vs. function, block, or built-in variant based opt-out, I think the variable/expression attribute fits into the uniformity analysis story most cleanly, and potentially provides many benefits down the line. I think that provably uniform can be tracked separately from user-asserted assumptions as well.

The downside is that if the different needed uniformity scopes are not properly captured/expressed with the attribute, a user would be required to "lie" about uniformity in some cases (promising a value is fully uniform rather than subgroup or quad uniform which might be all that was true and required). This "lie" could lead to invalid assumptions about uniformity on anything dependent on this value, and failures if those assumptions are checked or enforced at runtime. So this approach requires a lot more work to achieve a full solution.

The function/block based opt-out seems to just turn off the error for sensitive operations in a given scope, and built-in variants are similar for the operations themselves. This means either a) uniformity analysis will still be wrong about some unknown number of things, uncorrected by the opt-out mechanism, or b) we just give a blanket pass to code that actually violates the rules. It doesn't leave behind the detailed information required to document where precisely the uniform analysis is too conservative (or just can't know), allow the runtime to catch the introduction of a new source of non-uniformity that wasn't intentional (help catch real bugs), or allow for uniform-aware codegen / optimizations as highlighted by prior comments.

Given these trade-offs, I wonder if a global opt-out could be provided until a proper attribute-on-value based approach was ready and robust, after which the global opt-out could eventually be deprecated. Code using global opt-out could be prevented from running in some less trusted mode, or without some compatibility setting in some future, or assumptions could be verified/enforced with extra runtime cost. Putting properly-scoped attributes on the originating source of (analysis-based) potential non-uniformity could reduce assumption validation/enforcement overhead for a "safe" mode as well.

[kdashg]

WGSL 2022-11-01 Minutes

• MM: What about a uniformity opt-in?
• KG: Want developers to have to think about it first. Don’t want to start with the confusing behavior to start with. Want the guardrails on by default.
• KG: would like to see a proposal for which functions would need to be marked with nonuniform options. Eg. what entry points are modified. Generally I like if annotations for builtins for cleanliness.
• DN: Ben made comment that one thing he doesn’t like with new names for builtins, is that they don’t actually do different things.
• DN: Also, in v1, we have just one level of scopes, but we expect to have more scopes in the future. The number of builtin names you’d need would multiply.
• MM: I believe it’s group opinion that we need UAn for compute shaders for safety.
• DN: I do think that question is still open, e.g. other ways to clog the gpu
• JB: Does WebGL have this problem?
• KG: You can write very slow shaders in webgl
• MM:
• JB: I just wonder if we’re adding any new risks here or not, vs webgl?
• DN:
• TR: Just an idea, could you have a “more trusted” mode?
• JB: One nice thing about AB’s thing, attributes on vars and argos, plus attribs on statements are both useful things. I could imagine you want both. We maybe don’t need to choose, it might be useful to do both.
• KG: Whatever gets us consensus faster.
• JB: Yes. Also btw we do have a “collective “ section
• AB: Won’t always be builtin functions at fault, e.g. descriptor arrays (when we get to that feature). Attributes are the lowest cost thing right now. Gives you the freedom to modify your code in the smallest scope way. E.g. putting it on a comparison, rather than the whole if-statement. Once it’s a statement, you can’t get below a certain level of granularity without refactoring your code. Attr
• DN: I added an example of a user experience where the user knows better than UAn, and diags tell them what’s wrong, and they say “no that’s fine actually” when they mark with an attrib.
• DN: Flipside is when an author is instead merely willing to accept bad/wrong derivatives, and annotates instead acceptability rather than actually correct. Like “pretend uniform”.
• AB: SPIR-V has prior art: has Uniform decoration and UniformID (for scoped). And for descriptor arrays it has a “non-uniform” control for descriptor arrays.
• KG: Have several options here. My dream was how Rust has unsafe block and unsafe expressions, and also unsafe annotations on functions. The MVP for what we need would be annotations on functions. What would make me super satisfied is function-level annotation. That’s tolerable and minimal.
• JB: Have a question. Talking about attributes. But putting it things on stuff we don’t annotate now. Types, return types, parameters, struct members. We don’t put them on expressions yet. When someone says we put attributes I think of putting them on things we already have. Are you suggesting putting on expressions.
• AB: Yes, that’s a possibility.
• JB: Seems to be roughly isomorphic between unsafe-expr and attr.
• JB: Surgical approach seems not acceptable to Unity. Need to have “a” opt-out which can be used with cooperation with the code generator in the bulk way. Would prefer surgical,...
• KG: In extrema could have unsafe-annotations on all functions.
• AB: Unity’s only requirement is that tooling can generate it for them.
• JB: But don’t want it strewn all over.
• KG: You only really need those when certain spots. If your tooling is smart enough about which have to be non-uniform calls, then it can emit more surgical annotations/wrappers as needed.
• TR: Could have subgroup operations assumeUniform and makeUniform (latter would pick first lane’s value), to make it actually uniform.
• TR: Don’t like mangling the functions.
• KG: Most people don’t like those.
• EG: Think the premise of wider opt-out is needing a standard mechanism. Ok… resetting understanding.
• MM: Have said before. Jury is still out on what the security requirements are. It may actually be the case that both derivatives and barriers require the analysis. That could be the outcome.
• AB: Touching on makeUniform. We have workgroupBroadcast proposal for compute shaders. But don’t think it’s clear for fragment.
• TR: For fragment, need quads to cooperate.
• AB: This is where we have the issue of whether we have the invocations in the quad actually getting together again to supply their values.
• MM: Is there a quad barrier?
• TR: No.
• AB: There is a vulkan extension for subgroup uniform control flow that moves the reconvergence requirement down to subgroup level rather than whole invocation group. But that’s not universally available to where we’re targeting WebGPU.
  • VK_KHR_shader_subgroup_uniform_control_flow
• KG: You can kinda do certain transforms to incur total reconvergence but it’s not trivial or great.
• DN: Concretely how about:
  • @uniform and @fake_uniform (bikeshed name). Applicable to various things, and captures programmer intent.
• EG: In Rust practice, you use a single ‘unsafe’ construct but document it as you go (in comment?)
• TR: ‘assume_uniform’ does the same thing.
• JB: Watch out for Rust’s unsafe. Rust unsafe functions and blocks were incorrectly unified in first versions. Two different concepts actually. Let’s be cautious about unsafe functions and unsafe blocks.
• KG: That’s why I want to look to someone else. Otherwise we may fall into the same trap. Like, they had troubles with this, but then they fixed it.
• AB: If nobody is allergic to using an attribute as the solution, then @assume_uniform can capture both cases. Then I can take this and make it concrete. Then the discussion is what things to annotate.
• JB: Consider future direction. The analysis can report an error whose cause threads back through several functions. When that occurs, we can provide detailed diagnostics. However, it will be annoying for users. Programmes know where they thought things were going to be uniform. They have no way to say their expectations about their program. Because we can’t capture programmer intent at that granularity, we’ll cause difficulty for those difficult spots. Think we’ll want to put annotations on function-level boundary. User will want to say how each function will behave. Produces a summary on each function of effect. Will want the analysis to check their expectations. That will be useful just for the sake of being able to debug. Let’s keep that in mind, maybe not necessary for V1.
• KG: Namewise, I want the name to be negative and scary. Really like “unsafe” for that reason. Negative connotation is really important. Want people to not want to apply this opt-out.
• MM: +1 on what Jim said earlier. Historically expected an author to assert that a particular variable is uniform. Would cause an error. Good for cutting off the length of action-at-a-distance.
• JB: +1 Kelsey’s point about negative connotation. People are swayed by connotations of ‘unsafe’ and ‘use strict’ in JavaScript.
• AB: Understand that argument. Where you end up changing your code differs depending on positive or negative connotation.
• EG (in chat): take_off_kneeepads
• ZJ (in chat): ignore_non_uniform

[tex3d]

Follow up comment: I recently found out about Vulkan's uniformity requirement for barrier being "dynamic uniform". My previous assumption that attributes on values would be sufficient is not valid for the "dynamic uniform" case.

Here's a pseudocode example for "dynamic uniform"

if (a()) {
  // a() is not uniform, and exec here is not uniform
} else {
  // exec here is not uniform
  if (b()) {
    // b() is not uniform, and exec here is not uniform
  } else {
    // exec is dynamic uniform here due to (compiler unknown) dependency between a() and b()
    // uniform if and only if a() is false and b() is false.
    // if (!a() && !b()) on any thread, then this is true for *all* threads.
    barrier();
  }
  // exec here is not uniform
}

So if "dynamic uniform" is to be supported, it seems that some form of attribute on the barrier() call, or on the block, or on a function wrapping the barrier() call, is necessary. No value or expression attribute will suffice in this case.

Note: FXC may not allow this case for HLSL, and I know of at least one implementation that cannot run this code as expected on D3D.

[kdashg]

WGSL 2022-11-08 Minutes

• KG: What I really like about using the negative term instead of the positive term is that it suggests that things are normally safe but we’re using an escape hatch. Flags code for extra-careful review.
• KG: When reviewing an unsafe block, the goal is to become sure that what’s in the unsafe block is indeed safe - it’s just that the compiler can’t see that
• AB: If we were to write an attribute-based proposal with a placeholder name, would that be all right? Would that let us move forward on substantive questions like, where can this be used? Separately from naming
• DG: re: attribute: In one of the issues, someone was porting rendering and there was a variable input to the shader that was uniform, but the analysis had no way of knowing that. I support the uniform if it supplies additional information about the program that the compiler can’t determine.
• KG: And if that’s violated, we say, “that’s on you”?
• DG: yes
• DG: Unless people think we can get away without the ability to mark knowledge that the programmer has
• KG: We should go as far as we can in that direction
• TR: I made a long comment on the issue, but I recently found out about Vulkan’s “dynamic uniformity” requirement for barriers - they’re not quite the same as what I expected. This kind of broke the idea that you could simply attribute values and get a guarantee of uniformity in control flow for barriers. Marking individual values with attributes wouldn’t be sufficient, but you can create control flow that can’t be marked up sufficiently.
• AB: It depends on the use case. If you could put an attribute on an expression, you could put it on the condition of an if, or put it on a let and then use that in the if
• TR: It could be non-uniform when false, but uniform when true. You could have two if statements, and in the … (See comment on issue)
• AB: If you could put the attribute on the function call, then you could put it on the builtin itself, avoiding a proliferation of attributes on builtins. This gives it a lot of flexibility
• TR: There’s a difference between a value being uniform, and the control flow operating on that value being uniform
• AB: The key for the analysis is that what causes a problem is a collective operation called from non-uniform control flow. That could depend on the uniformity of a value. But with an attribute on expressions, you can always make it work.
• TR: But also uniformity scope matters. So I don’t know if we can capture everything we need. We need a simple escape hatch for v1.
• AB: The bare minimum is to put an attribute on function calls. That would solve everyone’s problems, it might not always be idea but it would suffice
• KG: For 1.0, messy suffixes might not be so bad
• AB: That would also suffice
• KG: Nobody prefers it, but it’s shippable
• BC: We would never be able to get rid of it
• KG: We would have to support them forever, but we could discourage people from using them once we have something better. And they’re not hard to implement
• BC: Is that preferable to module scope?
• KG: Yes, preferable, even for v1. Module scope will become something people slap on as a habit, without thinking about it.
• BC: I prefer function-scope annotations to littering the language with functions we don’t want to keep. That would just turn off errors on a particular function
• TR: How would that affect the rest of the analysis, using the results of those functions?
• BC: You could proceed with the analysis, but simply silence the warnings. Or, you could avoid propagating the non-uniformity
• AB: If I branch on a non-uniform value in a flagged function and use a barrier, would I get an error?
• ZJ: Is this an annotation that works within a function, or does it affect the function’s callers as well?
• KG: If you mark a function, you’re saying, just pretend this is uniform. So you would not get errors if you used its value
• ZJ: Maybe if we are going to have function-scope opt-out, where we put an attribute on the whole function, then does that mean that any expressions in the body of that function will not produce errors?
• KG: yes
• TR: It makes much more sense to preserve the analysis within the function, but mute the errors
• BC: I’m worried about the implementation complexity of determining where to place the attributes, if you wanted to avoid placing them unnecessarily. Right now, we only return one error and then bail, so having to identify all the places where it’s needed would be a major effort
• AB: Having an attribute on a function call would clearly work just as well as having variant functions with ugly names. Unity’s feedback is, they don’t care how we do it - just as long as the tool does it for them
• BC: Let’s discuss internally
• KG: I do like the equivalence between the attribute on a function call and calling a function that has a suffix, it’s nice to have those two equivalents. I don’t quite understand the implementation complexity concerns around collecting multiple errors
• KG: Can we do other things to the language, like the uniformRead primitive (Uniformity analysis: find a way to load from a storage buffer uniformly. #2321), that would make these annotations less necessary in the first place?

[dneto0]

Follow up comment: I recently found out about Vulkan's uniformity requirement for barrier being "dynamic uniform". My previous assumption that attributes on values would be sufficient is not valid for the "dynamic uniform" case.

Here's a pseudocode example for "dynamic uniform"

if (a()) {
  // a() is not uniform, and exec here is not uniform
} else {
  // exec here is not uniform
  if (b()) {
    // b() is not uniform, and exec here is not uniform
  } else {
    // exec is dynamic uniform here due to (compiler unknown) dependency between a() and b()
    // uniform if and only if a() is false and b() is false.
    // if (!a() && !b()) on any thread, then this is true for *all* threads.
    barrier();
  }
  // exec here is not uniform
}

To make sure I understand the concern, I wanted to make this more concrete.

The example is meaningful if some invocations reach the barrier.
A good meaningful case is if all invocations reach the barrier.

So a() and b() should both be false, but the compiler must believe they are MayBeNonUniform.
Then at runtime make w = 1, so that both a() and b() return false.

So:

     var<workgroup> w: i32;
     fn a() -> bool  { return  (w & 1) == 0; } // w is even
     fn b() -> bool  { return  (w & 3) == 0; } // w is a multiple of 4

     @compute @workgroup(3)
     fn main(@builtin(local_invocation_index) lid: u32) {
        if (lid==0) { w = 1; }  // falsify a() and b()
        workgroupBarrier();

        // now copy the example code fragment from above.
        if (a()) {
          // a() is not uniform, and exec here is not uniform.  [[1]]
        } else {
           // exec here is not uniform [[2]]
           if (b()) {
             // b() is not uniform, and exec here is not uniform.  [[3]]
           } else {
             // [[4]]
             // exec is dynamic uniform here due to (compiler unknown) dependency between a() and b()
             // uniform if and only if a() is false and b() is false.
             // if (!a() && !b()) on any thread, then this is true for *all* threads.
             workgroupBarrier();
           }
           // exec here is not uniform.  [[5]]
        }
     }

And now I see @tex3d's point.
Any set of attributes applied to the conditions or the 'if' constructs themselves to promise uniformity ends up being either a lie, or not powerful enough to give us the desired ability to "see" that the inner barrier call is uniform.

• Let's assume any attribute on a condition is a promise that a condition or flow is uniform.
• Then any attributes we apply individually on the evaluations of conditions a() and b() must get us far enough to say that execution at [[4]] is uniform.
• But then that would also either say that:
  • that [[1]] and [[2]] are executing uniformly because we promised evaluation of a() is uniform, which is a lie,
  • or that [[3]] is executing uniformly, because we promised evaluating b() is uniform, which is also a lie

That's a very nice demonstration that getting what we want when we only have "BLAH IS UNIFORM" attributes forces us to lie at least a little bit.

[dneto0]

In internal discussions a 4th option has come been expressed (that I wasn't aware of until recently), which is a mashup of the other three:

4. An attribute on a call to a builtin to ignore uniformity problems with an invocation of the scope.

Formally this is like option 3's uniform_textureSample but where it's spelled as two tokens instead of one: @unchecked_uniformity textureSample That avoids the concern about future combinatorial blowup of number of builtins.

Google doesn't have an opinion on this yet. We're still talking.

[dneto0]

After internal discussion, Google requires a coarse grain control to be present. Our preference for the coarse grain control is an option on CreateShaderModule (i.e. not in the shader, and so it's out of band). We suggest a name such as assumeDynamicUniformityForDerivatives.