Bug 201008 – [WHLSL] Remove getters/setters/anders

• Component: WebGPU
• Status: RESOLVED FIXED

[Saam Barati]

...

[Saam Barati]

Gonna start working on this

[Saam Barati]

Created attachment 377191 [details]
WIP

it begins

[Saam Barati]

Created attachment 377193 [details]
WIP

[Saam Barati]

This is going to entail changing metal code gen to emit expressions everywhere too

[Saam Barati]

Created attachment 377311 [details]
WIP

[Saam Barati]

Created attachment 377312 [details]
WIP

It can run compute_boids successfully. Metal compile times are down from ~32ms to ~20ms on my MBP

[Saam Barati]

Created attachment 377317 [details]
WIP

[Saam Barati]

Created attachment 377390 [details]
WIP

More things are working.

[Saam Barati]

Created attachment 377410 [details]
WIP

[Saam Barati]

Created attachment 377415 [details]
WIP

[Saam Barati]

Created attachment 377416 [details]
WIP

[Saam Barati]

Created attachment 377429 [details]
WIP

One remaining test that is failing, and then I need to add implement the texture native function inlining stuff in the new expression framework, and then I think this patch is done.

[Saam Barati]

Created attachment 377537 [details]
WIP

[Saam Barati]

Created attachment 377539 [details]
WIP

[Saam Barati]

Created attachment 377556 [details]
WIP

rebased

[Saam Barati]

Created attachment 377673 [details]
patch

[Saam Barati]

I think I need to change the stdlib code for matrices a bit. I forgot to reorder a few of the matrices.

[Saam Barati]

Created attachment 377738 [details]
patch

[Saam Barati]

Created attachment 377741 [details]
patch

[Robin Morisset]

Comment on attachment 377673 [details]
patch

View in context: https://bugs.webkit.org/attachment.cgi?id=377673&action=review

I've not yet finished reviewing, but I have a few comments.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp:-333
> -                return is<AST::NullLiteralType>(resolvableTypeReference->resolvableType()) ? Acceptability::Maybe : Acceptability::No;

I think this was the last use of Acceptability::Maybe, it could probably be removed.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp:483
> +    bool m_isVisitingParameters { false };

Another way of doing that would have been to add the check in VariableDeclarationsStatement, but I am fine with your approach.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp:1006
> +        auto addressSpace = baseInfo->typeAnnotation.isRightValue() ? AST::TypeAnnotation { AST::RightValue() } : AST::TypeAnnotation { AST::AbstractLeftValue() };

The name of this variable does really fit its content.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp:1059
> +        auto addressSpace = baseInfo->typeAnnotation.isRightValue() ? AST::TypeAnnotation { AST::RightValue() } : AST::TypeAnnotation { AST::AbstractLeftValue() };

ditto.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLHighZombieFinder.cpp:37
> +void findHighZombies(Program&)

Why not remove the file/pass altogether? It does not seem useful anymore.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLHighZombieFinder.cpp:40
> +    // https://bugs.webkit.org/show_bug.cgi?id=201251

Maybe mention this bug in the Changelog, if I understand correctly your current patch is not sound until this is done.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLProgram.cpp:380
> +    }

This is rather gross.
Is there a reason not to check simply that the length is >= 1, <= 4 and all letters are in {x, y, z, w} ?
It would additionally make it easier to allow letters in {r, g, b, a} later on (as in HLSL).

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp:38
> +    // https://bugs.webkit.org/show_bug.cgi?id=201251

If you make a pass do nothing, you may as well remove it.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt:258
> +    result[2][1] = x[1][2];

I find this super surprising, that result[2] is valid for float2x3.
Is it really what GLSL and MSL do?

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt:5396
> +    bool result = true;

The all function on matrices can probably be slightly simplified by relying on the all function on vectors.
e.g. return all(x[0]) && all(x[1]).
Not clear if it is worth the trouble though.

> Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt:5631
> +uint firstbitlow(uint x) {

This is not at all the HLSL semantics for this function. We should either change its name or change its definition. Matching the name from HLSL and having a silently different behavior is asinine.
But it might be from before and not something you changed.

> Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLDotExpression.h:60
> +    char m_padding[12];

What is this for? It at least needs some comment.

> Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h:-122
> -        if (auto* typeAnnotation = maybeTypeAnnotation())

How is this change connected to the rest of this patch? It seems potentially risky.

[Saam Barati]

Comment on attachment 377673 [details]
patch

View in context: https://bugs.webkit.org/attachment.cgi?id=377673&action=review

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp:-333
>> -                return is<AST::NullLiteralType>(resolvableTypeReference->resolvableType()) ? Acceptability::Maybe : Acceptability::No;
> 
> I think this was the last use of Acceptability::Maybe, it could probably be removed.

will do

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp:483
>> +    bool m_isVisitingParameters { false };
> 
> Another way of doing that would have been to add the check in VariableDeclarationsStatement, but I am fine with your approach.

I think your proposal is nicer, and I thought of doing it that way, but I didn't want to add another field just for this.

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp:1006
>> +        auto addressSpace = baseInfo->typeAnnotation.isRightValue() ? AST::TypeAnnotation { AST::RightValue() } : AST::TypeAnnotation { AST::AbstractLeftValue() };
> 
> The name of this variable does really fit its content.

will fix and name it as typeAnnotation

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLHighZombieFinder.cpp:37
>> +void findHighZombies(Program&)
> 
> Why not remove the file/pass altogether? It does not seem useful anymore.

I'm adding rules back here in:
https://bugs.webkit.org/show_bug.cgi?id=201251

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLHighZombieFinder.cpp:40
>> +    // https://bugs.webkit.org/show_bug.cgi?id=201251
> 
> Maybe mention this bug in the Changelog, if I understand correctly your current patch is not sound until this is done.

Will do.

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLProgram.cpp:380
>> +    }
> 
> This is rather gross.
> Is there a reason not to check simply that the length is >= 1, <= 4 and all letters are in {x, y, z, w} ?
> It would additionally make it easier to allow letters in {r, g, b, a} later on (as in HLSL).

yes. Good call. Will do

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp:38
>> +    // https://bugs.webkit.org/show_bug.cgi?id=201251
> 
> If you make a pass do nothing, you may as well remove it.

it's nicer to not delete the file and bring it back to life IMO so the version history lives on in a nicer way

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt:258
>> +    result[2][1] = x[1][2];
> 
> I find this super surprising, that result[2] is valid for float2x3.
> Is it really what GLSL and MSL do?

I fixed this in the latest patches. I made some bugs in my find/replacing in the standard library

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt:5396
>> +    bool result = true;
> 
> The all function on matrices can probably be slightly simplified by relying on the all function on vectors.
> e.g. return all(x[0]) && all(x[1]).
> Not clear if it is worth the trouble though.

I'll skip for now but I agree it's a nice change to make.

>> Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt:5631
>> +uint firstbitlow(uint x) {
> 
> This is not at all the HLSL semantics for this function. We should either change its name or change its definition. Matching the name from HLSL and having a silently different behavior is asinine.
> But it might be from before and not something you changed.

Unfortunately the diff for this file is really bad. I didn't change this.

Feel free to file a bug to fix it since you know more about this than I do.

>> Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLDotExpression.h:60
>> +    char m_padding[12];
> 
> What is this for? It at least needs some comment.

for one of the replaceWith. Not sure we really need a comment as we do this elsewhere

>> Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h:-122
>> -        if (auto* typeAnnotation = maybeTypeAnnotation())
> 
> How is this change connected to the rest of this patch? It seems potentially risky.

otherwise we run into a crash where we set the type annotation twice for a cloned constant.

This method is only used from constant cloning so I don't think there is any risk.

[Robin Morisset]

Comment on attachment 377741 [details]
patch

View in context: https://bugs.webkit.org/attachment.cgi?id=377741&action=review

(review still ongoing)

The one thing I am a bit wary of so far is the way that parentheses are added around expressions in the Metal generator.
It is not clear at all if they are put by the caller of checkErrorAndVisit, or by visit itself. Picking one and documenting it (and making it the rule uniformly) would probably save us from redundant parentheses and (more importantly) missing ones.

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:430
> +        auto tempReturnName = generateNextVariableName(); 

What is the point of this extra variable ?

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:551
> +        } else {

When do we ever hit this case?

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:647
> +        checkErrorAndVisit(callExpression.arguments()[i]);

You put the expression on the right of the '=' in parentheses for AssignmentExpression, but not here.
I don't see why it would be needed in one case but not the other. I think it is not needed, but am not 100% sure.
Maybe there should instead be parentheses around the assignment as a whole (since there are some around the callExpression as a whole).

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:650
> +    if (callExpression.arguments().size())

Nit: This and the "if(i) m_stringBuilder.append(", ");" above can be replaced by unconditionally appending ", " at the end of the loop iteration.

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:724
> +        checkErrorAndVisit(makeArrayReferenceExpression.leftValue());

Ah right, we have some special case for @ of a pointer.
I'll try to remember to add it to the spec (we went back and forth on this, and I thought it was out of the language).

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:113
> +void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDeclaration& nativeFunctionDeclaration, const Vector<MangledVariableName>& args, MangledVariableName resultName, Intrinsics& /*intrinsics*/, TypeNamer& typeNamer, Indentation<4>)

This "Intrinsics& /*intrinsics*/" is very weird. Why not remove it altogether if intrinsics is unused?

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:312
> +            "threadgroup_barrier(mem_flags::mem_device)",

Shouldn't there be some comma between these function calls?

[Robin Morisset]

> >> Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt:5631
> >> +uint firstbitlow(uint x) {
> > 
> > This is not at all the HLSL semantics for this function. We should either change its name or change its definition. Matching the name from HLSL and having a silently different behavior is asinine.
> > But it might be from before and not something you changed.
> 
> Unfortunately the diff for this file is really bad. I didn't change this.
> 
> Feel free to file a bug to fix it since you know more about this than I do.

Will do
 
> >> Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLDotExpression.h:60
> >> +    char m_padding[12];
> > 
> > What is this for? It at least needs some comment.
> 
> for one of the replaceWith. Not sure we really need a comment as we do this
> elsewhere

Thanks for all the answers. This one is the only one I disagree with, I'd like some comment everywhere we have this kind of padding. It does not have to be long, just to point to replaceWith.

[Saam Barati]

Comment on attachment 377741 [details]
patch

View in context: https://bugs.webkit.org/attachment.cgi?id=377741&action=review

>> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:430
>> +        auto tempReturnName = generateNextVariableName(); 
> 
> What is the point of this extra variable ?

So I didn't have to rewrite the emitPack code.

>> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:551
>> +        } else {
> 
> When do we ever hit this case?

this is the swizzle case

>> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:647
>> +        checkErrorAndVisit(callExpression.arguments()[i]);
> 
> You put the expression on the right of the '=' in parentheses for AssignmentExpression, but not here.
> I don't see why it would be needed in one case but not the other. I think it is not needed, but am not 100% sure.
> Maybe there should instead be parentheses around the assignment as a whole (since there are some around the callExpression as a whole).

I should probably just do it here too. It can't hurt and we don't need to worry about it

>> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:650
>> +    if (callExpression.arguments().size())
> 
> Nit: This and the "if(i) m_stringBuilder.append(", ");" above can be replaced by unconditionally appending ", " at the end of the loop iteration.

👍🏼

>> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:724
>> +        checkErrorAndVisit(makeArrayReferenceExpression.leftValue());
> 
> Ah right, we have some special case for @ of a pointer.
> I'll try to remember to add it to the spec (we went back and forth on this, and I thought it was out of the language).

I'm not against removing it. It seems useless. But I think we can leave it to another patch if we decide to

[Saam Barati]

The rule is for the caller to place parens to avoid ambiguity. It's never the callee's responsibility.

[Saam Barati]

Comment on attachment 377741 [details]
patch

View in context: https://bugs.webkit.org/attachment.cgi?id=377741&action=review

>> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:113
>> +void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDeclaration& nativeFunctionDeclaration, const Vector<MangledVariableName>& args, MangledVariableName resultName, Intrinsics& /*intrinsics*/, TypeNamer& typeNamer, Indentation<4>)
> 
> This "Intrinsics& /*intrinsics*/" is very weird. Why not remove it altogether if intrinsics is unused?

will remove

>> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:312
>> +            "threadgroup_barrier(mem_flags::mem_device)",
> 
> Shouldn't there be some comma between these function calls?

good catch. Will fix

[Robin Morisset]

Comment on attachment 377741 [details]
patch

r=me

[Saam Barati]

Created attachment 377762 [details]
patch for landing

[WebKit Commit Bot]

Comment on attachment 377762 [details]
patch for landing

Clearing flags on attachment: 377762

Committed r249351: <https://trac.webkit.org/changeset/249351>

[WebKit Commit Bot]

All reviewed patches have been landed.  Closing bug.

[Radar WebKit Bug Importer]

<rdar://problem/54902883>

[Myles C. Maxfield]

Comment on attachment 377762 [details]
patch for landing

View in context: https://bugs.webkit.org/attachment.cgi?id=377762&action=review

> Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp:57
> +        "struct WSLMatrix\n"

This looks like it's copied from the Metal headers, are we sure the original is licensed appropriately for inclusion in WebKit?

Also, I thought we adopted MSL's matrix semantics, so why can't we just use MSL's matrices directly?

[Saam Barati]

(In reply to Myles C. Maxfield from comment #32)
> Comment on attachment 377762 [details]
> patch for landing
> 
> View in context:
> https://bugs.webkit.org/attachment.cgi?id=377762&action=review
> 
> > Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp:57
> > +        "struct WSLMatrix\n"
> 
> This looks like it's copied from the Metal headers, are we sure the original
> is licensed appropriately for inclusion in WebKit?
> 
> Also, I thought we adopted MSL's matrix semantics, so why can't we just use
> MSL's matrices directly?
The reason is this:
```
template <typename T, int Cols, int Rows = Cols, typename _E = typename enable_if<is_scalar<T>::value && is_floating_point<T>::value && Cols >=2 && Rows >= 2>::type>
struct matrix
```

The key point being that bool matrices are not allowed. (And int matrices for that matter too, if we decide to add them in the future.) So when I tried using metal's builtin matrix class, I got a compile error for our bool matrix.

[Anders Carlsson]

> [WHLSL] Remove getters/setters/anders
                                 ^^^^^^
😦

[Sam Weinig]

(In reply to Anders Carlsson from comment #34)
> > [WHLSL] Remove getters/setters/anders
>                                  ^^^^^^
> 😦

You know what you did.

[Saam Barati]

(In reply to Anders Carlsson from comment #34)
> > [WHLSL] Remove getters/setters/anders
>                                  ^^^^^^
> 😦

🥺