[webkit-changes] [WebKit/WebKit] 9e532e: [WGSL] Don't rewrite global variable accesses

[Tadeu Zagallo]

  Branch: refs/heads/main
  Home:   https://github.com/WebKit/WebKit
  Commit: 9e532e28386370bd19b8bf5269c6ffec649cc5ed
      https://github.com/WebKit/WebKit/commit/9e532e28386370bd19b8bf5269c6ffec649cc5ed
  Author: Tadeu Zagallo <tzagallo at apple.com>
  Date:   2023-05-23 (Tue, 23 May 2023)

  Changed paths:
    M Source/WebGPU/WGSL/AST/ASTIdentifier.h
    M Source/WebGPU/WGSL/GlobalVariableRewriter.cpp
    M Source/WebGPU/WGSL/Metal/MetalCodeGenerator.cpp
    M Source/WebGPU/WGSL/Metal/MetalFunctionWriter.cpp
    M Source/WebGPU/WGSL/WGSLShaderModule.h

  Log Message:
  -----------
  [WGSL] Don't rewrite global variable accesses
https://bugs.webkit.org/show_bug.cgi?id=257135
rdar://109667180

Reviewed by Myles C. Maxfield.

Global variables get moved into a struct that represents the ArgumentBuffer layout.
Currently, we rewrite the entrypoint to take an additional parameter, and rewrite
all references to the global into a field access into the parameter. e.g.

@group(0) @binding(0) var s: sampler;
@compute main() { f(s); }

gets rewritten into:

struct AB { s: sampler }
@compute main(ab: AB) { f(ab.s); }

The problem with this is when add we another user-defined function called from multiple
entry points. The different entry points may not share the same argument buffer layout,
so we cannot pass the struct directly. There are a few alternative options on how
we can pass such values, but this patch implements this by materializing all the globals
in the entry point, so that we no longer need to to rewrite all access. The same example
will look as follows:

struct AB { s: sampler }
@compute main(ab: AB) { let s = ab.s; f(s); }

We introduce a new variable declaration for each global, bringing it into scope, rather
than rewrite all the references to all globals.

* Source/WebGPU/WGSL/AST/ASTIdentifier.h:
(WGSL::AST::Identifier::make):
* Source/WebGPU/WGSL/GlobalVariableRewriter.cpp:
(WGSL::RewriteGlobalVariables::visit):
(WGSL::RewriteGlobalVariables::visitEntryPoint):
(WGSL::RewriteGlobalVariables::determineUsedGlobals):
(WGSL::RewriteGlobalVariables::insertStructs):
(WGSL::RewriteGlobalVariables::insertMaterializations):
(WGSL::RewriteGlobalVariables::shouldBeReference const):
* Source/WebGPU/WGSL/Metal/MetalCodeGenerator.cpp:
* Source/WebGPU/WGSL/Metal/MetalFunctionWriter.cpp:
(WGSL::Metal::FunctionDefinitionWriter::write):
(WGSL::Metal::FunctionDefinitionWriter::visit):
(WGSL::Metal::FunctionDefinitionWriter::serializeVariable):
* Source/WebGPU/WGSL/WGSLShaderModule.h:
(WGSL::ShaderModule::usesExternalTextures const):
(WGSL::ShaderModule::setUsesExternalTextures):
(WGSL::ShaderModule::clearUsesExternalTextures):

Canonical link: https://commits.webkit.org/264417@main