...
Gonna start working on this
Created attachment 377191 [details] WIP it begins
Created attachment 377193 [details] WIP
This is going to entail changing metal code gen to emit expressions everywhere too
Created attachment 377311 [details] WIP
Created attachment 377312 [details] WIP It can run compute_boids successfully. Metal compile times are down from ~32ms to ~20ms on my MBP
Created attachment 377317 [details] WIP
Created attachment 377390 [details] WIP More things are working.
Created attachment 377410 [details] WIP
Created attachment 377415 [details] WIP
Created attachment 377416 [details] WIP
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.
Created attachment 377537 [details] WIP
Created attachment 377539 [details] WIP
Created attachment 377556 [details] WIP rebased
Created attachment 377673 [details] patch
I think I need to change the stdlib code for matrices a bit. I forgot to reorder a few of the matrices.
Created attachment 377738 [details] patch
Created attachment 377741 [details] patch
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.
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.
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?
> >> 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.
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
The rule is for the caller to place parens to avoid ambiguity. It's never the callee's responsibility.
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
Comment on attachment 377741 [details] patch r=me
Created attachment 377762 [details] patch for landing
Comment on attachment 377762 [details] patch for landing Clearing flags on attachment: 377762 Committed r249351: <https://trac.webkit.org/changeset/249351>
All reviewed patches have been landed. Closing bug.
<rdar://problem/54902883>
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?
(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.
> [WHLSL] Remove getters/setters/anders ^^^^^^ 😦
(In reply to Anders Carlsson from comment #34) > > [WHLSL] Remove getters/setters/anders > ^^^^^^ > 😦 You know what you did.
(In reply to Anders Carlsson from comment #34) > > [WHLSL] Remove getters/setters/anders > ^^^^^^ > 😦 🥺