Google Git
Sign in
android / platform / external / skia / refs/heads/android15-tests-dev / . / src / sksl / codegen / SkSLWGSLCodeGenerator.cpp
blob: 023155595bc684bdbf1ead10b1f04ed09f98512a [file] [log] [blame] [edit]
/*
* Copyright 2022 Google Inc.
*
* Use of this source code is governed by a BSD-style license that can be
* found in the LICENSE file.
*/
#include "src/sksl/codegen/SkSLWGSLCodeGenerator.h"
#include "include/core/SkSpan.h"
#include "include/core/SkTypes.h"
#include "include/private/base/SkTArray.h"
#include "include/private/base/SkTo.h"
#include "src/base/SkEnumBitMask.h"
#include "src/base/SkStringView.h"
#include "src/core/SkTHash.h"
#include "src/core/SkTraceEvent.h"
#include "src/sksl/SkSLAnalysis.h"
#include "src/sksl/SkSLBuiltinTypes.h"
#include "src/sksl/SkSLCompiler.h"
#include "src/sksl/SkSLConstantFolder.h"
#include "src/sksl/SkSLContext.h"
#include "src/sksl/SkSLDefines.h"
#include "src/sksl/SkSLErrorReporter.h"
#include "src/sksl/SkSLIntrinsicList.h"
#include "src/sksl/SkSLMemoryLayout.h"
#include "src/sksl/SkSLOperator.h"
#include "src/sksl/SkSLOutputStream.h"
#include "src/sksl/SkSLPosition.h"
#include "src/sksl/SkSLProgramSettings.h"
#include "src/sksl/SkSLString.h"
#include "src/sksl/SkSLStringStream.h"
#include "src/sksl/SkSLUtil.h"
#include "src/sksl/analysis/SkSLProgramUsage.h"
#include "src/sksl/analysis/SkSLProgramVisitor.h"
#include "src/sksl/codegen/SkSLCodeGenerator.h"
#include "src/sksl/ir/SkSLBinaryExpression.h"
#include "src/sksl/ir/SkSLBlock.h"
#include "src/sksl/ir/SkSLConstructor.h"
#include "src/sksl/ir/SkSLConstructorArrayCast.h"
#include "src/sksl/ir/SkSLConstructorCompound.h"
#include "src/sksl/ir/SkSLConstructorDiagonalMatrix.h"
#include "src/sksl/ir/SkSLConstructorMatrixResize.h"
#include "src/sksl/ir/SkSLDoStatement.h"
#include "src/sksl/ir/SkSLExpression.h"
#include "src/sksl/ir/SkSLExpressionStatement.h"
#include "src/sksl/ir/SkSLFieldAccess.h"
#include "src/sksl/ir/SkSLForStatement.h"
#include "src/sksl/ir/SkSLFunctionCall.h"
#include "src/sksl/ir/SkSLFunctionDeclaration.h"
#include "src/sksl/ir/SkSLFunctionDefinition.h"
#include "src/sksl/ir/SkSLIRNode.h"
#include "src/sksl/ir/SkSLIfStatement.h"
#include "src/sksl/ir/SkSLIndexExpression.h"
#include "src/sksl/ir/SkSLInterfaceBlock.h"
#include "src/sksl/ir/SkSLLayout.h"
#include "src/sksl/ir/SkSLLiteral.h"
#include "src/sksl/ir/SkSLModifierFlags.h"
#include "src/sksl/ir/SkSLModifiersDeclaration.h"
#include "src/sksl/ir/SkSLPostfixExpression.h"
#include "src/sksl/ir/SkSLPrefixExpression.h"
#include "src/sksl/ir/SkSLProgram.h"
#include "src/sksl/ir/SkSLProgramElement.h"
#include "src/sksl/ir/SkSLReturnStatement.h"
#include "src/sksl/ir/SkSLSetting.h"
#include "src/sksl/ir/SkSLStatement.h"
#include "src/sksl/ir/SkSLStructDefinition.h"
#include "src/sksl/ir/SkSLSwitchCase.h"
#include "src/sksl/ir/SkSLSwitchStatement.h"
#include "src/sksl/ir/SkSLSwizzle.h"
#include "src/sksl/ir/SkSLTernaryExpression.h"
#include "src/sksl/ir/SkSLType.h"
#include "src/sksl/ir/SkSLVarDeclarations.h"
#include "src/sksl/ir/SkSLVariable.h"
#include "src/sksl/ir/SkSLVariableReference.h"
#include "src/sksl/spirv.h"
#include "src/sksl/transform/SkSLTransform.h"
#include <algorithm>
#include <cstddef>
#include <cstdint>
#include <initializer_list>
#include <iterator>
#include <memory>
#include <optional>
#include <string>
#include <string_view>
#include <utility>
#ifdef SK_ENABLE_WGSL_VALIDATION
#include "tint/tint.h"
#include "src/tint/lang/wgsl/reader/options.h"
#include "src/tint/lang/wgsl/extension.h"
#endif
using namespace skia_private;
namespace {
// Represents a function's dependencies that are not accessible in global scope. For instance,
// pipeline stage input and output parameters must be passed in as an argument.
//
// This is a bitmask enum. (It would be inside `class WGSLCodeGenerator`, but this leads to build
// errors in MSVC.)
enum class WGSLFunctionDependency : uint8_t {
kNone = 0,
kPipelineInputs = 1 << 0,
kPipelineOutputs = 1 << 1,
};
using WGSLFunctionDependencies = SkEnumBitMask<WGSLFunctionDependency>;
SK_MAKE_BITMASK_OPS(WGSLFunctionDependency)
} // namespace
namespace SkSL {
class WGSLCodeGenerator : public CodeGenerator {
public:
// See https://www.w3.org/TR/WGSL/#builtin-values
enum class Builtin {
// Vertex stage:
kVertexIndex, // input
kInstanceIndex, // input
kPosition, // output, fragment stage input
// Fragment stage:
kLastFragColor, // input
kFrontFacing, // input
kSampleIndex, // input
kFragDepth, // output
kSampleMaskIn, // input
kSampleMask, // output
// Compute stage:
kLocalInvocationId, // input
kLocalInvocationIndex, // input
kGlobalInvocationId, // input
kWorkgroupId, // input
kNumWorkgroups, // input
};
// Variable declarations can be terminated by:
// - comma (","), e.g. in struct member declarations or function parameters
// - semicolon (";"), e.g. in function scope variables
// A "none" option is provided to skip the delimiter when not needed, e.g. at the end of a list
// of declarations.
enum class Delimiter {
kComma,
kSemicolon,
kNone,
};
struct ProgramRequirements {
using DepsMap = skia_private::THashMap<const FunctionDeclaration*,
WGSLFunctionDependencies>;
// Mappings used to synthesize function parameters according to dependencies on pipeline
// input/output variables.
DepsMap fDependencies;
// These flags track extensions that will need to be enabled.
bool fPixelLocalExtension = false;
};
WGSLCodeGenerator(const Context* context,
const ShaderCaps* caps,
const Program* program,
OutputStream* out)
: INHERITED(context, caps, program, out) {}
bool generateCode() override;
private:
using INHERITED = CodeGenerator;
using Precedence = OperatorPrecedence;
// Called by generateCode() as the first step.
void preprocessProgram();
// Write output content while correctly handling indentation.
void write(std::string_view s);
void writeLine(std::string_view s = std::string_view());
void finishLine();
// Helpers to declare a pipeline stage IO parameter declaration.
void writePipelineIODeclaration(const Layout& layout,
const Type& type,
std::string_view name,
Delimiter delimiter);
void writeUserDefinedIODecl(const Layout& layout,
const Type& type,
std::string_view name,
Delimiter delimiter);
void writeBuiltinIODecl(const Type& type,
std::string_view name,
Builtin builtin,
Delimiter delimiter);
void writeVariableDecl(const Layout& layout,
const Type& type,
std::string_view name,
Delimiter delimiter);
// Write a function definition.
void writeFunction(const FunctionDefinition& f);
void writeFunctionDeclaration(const FunctionDeclaration& f,
SkSpan<const bool> paramNeedsDedicatedStorage);
// Write the program entry point.
void writeEntryPoint(const FunctionDefinition& f);
// Writers for supported statement types.
void writeStatement(const Statement& s);
void writeStatements(const StatementArray& statements);
void writeBlock(const Block& b);
void writeDoStatement(const DoStatement& expr);
void writeExpressionStatement(const Expression& expr);
void writeForStatement(const ForStatement& s);
void writeIfStatement(const IfStatement& s);
void writeReturnStatement(const ReturnStatement& s);
void writeSwitchStatement(const SwitchStatement& s);
void writeSwitchCases(SkSpan<const SwitchCase* const> cases);
void writeEmulatedSwitchFallthroughCases(SkSpan<const SwitchCase* const> cases,
std::string_view switchValue);
void writeSwitchCaseList(SkSpan<const SwitchCase* const> cases);
void writeVarDeclaration(const VarDeclaration& varDecl);
// Synthesizes an LValue for an expression.
class LValue;
class PointerLValue;
class SwizzleLValue;
class VectorComponentLValue;
std::unique_ptr<LValue> makeLValue(const Expression& e);
std::string variableReferenceNameForLValue(const VariableReference& r);
std::string variablePrefix(const Variable& v);
bool binaryOpNeedsComponentwiseMatrixPolyfill(const Type& left, const Type& right, Operator op);
// Writers for expressions. These return the final expression text as a string, and emit any
// necessary setup code directly into the program as necessary. The returned expression may be
// a `let`-alias that cannot be assigned-into; use `makeLValue` for an assignable expression.
std::string assembleExpression(const Expression& e, Precedence parentPrecedence);
std::string assembleBinaryExpression(const BinaryExpression& b, Precedence parentPrecedence);
std::string assembleBinaryExpression(const Expression& left,
Operator op,
const Expression& right,
const Type& resultType,
Precedence parentPrecedence);
std::string assembleFieldAccess(const FieldAccess& f);
std::string assembleFunctionCall(const FunctionCall& call, Precedence parentPrecedence);
std::string assembleIndexExpression(const IndexExpression& i);
std::string assembleLiteral(const Literal& l);
std::string assemblePostfixExpression(const PostfixExpression& p, Precedence parentPrecedence);
std::string assemblePrefixExpression(const PrefixExpression& p, Precedence parentPrecedence);
std::string assembleSwizzle(const Swizzle& swizzle);
std::string assembleTernaryExpression(const TernaryExpression& t, Precedence parentPrecedence);
std::string assembleVariableReference(const VariableReference& r);
std::string assembleName(std::string_view name);
std::string assembleIncrementExpr(const Type& type);
// Intrinsic helper functions.
std::string assembleIntrinsicCall(const FunctionCall& call,
IntrinsicKind kind,
Precedence parentPrecedence);
std::string assembleSimpleIntrinsic(std::string_view intrinsicName, const FunctionCall& call);
std::string assembleUnaryOpIntrinsic(Operator op,
const FunctionCall& call,
Precedence parentPrecedence);
std::string assembleBinaryOpIntrinsic(Operator op,
const FunctionCall& call,
Precedence parentPrecedence);
std::string assembleVectorizedIntrinsic(std::string_view intrinsicName,
const FunctionCall& call);
std::string assembleOutAssignedIntrinsic(std::string_view intrinsicName,
std::string_view returnFieldName,
std::string_view outFieldName,
const FunctionCall& call);
std::string assemblePartialSampleCall(std::string_view intrinsicName,
const Expression& sampler,
const Expression& coords);
std::string assembleInversePolyfill(const FunctionCall& call);
std::string assembleComponentwiseMatrixBinary(const Type& leftType,
const Type& rightType,
const std::string& left,
const std::string& right,
Operator op);
// Constructor expressions
std::string assembleAnyConstructor(const AnyConstructor& c);
std::string assembleConstructorCompound(const ConstructorCompound& c);
std::string assembleConstructorCompoundVector(const ConstructorCompound& c);
std::string assembleConstructorCompoundMatrix(const ConstructorCompound& c);
std::string assembleConstructorDiagonalMatrix(const ConstructorDiagonalMatrix& c);
std::string assembleConstructorMatrixResize(const ConstructorMatrixResize& ctor);
// Synthesized helper functions for comparison operators that are not supported by WGSL.
std::string assembleEqualityExpression(const Type& left,
const std::string& leftName,
const Type& right,
const std::string& rightName,
Operator op,
Precedence parentPrecedence);
std::string assembleEqualityExpression(const Expression& left,
const Expression& right,
Operator op,
Precedence parentPrecedence);
// Writes a scratch variable into the program and returns its name (e.g. `_skTemp123`).
std::string writeScratchVar(const Type& type, const std::string& value = "");
// Writes a scratch let-variable into the program, gives it the value of `expr`, and returns its
// name (e.g. `_skTemp123`).
std::string writeScratchLet(const std::string& expr);
std::string writeScratchLet(const Expression& expr, Precedence parentPrecedence);
// Converts `expr` into a string and returns a scratch let-variable associated with the
// expression. Compile-time constants and plain variable references will return the expression
// directly and omit the let-variable.
std::string writeNontrivialScratchLet(const Expression& expr, Precedence parentPrecedence);
// Generic recursive ProgramElement visitor.
void writeProgramElement(const ProgramElement& e);
void writeGlobalVarDeclaration(const GlobalVarDeclaration& d);
void writeStructDefinition(const StructDefinition& s);
void writeModifiersDeclaration(const ModifiersDeclaration&);
// Writes the WGSL struct fields for SkSL structs and interface blocks. Enforces WGSL address
// space layout constraints
// (https://www.w3.org/TR/WGSL/#address-space-layout-constraints) if a `layout` is
// provided. A struct that does not need to be host-shareable does not require a `layout`.
void writeFields(SkSpan<const Field> fields, const MemoryLayout* memoryLayout = nullptr);
// We bundle uniforms, and all varying pipeline stage inputs and outputs, into separate structs.
bool needsStageInputStruct() const;
void writeStageInputStruct();
bool needsStageOutputStruct() const;
void writeStageOutputStruct();
void writeUniformsAndBuffers();
void prepareUniformPolyfillsForInterfaceBlock(const InterfaceBlock* interfaceBlock,
std::string_view instanceName,
MemoryLayout::Standard nativeLayout);
void writeEnables();
void writeUniformPolyfills();
void writeTextureOrSampler(const Variable& var,
int bindingLocation,
std::string_view suffix,
std::string_view wgslType);
// Writes all top-level non-opaque global uniform declarations (i.e. not part of an interface
// block) into a single uniform block binding.
//
// In complete fragment/vertex/compute programs, uniforms will be declared only as interface
// blocks and global opaque types (like textures and samplers) which we expect to be declared
// with a unique binding and descriptor set index. However, test files that are declared as RTE
// programs may contain OpenGL-style global uniform declarations with no clear binding index to
// use for the containing synthesized block.
//
// Since we are handling these variables only to generate gold files from RTEs and never run
// them, we always declare them at the default bind group and binding index.
void writeNonBlockUniformsForTests();
// For a given function declaration, writes out any implicitly required pipeline stage arguments
// based on the function's pre-determined dependencies. These are expected to be written out as
// the first parameters for a function that requires them. Returns true if any arguments were
// written.
std::string functionDependencyArgs(const FunctionDeclaration&);
bool writeFunctionDependencyParams(const FunctionDeclaration&);
// Code in the header appears before the main body of code.
StringStream fHeader;
// We assign unique names to anonymous interface blocks based on the type.
skia_private::THashMap<const Type*, std::string> fInterfaceBlockNameMap;
// Stores the functions which use stage inputs/outputs as well as required WGSL extensions.
ProgramRequirements fRequirements;
skia_private::TArray<const Variable*> fPipelineInputs;
skia_private::TArray<const Variable*> fPipelineOutputs;
// These fields track whether we have written the polyfill for `inverse()` for a given matrix
// type.
bool fWrittenInverse2 = false;
bool fWrittenInverse3 = false;
bool fWrittenInverse4 = false;
// These fields control uniform polyfill support in cases where WGSL and std140 disagree.
// In std140 layout, matrices need to be represented as arrays of @size(16)-aligned vectors, and
// array elements are wrapped in a struct containing a single @size(16)-aligned element. Arrays
// of matrices combine both wrappers. These wrapper structs are unpacked into natively-typed
// globals at the shader entrypoint.
struct FieldPolyfillInfo {
const InterfaceBlock* fInterfaceBlock;
std::string fReplacementName;
bool fIsArray = false;
bool fIsMatrix = false;
bool fWasAccessed = false;
};
using FieldPolyfillMap = skia_private::THashMap<const Field*, FieldPolyfillInfo>;
FieldPolyfillMap fFieldPolyfillMap;
// Output processing state.
int fIndentation = 0;
bool fAtLineStart = false;
bool fHasUnconditionalReturn = false;
bool fAtFunctionScope = false;
int fConditionalScopeDepth = 0;
int fLocalSizeX = 1;
int fLocalSizeY = 1;
int fLocalSizeZ = 1;
int fScratchCount = 0;
};
enum class ProgramKind : int8_t;
namespace {
static constexpr char kSamplerSuffix[] = "_Sampler";
static constexpr char kTextureSuffix[] = "_Texture";
// See https://www.w3.org/TR/WGSL/#memory-view-types
enum class PtrAddressSpace {
kFunction,
kPrivate,
kStorage,
};
const char* operator_name(Operator op) {
switch (op.kind()) {
case Operator::Kind::LOGICALXOR: return " != ";
default: return op.operatorName();
}
}
bool is_reserved_word(std::string_view word) {
static const THashSet<std::string_view> kReservedWords{
// Used by SkSL:
"FSIn",
"FSOut",
"VSIn",
"VSOut",
"CSIn",
"_globalUniforms",
"_GlobalUniforms",
"_return",
"_stageIn",
"_stageOut",
// Keywords: https://www.w3.org/TR/WGSL/#keyword-summary
"alias",
"break",
"case",
"const",
"const_assert",
"continue",
"continuing",
"default",
"diagnostic",
"discard",
"else",
"enable",
"false",
"fn",
"for",
"if",
"let",
"loop",
"override",
"requires",
"return",
"struct",
"switch",
"true",
"var",
"while",
// Pre-declared types: https://www.w3.org/TR/WGSL/#predeclared-types
"bool",
"f16",
"f32",
"i32",
"u32",
// ... and pre-declared type generators:
"array",
"atomic",
"mat2x2",
"mat2x3",
"mat2x4",
"mat3x2",
"mat3x3",
"mat3x4",
"mat4x2",
"mat4x3",
"mat4x4",
"ptr",
"texture_1d",
"texture_2d",
"texture_2d_array",
"texture_3d",
"texture_cube",
"texture_cube_array",
"texture_multisampled_2d",
"texture_storage_1d",
"texture_storage_2d",
"texture_storage_2d_array",
"texture_storage_3d",
"vec2",
"vec3",
"vec4",
// Pre-declared enumerants: https://www.w3.org/TR/WGSL/#predeclared-enumerants
"read",
"write",
"read_write",
"function",
"private",
"workgroup",
"uniform",
"storage",
"perspective",
"linear",
"flat",
"center",
"centroid",
"sample",
"vertex_index",
"instance_index",
"position",
"front_facing",
"frag_depth",
"local_invocation_id",
"local_invocation_index",
"global_invocation_id",
"workgroup_id",
"num_workgroups",
"sample_index",
"sample_mask",
"rgba8unorm",
"rgba8snorm",
"rgba8uint",
"rgba8sint",
"rgba16uint",
"rgba16sint",
"rgba16float",
"r32uint",
"r32sint",
"r32float",
"rg32uint",
"rg32sint",
"rg32float",
"rgba32uint",
"rgba32sint",
"rgba32float",
"bgra8unorm",
// Reserved words: https://www.w3.org/TR/WGSL/#reserved-words
"_",
"NULL",
"Self",
"abstract",
"active",
"alignas",
"alignof",
"as",
"asm",
"asm_fragment",
"async",
"attribute",
"auto",
"await",
"become",
"binding_array",
"cast",
"catch",
"class",
"co_await",
"co_return",
"co_yield",
"coherent",
"column_major",
"common",
"compile",
"compile_fragment",
"concept",
"const_cast",
"consteval",
"constexpr",
"constinit",
"crate",
"debugger",
"decltype",
"delete",
"demote",
"demote_to_helper",
"do",
"dynamic_cast",
"enum",
"explicit",
"export",
"extends",
"extern",
"external",
"fallthrough",
"filter",
"final",
"finally",
"friend",
"from",
"fxgroup",
"get",
"goto",
"groupshared",
"highp",
"impl",
"implements",
"import",
"inline",
"instanceof",
"interface",
"layout",
"lowp",
"macro",
"macro_rules",
"match",
"mediump",
"meta",
"mod",
"module",
"move",
"mut",
"mutable",
"namespace",
"new",
"nil",
"noexcept",
"noinline",
"nointerpolation",
"noperspective",
"null",
"nullptr",
"of",
"operator",
"package",
"packoffset",
"partition",
"pass",
"patch",
"pixelfragment",
"precise",
"precision",
"premerge",
"priv",
"protected",
"pub",
"public",
"readonly",
"ref",
"regardless",
"register",
"reinterpret_cast",
"require",
"resource",
"restrict",
"self",
"set",
"shared",
"sizeof",
"smooth",
"snorm",
"static",
"static_assert",
"static_cast",
"std",
"subroutine",
"super",
"target",
"template",
"this",
"thread_local",
"throw",
"trait",
"try",
"type",
"typedef",
"typeid",
"typename",
"typeof",
"union",
"unless",
"unorm",
"unsafe",
"unsized",
"use",
"using",
"varying",
"virtual",
"volatile",
"wgsl",
"where",
"with",
"writeonly",
"yield",
};
return kReservedWords.contains(word);
}
std::string_view pipeline_struct_prefix(ProgramKind kind) {
if (ProgramConfig::IsVertex(kind)) {
return "VS";
}
if (ProgramConfig::IsFragment(kind)) {
return "FS";
}
if (ProgramConfig::IsCompute(kind)) {
return "CS";
}
// Compute programs don't have stage-in/stage-out pipeline structs.
return "";
}
std::string_view address_space_to_str(PtrAddressSpace addressSpace) {
switch (addressSpace) {
case PtrAddressSpace::kFunction:
return "function";
case PtrAddressSpace::kPrivate:
return "private";
case PtrAddressSpace::kStorage:
return "storage";
}
SkDEBUGFAIL("unsupported ptr address space");
return "unsupported";
}
std::string_view to_scalar_type(const Type& type) {
SkASSERT(type.typeKind() == Type::TypeKind::kScalar);
switch (type.numberKind()) {
// Floating-point numbers in WebGPU currently always have 32-bit footprint and
// relaxed-precision is not supported without extensions. f32 is the only floating-point
// number type in WGSL (see the discussion on https://github.com/gpuweb/gpuweb/issues/658).
case Type::NumberKind::kFloat:
return "f32";
case Type::NumberKind::kSigned:
return "i32";
case Type::NumberKind::kUnsigned:
return "u32";
case Type::NumberKind::kBoolean:
return "bool";
case Type::NumberKind::kNonnumeric:
[[fallthrough]];
default:
break;
}
return type.name();
}
// Convert a SkSL type to a WGSL type. Handles all plain types except structure types
// (see https://www.w3.org/TR/WGSL/#plain-types-section).
std::string to_wgsl_type(const Context& context, const Type& raw, const Layout* layout = nullptr) {
const Type& type = raw.resolve().scalarTypeForLiteral();
switch (type.typeKind()) {
case Type::TypeKind::kScalar:
return std::string(to_scalar_type(type));
case Type::TypeKind::kAtomic:
SkASSERT(type.matches(*context.fTypes.fAtomicUInt));
return "atomic<u32>";
case Type::TypeKind::kVector: {
std::string_view ct = to_scalar_type(type.componentType());
return String::printf("vec%d<%.*s>", type.columns(), (int)ct.length(), ct.data());
}
case Type::TypeKind::kMatrix: {
std::string_view ct = to_scalar_type(type.componentType());
return String::printf("mat%dx%d<%.*s>",
type.columns(), type.rows(), (int)ct.length(), ct.data());
}
case Type::TypeKind::kArray: {
std::string result = "array<" + to_wgsl_type(context, type.componentType(), layout);
if (!type.isUnsizedArray()) {
result += ", ";
result += std::to_string(type.columns());
}
return result + '>';
}
case Type::TypeKind::kTexture: {
if (type.matches(*context.fTypes.fWriteOnlyTexture2D)) {
std::string result = "texture_storage_2d<";
// Write-only storage texture types require a pixel format, which is in the layout.
SkASSERT(layout);
LayoutFlags pixelFormat = layout->fFlags & LayoutFlag::kAllPixelFormats;
switch (pixelFormat.value()) {
case (int)LayoutFlag::kRGBA8:
return result + "rgba8unorm, write>";
case (int)LayoutFlag::kRGBA32F:
return result + "rgba32float, write>";
case (int)LayoutFlag::kR32F:
return result + "r32float, write>";
default:
// The front-end should have rejected this.
return result + "write>";
}
}
if (type.matches(*context.fTypes.fReadOnlyTexture2D)) {
return "texture_2d<f32>";
}
break;
}
default:
break;
}
return std::string(type.name());
}
std::string to_ptr_type(const Context& context,
const Type& type,
const Layout* layout,
PtrAddressSpace addressSpace = PtrAddressSpace::kFunction) {
return "ptr<" + std::string(address_space_to_str(addressSpace)) + ", " +
to_wgsl_type(context, type, layout) + '>';
}
std::string_view wgsl_builtin_name(WGSLCodeGenerator::Builtin builtin) {
using Builtin = WGSLCodeGenerator::Builtin;
switch (builtin) {
case Builtin::kVertexIndex:
return "@builtin(vertex_index)";
case Builtin::kInstanceIndex:
return "@builtin(instance_index)";
case Builtin::kPosition:
return "@builtin(position)";
case Builtin::kLastFragColor:
return "@color(0)";
case Builtin::kFrontFacing:
return "@builtin(front_facing)";
case Builtin::kSampleIndex:
return "@builtin(sample_index)";
case Builtin::kFragDepth:
return "@builtin(frag_depth)";
case Builtin::kSampleMask:
case Builtin::kSampleMaskIn:
return "@builtin(sample_mask)";
case Builtin::kLocalInvocationId:
return "@builtin(local_invocation_id)";
case Builtin::kLocalInvocationIndex:
return "@builtin(local_invocation_index)";
case Builtin::kGlobalInvocationId:
return "@builtin(global_invocation_id)";
case Builtin::kWorkgroupId:
return "@builtin(workgroup_id)";
case Builtin::kNumWorkgroups:
return "@builtin(num_workgroups)";
default:
break;
}
SkDEBUGFAIL("unsupported builtin");
return "unsupported";
}
std::string_view wgsl_builtin_type(WGSLCodeGenerator::Builtin builtin) {
using Builtin = WGSLCodeGenerator::Builtin;
switch (builtin) {
case Builtin::kVertexIndex:
return "u32";
case Builtin::kInstanceIndex:
return "u32";
case Builtin::kPosition:
return "vec4<f32>";
case Builtin::kLastFragColor:
return "vec4<f32>";
case Builtin::kFrontFacing:
return "bool";
case Builtin::kSampleIndex:
return "u32";
case Builtin::kFragDepth:
return "f32";
case Builtin::kSampleMask:
return "u32";
case Builtin::kSampleMaskIn:
return "u32";
case Builtin::kLocalInvocationId:
return "vec3<u32>";
case Builtin::kLocalInvocationIndex:
return "u32";
case Builtin::kGlobalInvocationId:
return "vec3<u32>";
case Builtin::kWorkgroupId:
return "vec3<u32>";
case Builtin::kNumWorkgroups:
return "vec3<u32>";
default:
break;
}
SkDEBUGFAIL("unsupported builtin");
return "unsupported";
}
// Some built-in variables have a type that differs from their SkSL counterpart (e.g. signed vs
// unsigned integer). We handle these cases with an explicit type conversion during a variable
// reference. Returns the WGSL type of the conversion target if conversion is needed, otherwise
// returns std::nullopt.
std::optional<std::string_view> needs_builtin_type_conversion(const Variable& v) {
switch (v.layout().fBuiltin) {
case SK_VERTEXID_BUILTIN:
case SK_INSTANCEID_BUILTIN:
return {"i32"};
default:
break;
}
return std::nullopt;
}
// Map a SkSL builtin flag to a WGSL builtin kind. Returns std::nullopt if `builtin` is not
// not supported for WGSL.
//
// Also see //src/sksl/sksl_vert.sksl and //src/sksl/sksl_frag.sksl for supported built-ins.
std::optional<WGSLCodeGenerator::Builtin> builtin_from_sksl_name(int builtin) {
using Builtin = WGSLCodeGenerator::Builtin;
switch (builtin) {
case SK_POSITION_BUILTIN:
[[fallthrough]];
case SK_FRAGCOORD_BUILTIN:
return Builtin::kPosition;
case SK_VERTEXID_BUILTIN:
return Builtin::kVertexIndex;
case SK_INSTANCEID_BUILTIN:
return Builtin::kInstanceIndex;
case SK_LASTFRAGCOLOR_BUILTIN:
return Builtin::kLastFragColor;
case SK_CLOCKWISE_BUILTIN:
// TODO(skia:13092): While `front_facing` is the corresponding built-in, it does not
// imply a particular winding order. We correctly compute the face orientation based
// on how Skia configured the render pipeline for all references to this built-in
// variable (see `SkSL::Program::Interface::fRTFlipUniform`).
return Builtin::kFrontFacing;
case SK_SAMPLEMASKIN_BUILTIN:
return Builtin::kSampleMaskIn;
case SK_SAMPLEMASK_BUILTIN:
return Builtin::kSampleMask;
case SK_NUMWORKGROUPS_BUILTIN:
return Builtin::kNumWorkgroups;
case SK_WORKGROUPID_BUILTIN:
return Builtin::kWorkgroupId;
case SK_LOCALINVOCATIONID_BUILTIN:
return Builtin::kLocalInvocationId;
case SK_GLOBALINVOCATIONID_BUILTIN:
return Builtin::kGlobalInvocationId;
case SK_LOCALINVOCATIONINDEX_BUILTIN:
return Builtin::kLocalInvocationIndex;
default:
break;
}
return std::nullopt;
}
const char* delimiter_to_str(WGSLCodeGenerator::Delimiter delimiter) {
using Delim = WGSLCodeGenerator::Delimiter;
switch (delimiter) {
case Delim::kComma:
return ",";
case Delim::kSemicolon:
return ";";
case Delim::kNone:
default:
break;
}
return "";
}
// FunctionDependencyResolver visits the IR tree rooted at a particular function definition and
// computes that function's dependencies on pipeline stage IO parameters. These are later used to
// synthesize arguments when writing out function definitions.
class FunctionDependencyResolver : public ProgramVisitor {
public:
using Deps = WGSLFunctionDependencies;
using DepsMap = WGSLCodeGenerator::ProgramRequirements::DepsMap;
FunctionDependencyResolver(const Program* p,
const FunctionDeclaration* f,
DepsMap* programDependencyMap)
: fProgram(p), fFunction(f), fDependencyMap(programDependencyMap) {}
Deps resolve() {
fDeps = WGSLFunctionDependency::kNone;
this->visit(*fProgram);
return fDeps;
}
private:
bool visitProgramElement(const ProgramElement& p) override {
// Only visit the program that matches the requested function.
if (p.is<FunctionDefinition>() && &p.as<FunctionDefinition>().declaration() == fFunction) {
return INHERITED::visitProgramElement(p);
}
// Continue visiting other program elements.
return false;
}
bool visitExpression(const Expression& e) override {
if (e.is<VariableReference>()) {
const VariableReference& v = e.as<VariableReference>();
if (v.variable()->storage() == Variable::Storage::kGlobal) {
ModifierFlags flags = v.variable()->modifierFlags();
if (flags & ModifierFlag::kIn) {
fDeps |= WGSLFunctionDependency::kPipelineInputs;
}
if (flags & ModifierFlag::kOut) {
fDeps |= WGSLFunctionDependency::kPipelineOutputs;
}
}
} else if (e.is<FunctionCall>()) {
// The current function that we're processing (`fFunction`) inherits the dependencies of
// functions that it makes calls to, because the pipeline stage IO parameters need to be
// passed down as an argument.
const FunctionCall& callee = e.as<FunctionCall>();
// Don't process a function again if we have already resolved it.
Deps* found = fDependencyMap->find(&callee.function());
if (found) {
fDeps |= *found;
} else {
// Store the dependencies that have been discovered for the current function so far.
// If `callee` directly or indirectly calls the current function, then this value
// will prevent an infinite recursion.
fDependencyMap->set(fFunction, fDeps);
// Separately traverse the called function's definition and determine its
// dependencies.
FunctionDependencyResolver resolver(fProgram, &callee.function(), fDependencyMap);
Deps calleeDeps = resolver.resolve();
// Store the callee's dependencies in the global map to avoid processing
// the function again for future calls.
fDependencyMap->set(&callee.function(), calleeDeps);
// Add to the current function's dependencies.
fDeps |= calleeDeps;
}
}
return INHERITED::visitExpression(e);
}
const Program* const fProgram;
const FunctionDeclaration* const fFunction;
DepsMap* const fDependencyMap;
Deps fDeps = WGSLFunctionDependency::kNone;
using INHERITED = ProgramVisitor;
};
WGSLCodeGenerator::ProgramRequirements resolve_program_requirements(const Program* program) {
WGSLCodeGenerator::ProgramRequirements requirements;
for (const ProgramElement* e : program->elements()) {
switch (e->kind()) {
case ProgramElement::Kind::kFunction: {
const FunctionDeclaration& decl = e->as<FunctionDefinition>().declaration();
FunctionDependencyResolver resolver(program, &decl, &requirements.fDependencies);
requirements.fDependencies.set(&decl, resolver.resolve());
break;
}
case ProgramElement::Kind::kGlobalVar: {
const GlobalVarDeclaration& decl = e->as<GlobalVarDeclaration>();
if (decl.varDeclaration().var()->modifierFlags().isPixelLocal()) {
requirements.fPixelLocalExtension = true;
}
break;
}
default:
break;
}
}
return requirements;
}
void collect_pipeline_io_vars(const Program* program,
TArray<const Variable*>* ioVars,
ModifierFlag ioType) {
for (const ProgramElement* e : program->elements()) {
if (e->is<GlobalVarDeclaration>()) {
const Variable* v = e->as<GlobalVarDeclaration>().varDeclaration().var();
if (v->modifierFlags() & ioType) {
ioVars->push_back(v);
}
} else if (e->is<InterfaceBlock>()) {
const Variable* v = e->as<InterfaceBlock>().var();
if (v->modifierFlags() & ioType) {
ioVars->push_back(v);
}
}
}
}
bool is_in_global_uniforms(const Variable& var) {
SkASSERT(var.storage() == VariableStorage::kGlobal);
return var.modifierFlags().isUniform() &&
!var.type().isOpaque() &&
!var.interfaceBlock();
}
} // namespace
class WGSLCodeGenerator::LValue {
public:
virtual ~LValue() = default;
// Returns a WGSL expression that loads from the lvalue with no side effects.
// (e.g. `array[index].field`)
virtual std::string load() = 0;
// Returns a WGSL statement that stores into the lvalue with no side effects.
// (e.g. `array[index].field = the_passed_in_value_string;`)
virtual std::string store(const std::string& value) = 0;
};
class WGSLCodeGenerator::PointerLValue : public WGSLCodeGenerator::LValue {
public:
// `name` must be a WGSL expression with no side-effects, which we can safely take the address
// of. (e.g. `array[index].field` would be valid, but `array[Func()]` or `vector.x` are not.)
PointerLValue(std::string name) : fName(std::move(name)) {}
std::string load() override {
return fName;
}
std::string store(const std::string& value) override {
return fName + " = " + value + ";";
}
private:
std::string fName;
};
class WGSLCodeGenerator::VectorComponentLValue : public WGSLCodeGenerator::LValue {
public:
// `name` must be a WGSL expression with no side-effects that points to a single component of a
// WGSL vector.
VectorComponentLValue(std::string name) : fName(std::move(name)) {}
std::string load() override {
return fName;
}
std::string store(const std::string& value) override {
return fName + " = " + value + ";";
}
private:
std::string fName;
};
class WGSLCodeGenerator::SwizzleLValue : public WGSLCodeGenerator::LValue {
public:
// `name` must be a WGSL expression with no side-effects that points to a WGSL vector.
SwizzleLValue(const Context& ctx, std::string name, const Type& t, const ComponentArray& c)
: fContext(ctx)
, fName(std::move(name))
, fType(t)
, fComponents(c) {
// If the component array doesn't cover the entire value, we need to create masks for
// writing back into the lvalue. For example, if the type is vec4 and the component array
// holds `zx`, a GLSL assignment would look like:
// name.zx = new_value;
//
// The equivalent WGSL assignment statement would look like:
// name = vec4<f32>(new_value, name.xw).yzxw;
//
// This replaces name.zy with new_value.xy, and leaves name.xw at their original values.
// By convention, we always put the new value first and the original values second; it might
// be possible to find better arrangements which simplify the assignment overall, but we
// don't attempt this.
int fullSlotCount = fType.slotCount();
SkASSERT(fullSlotCount <= 4);
// First, see which components are used.
// The assignment swizzle must not reuse components.
bool used[4] = {};
for (int8_t component : fComponents) {
SkASSERT(!used[component]);
used[component] = true;
}
// Any untouched components will need to be fetched from the original value.
for (int index = 0; index < fullSlotCount; ++index) {
if (!used[index]) {
fUntouchedComponents.push_back(index);
}
}
// The reintegration swizzle needs to move the components back into their proper slots.
fReintegrationSwizzle.resize(fullSlotCount);
int reintegrateIndex = 0;
// This refills the untouched slots with the original values.
auto refillUntouchedSlots = [&] {
for (int index = 0; index < fullSlotCount; ++index) {
if (!used[index]) {
fReintegrationSwizzle[index] = reintegrateIndex++;
}
}
};
// This places the new-value components into the proper slots.
auto insertNewValuesIntoSlots = [&] {
for (int index = 0; index < fComponents.size(); ++index) {
fReintegrationSwizzle[fComponents[index]] = reintegrateIndex++;
}
};
// When reintegrating the untouched and new values, if the `x` slot is overwritten, we
// reintegrate the new value first. Otherwise, we reintegrate the original value first.
// This increases our odds of getting an identity swizzle for the reintegration.
if (used[0]) {
fReintegrateNewValueFirst = true;
insertNewValuesIntoSlots();
refillUntouchedSlots();
} else {
fReintegrateNewValueFirst = false;
refillUntouchedSlots();
insertNewValuesIntoSlots();
}
}
std::string load() override {
return fName + "." + Swizzle::MaskString(fComponents);
}
std::string store(const std::string& value) override {
// `variable = `
std::string result = fName;
result += " = ";
if (fUntouchedComponents.empty()) {
// `(new_value);`
result += '(';
result += value;
result += ")";
} else if (fReintegrateNewValueFirst) {
// `vec4<f32>((new_value), `
result += to_wgsl_type(fContext, fType);
result += "((";
result += value;
result += "), ";
// `variable.yz)`
result += fName;
result += '.';
result += Swizzle::MaskString(fUntouchedComponents);
result += ')';
} else {
// `vec4<f32>(variable.yz`
result += to_wgsl_type(fContext, fType);
result += '(';
result += fName;
result += '.';
result += Swizzle::MaskString(fUntouchedComponents);
// `, (new_value))`
result += ", (";
result += value;
result += "))";
}
if (!Swizzle::IsIdentity(fReintegrationSwizzle)) {
// `.wzyx`
result += '.';
result += Swizzle::MaskString(fReintegrationSwizzle);
}
return result + ';';
}
private:
const Context& fContext;
std::string fName;
const Type& fType;
ComponentArray fComponents;
ComponentArray fUntouchedComponents;
ComponentArray fReintegrationSwizzle;
bool fReintegrateNewValueFirst = false;
};
bool WGSLCodeGenerator::generateCode() {
// The resources of a WGSL program are structured in the following way:
// - Stage attribute inputs and outputs are bundled inside synthetic structs called
// VSIn/VSOut/FSIn/FSOut/CSIn.
// - All uniform and storage type resources are declared in global scope.
this->preprocessProgram();
{
AutoOutputStream outputToHeader(this, &fHeader, &fIndentation);
this->writeEnables();
this->writeStageInputStruct();
this->writeStageOutputStruct();
this->writeUniformsAndBuffers();
this->writeNonBlockUniformsForTests();
}
StringStream body;
{
// Emit the program body.
AutoOutputStream outputToBody(this, &body, &fIndentation);
const FunctionDefinition* mainFunc = nullptr;
for (const ProgramElement* e : fProgram.elements()) {
this->writeProgramElement(*e);
if (e->is<FunctionDefinition>()) {
const FunctionDefinition& func = e->as<FunctionDefinition>();
if (func.declaration().isMain()) {
mainFunc = &func;
}
}
}
// At the bottom of the program body, emit the entrypoint function.
// The entrypoint relies on state that has been collected while we emitted the rest of the
// program, so it's important to do it last to make sure we don't miss anything.
if (mainFunc) {
this->writeEntryPoint(*mainFunc);
}
}
write_stringstream(fHeader, *fOut);
write_stringstream(body, *fOut);
this->writeUniformPolyfills();
return fContext.fErrors->errorCount() == 0;
}
void WGSLCodeGenerator::writeUniformPolyfills() {
// If we didn't encounter any uniforms that need polyfilling, there is nothing to do.
if (fFieldPolyfillMap.empty()) {
return;
}
// We store the list of polyfilled fields as pointers in a hash-map, so the order can be
// inconsistent across runs. For determinism, we sort the polyfilled objects by name here.
TArray<const FieldPolyfillMap::Pair*> orderedFields;
orderedFields.reserve_exact(fFieldPolyfillMap.count());
fFieldPolyfillMap.foreach([&](const FieldPolyfillMap::Pair& pair) {
orderedFields.push_back(&pair);
});
std::sort(orderedFields.begin(),
orderedFields.end(),
[](const FieldPolyfillMap::Pair* a, const FieldPolyfillMap::Pair* b) {
return a->second.fReplacementName < b->second.fReplacementName;
});
THashSet<const Type*> writtenArrayElementPolyfill;
bool writtenUniformMatrixPolyfill[5][5] = {}; // m[column][row] for each matrix type
bool writtenUniformRowPolyfill[5] = {}; // for each matrix row-size
bool anyFieldAccessed = false;
for (const FieldPolyfillMap::Pair* pair : orderedFields) {
const auto& [field, info] = *pair;
const Type* fieldType = field->fType;
const Layout* fieldLayout = &field->fLayout;
if (info.fIsArray) {
fieldType = &fieldType->componentType();
if (!writtenArrayElementPolyfill.contains(fieldType)) {
writtenArrayElementPolyfill.add(fieldType);
this->write("struct _skArrayElement_");
this->write(fieldType->abbreviatedName());
this->writeLine(" {");
if (info.fIsMatrix) {
// Create a struct representing the array containing std140-padded matrices.
this->write(" e : _skMatrix");
this->write(std::to_string(fieldType->columns()));
this->writeLine(std::to_string(fieldType->rows()));
} else {
// Create a struct representing the array with extra padding between elements.
this->write(" @size(16) e : ");
this->writeLine(to_wgsl_type(fContext, *fieldType, fieldLayout));
}
this->writeLine("};");
}
}
if (info.fIsMatrix) {
// Create structs representing the matrix as an array of vectors, whether or not the
// matrix is ever accessed by the SkSL. (The struct itself is mentioned in the list of
// uniforms.)
int c = fieldType->columns();
int r = fieldType->rows();
if (!writtenUniformRowPolyfill[r]) {
writtenUniformRowPolyfill[r] = true;
this->write("struct _skRow");
this->write(std::to_string(r));
this->writeLine(" {");
this->write(" @size(16) r : vec");
this->write(std::to_string(r));
this->write("<");
this->write(to_wgsl_type(fContext, fieldType->componentType(), fieldLayout));
this->writeLine(">");
this->writeLine("};");
}
if (!writtenUniformMatrixPolyfill[c][r]) {
writtenUniformMatrixPolyfill[c][r] = true;
this->write("struct _skMatrix");
this->write(std::to_string(c));
this->write(std::to_string(r));
this->writeLine(" {");
this->write(" c : array<_skRow");
this->write(std::to_string(r));
this->write(", ");
this->write(std::to_string(c));
this->writeLine(">");
this->writeLine("};");
}
}
// We create a polyfill variable only if the uniform was actually accessed.
if (!info.fWasAccessed) {
continue;
}
anyFieldAccessed = true;
this->write("var<private> ");
this->write(info.fReplacementName);
this->write(": ");
const Type& interfaceBlockType = info.fInterfaceBlock->var()->type();
if (interfaceBlockType.isArray()) {
this->write("array<");
this->write(to_wgsl_type(fContext, *field->fType, fieldLayout));
this->write(", ");
this->write(std::to_string(interfaceBlockType.columns()));
this->write(">");
} else {
this->write(to_wgsl_type(fContext, *field->fType, fieldLayout));
}
this->writeLine(";");
}
// If no fields were actually accessed, _skInitializePolyfilledUniforms will not be called and
// we can avoid emitting an empty, dead function.
if (!anyFieldAccessed) {
return;
}
this->writeLine("fn _skInitializePolyfilledUniforms() {");
++fIndentation;
for (const FieldPolyfillMap::Pair* pair : orderedFields) {
// Only initialize a polyfill global if the uniform was actually accessed.
const auto& [field, info] = *pair;
if (!info.fWasAccessed) {
continue;
}
// Synthesize the name of this uniform variable
std::string_view instanceName = info.fInterfaceBlock->instanceName();
const Type& interfaceBlockType = info.fInterfaceBlock->var()->type();
if (instanceName.empty()) {
instanceName = fInterfaceBlockNameMap[&interfaceBlockType.componentType()];
}
// Initialize the global variable associated with this uniform.
// If the interface block is arrayed, the associated global will be arrayed as well.
int numIBElements = interfaceBlockType.isArray() ? interfaceBlockType.columns() : 1;
for (int ibIdx = 0; ibIdx < numIBElements; ++ibIdx) {
this->write(info.fReplacementName);
if (interfaceBlockType.isArray()) {
this->write("[");
this->write(std::to_string(ibIdx));
this->write("]");
}
this->write(" = ");
const Type* fieldType = field->fType;
const Layout* fieldLayout = &field->fLayout;
int numArrayElements;
if (info.fIsArray) {
this->write(to_wgsl_type(fContext, *fieldType, fieldLayout));
this->write("(");
numArrayElements = fieldType->columns();
fieldType = &fieldType->componentType();
} else {
numArrayElements = 1;
}
auto arraySeparator = String::Separator();
for (int arrayIdx = 0; arrayIdx < numArrayElements; arrayIdx++) {
this->write(arraySeparator());
std::string fieldName{instanceName};
if (interfaceBlockType.isArray()) {
fieldName += '[';
fieldName += std::to_string(ibIdx);
fieldName += ']';
}
fieldName += '.';
fieldName += this->assembleName(field->fName);
if (info.fIsArray) {
fieldName += '[';
fieldName += std::to_string(arrayIdx);
fieldName += "].e";
}
if (info.fIsMatrix) {
this->write(to_wgsl_type(fContext, *fieldType, fieldLayout));
this->write("(");
int numColumns = fieldType->columns();
auto matrixSeparator = String::Separator();
for (int column = 0; column < numColumns; column++) {
this->write(matrixSeparator());
this->write(fieldName);
this->write(".c[");
this->write(std::to_string(column));
this->write("].r");
}
this->write(")");
} else {
this->write(fieldName);
}
}
if (info.fIsArray) {
this->write(")");
}
this->writeLine(";");
}
}
--fIndentation;
this->writeLine("}");
}
void WGSLCodeGenerator::preprocessProgram() {
fRequirements = resolve_program_requirements(&fProgram);
collect_pipeline_io_vars(&fProgram, &fPipelineInputs, ModifierFlag::kIn);
collect_pipeline_io_vars(&fProgram, &fPipelineOutputs, ModifierFlag::kOut);
}
void WGSLCodeGenerator::write(std::string_view s) {
if (s.empty()) {
return;
}
#if defined(SK_DEBUG) || defined(SKSL_STANDALONE)
if (fAtLineStart) {
for (int i = 0; i < fIndentation; i++) {
fOut->writeText(" ");
}
}
#endif
fOut->writeText(std::string(s).c_str());
fAtLineStart = false;
}
void WGSLCodeGenerator::writeLine(std::string_view s) {
this->write(s);
fOut->writeText("\n");
fAtLineStart = true;
}
void WGSLCodeGenerator::finishLine() {
if (!fAtLineStart) {
this->writeLine();
}
}
std::string WGSLCodeGenerator::assembleName(std::string_view name) {
if (name.empty()) {
// WGSL doesn't allow anonymous function parameters.
return "_skAnonymous" + std::to_string(fScratchCount++);
}
// Add `R_` before reserved names to avoid any potential reserved-word conflict.
return (skstd::starts_with(name, "_sk") ||
skstd::starts_with(name, "R_") ||
is_reserved_word(name))
? std::string("R_") + std::string(name)
: std::string(name);
}
void WGSLCodeGenerator::writeVariableDecl(const Layout& layout,
const Type& type,
std::string_view name,
Delimiter delimiter) {
this->write(this->assembleName(name));
this->write(": " + to_wgsl_type(fContext, type, &layout));
this->writeLine(delimiter_to_str(delimiter));
}
void WGSLCodeGenerator::writePipelineIODeclaration(const Layout& layout,
const Type& type,
std::string_view name,
Delimiter delimiter) {
// In WGSL, an entry-point IO parameter is "one of either a built-in value or assigned a
// location". However, some SkSL declarations, specifically sk_FragColor, can contain both a
// location and a builtin modifier. In addition, WGSL doesn't have a built-in equivalent for
// sk_FragColor as it relies on the user-defined location for a render target.
//
// Instead of special-casing sk_FragColor, we just give higher precedence to a location modifier
// if a declaration happens to both have a location and it's a built-in.
//
// Also see:
// https://www.w3.org/TR/WGSL/#input-output-locations
// https://www.w3.org/TR/WGSL/#attribute-location
// https://www.w3.org/TR/WGSL/#builtin-inputs-outputs
if (layout.fLocation >= 0) {
this->writeUserDefinedIODecl(layout, type, name, delimiter);
return;
}
if (layout.fBuiltin >= 0) {
if (layout.fBuiltin == SK_POINTSIZE_BUILTIN) {
// WebGPU does not support the point-size builtin, but we silently replace it with a
// global variable when it is used, instead of reporting an error.
return;
}
auto builtin = builtin_from_sksl_name(layout.fBuiltin);
if (builtin.has_value()) {
this->writeBuiltinIODecl(type, name, *builtin, delimiter);
return;
}
}
fContext.fErrors->error(Position(), "declaration '" + std::string(name) + "' is not supported");
}
void WGSLCodeGenerator::writeUserDefinedIODecl(const Layout& layout,
const Type& type,
std::string_view name,
Delimiter delimiter) {
this->write("@location(" + std::to_string(layout.fLocation) + ") ");
// @blend_src is only allowed when doing dual-source blending, and only on color attachment 0.
if (layout.fLocation == 0 && layout.fIndex >= 0 && fProgram.fInterface.fOutputSecondaryColor) {
this->write("@blend_src(" + std::to_string(layout.fIndex) + ") ");
}
// "User-defined IO of scalar or vector integer type must always be specified as
// @interpolate(flat)" (see https://www.w3.org/TR/WGSL/#interpolation)
if (type.isInteger() || (type.isVector() && type.componentType().isInteger())) {
this->write("@interpolate(flat) ");
}
this->writeVariableDecl(layout, type, name, delimiter);
}
void WGSLCodeGenerator::writeBuiltinIODecl(const Type& type,
std::string_view name,
Builtin builtin,
Delimiter delimiter) {
this->write(wgsl_builtin_name(builtin));
this->write(" ");
this->write(this->assembleName(name));
this->write(": ");
this->write(wgsl_builtin_type(builtin));
this->writeLine(delimiter_to_str(delimiter));
}
void WGSLCodeGenerator::writeFunction(const FunctionDefinition& f) {
const FunctionDeclaration& decl = f.declaration();
fHasUnconditionalReturn = false;
fConditionalScopeDepth = 0;
SkASSERT(!fAtFunctionScope);
fAtFunctionScope = true;
// WGSL parameters are immutable and are considered as taking no storage, but SkSL parameters
// are real variables. To work around this, we make var-based copies of parameters. It's
// wasteful to make a copy of every single parameter--even if the compiler can eventually
// optimize them all away, that takes time and generates bloated code. So, we only make
// parameter copies if the variable is actually written-to.
STArray<32, bool> paramNeedsDedicatedStorage;
paramNeedsDedicatedStorage.push_back_n(decl.parameters().size(), true);
for (size_t index = 0; index < decl.parameters().size(); ++index) {
const Variable& param = *decl.parameters()[index];
if (param.type().isOpaque() || param.name().empty()) {
// Opaque-typed or anonymous parameters don't need dedicated storage.
paramNeedsDedicatedStorage[index] = false;
continue;
}
const ProgramUsage::VariableCounts counts = fProgram.fUsage->get(param);
if ((param.modifierFlags() & ModifierFlag::kOut) || counts.fWrite == 0) {
// Variables which are never written-to don't need dedicated storage.
// Out-parameters are passed as pointers; the pointer itself is never modified, so
// it doesn't need dedicated storage.
paramNeedsDedicatedStorage[index] = false;
}
}
this->writeFunctionDeclaration(decl, paramNeedsDedicatedStorage);
this->writeLine(" {");
++fIndentation;
// The parameters were given generic names like `_skParam1`, because WGSL parameters don't have
// storage and are immutable. If mutability is required, we create variables here; otherwise, we
// create properly-named `let` aliases.
for (size_t index = 0; index < decl.parameters().size(); ++index) {
if (paramNeedsDedicatedStorage[index]) {
const Variable& param = *decl.parameters()[index];
this->write("var ");
this->write(this->assembleName(param.mangledName()));
this->write(" = _skParam");
this->write(std::to_string(index));
this->writeLine(";");
}
}
this->writeBlock(f.body()->as<Block>());
// If fConditionalScopeDepth isn't zero, we have an unbalanced +1 or -1 when updating the depth.
SkASSERT(fConditionalScopeDepth == 0);
if (!fHasUnconditionalReturn && !f.declaration().returnType().isVoid()) {
this->write("return ");
this->write(to_wgsl_type(fContext, f.declaration().returnType()));
this->writeLine("();");
}
--fIndentation;
this->writeLine("}");
SkASSERT(fAtFunctionScope);
fAtFunctionScope = false;
}
void WGSLCodeGenerator::writeFunctionDeclaration(const FunctionDeclaration& decl,
SkSpan<const bool> paramNeedsDedicatedStorage) {
this->write("fn ");
if (decl.isMain()) {
this->write("_skslMain(");
} else {
this->write(this->assembleName(decl.mangledName()));
this->write("(");
}
auto separator = SkSL::String::Separator();
if (this->writeFunctionDependencyParams(decl)) {
separator(); // update the separator as parameters have been written
}
for (size_t index = 0; index < decl.parameters().size(); ++index) {
this->write(separator());
const Variable& param = *decl.parameters()[index];
if (param.type().isOpaque()) {
SkASSERT(!paramNeedsDedicatedStorage[index]);
if (param.type().isSampler()) {
// Create parameters for both the texture and associated sampler.
this->write(param.name());
this->write(kTextureSuffix);
this->write(": texture_2d<f32>, ");
this->write(param.name());
this->write(kSamplerSuffix);
this->write(": sampler");
} else {
// Create a parameter for the opaque object.
this->write(param.name());
this->write(": ");
this->write(to_wgsl_type(fContext, param.type(), &param.layout()));
}
} else {
if (paramNeedsDedicatedStorage[index] || param.name().empty()) {
// Create an unnamed parameter. If the parameter needs dedicated storage, it will
// later be assigned a `var` in the function body. (If it's anonymous, a var isn't
// needed.)
this->write("_skParam");
this->write(std::to_string(index));
} else {
// Use the name directly from the SkSL program.
this->write(this->assembleName(param.name()));
}
this->write(": ");
if (param.modifierFlags() & ModifierFlag::kOut) {
// Declare an "out" function parameter as a pointer.
this->write(to_ptr_type(fContext, param.type(), &param.layout()));
} else {
this->write(to_wgsl_type(fContext, param.type(), &param.layout()));
}
}
}
this->write(")");
if (!decl.returnType().isVoid()) {
this->write(" -> ");
this->write(to_wgsl_type(fContext, decl.returnType()));
}
}
void WGSLCodeGenerator::writeEntryPoint(const FunctionDefinition& main) {
SkASSERT(main.declaration().isMain());
const ProgramKind programKind = fProgram.fConfig->fKind;
#if defined(SKSL_STANDALONE)
if (ProgramConfig::IsRuntimeShader(programKind)) {
// Synthesize a basic entrypoint which just calls straight through to main.
// This is only used by skslc and just needs to pass the WGSL validator; Skia won't ever
// emit functions like this.
this->writeLine("@fragment fn main(@location(0) _coords: vec2<f32>) -> "
"@location(0) vec4<f32> {");
++fIndentation;
this->writeLine("return _skslMain(_coords);");
--fIndentation;
this->writeLine("}");
return;
}
#endif
// The input and output parameters for a vertex/fragment stage entry point function have the
// FSIn/FSOut/VSIn/VSOut/CSIn struct types that have been synthesized in generateCode(). An
// entrypoint always has a predictable signature and acts as a trampoline to the user-defined
// main function.
if (ProgramConfig::IsVertex(programKind)) {
this->write("@vertex");
} else if (ProgramConfig::IsFragment(programKind)) {
this->write("@fragment");
} else if (ProgramConfig::IsCompute(programKind)) {
this->write("@compute @workgroup_size(");
this->write(std::to_string(fLocalSizeX));
this->write(", ");
this->write(std::to_string(fLocalSizeY));
this->write(", ");
this->write(std::to_string(fLocalSizeZ));
this->write(")");
} else {
fContext.fErrors->error(Position(), "program kind not supported");
return;
}
this->write(" fn main(");
// The stage input struct is a parameter passed to main().
if (this->needsStageInputStruct()) {
this->write("_stageIn: ");
this->write(pipeline_struct_prefix(programKind));
this->write("In");
}
// The stage output struct is returned from main().
if (this->needsStageOutputStruct()) {
this->write(") -> ");
this->write(pipeline_struct_prefix(programKind));
this->writeLine("Out {");
} else {
this->writeLine(") {");
}
// Initialize polyfilled matrix uniforms if any were used.
fIndentation++;
for (const auto& [field, info] : fFieldPolyfillMap) {
if (info.fWasAccessed) {
this->writeLine("_skInitializePolyfilledUniforms();");
break;
}
}
// Declare the stage output struct.
if (this->needsStageOutputStruct()) {
this->write("var _stageOut: ");
this->write(pipeline_struct_prefix(programKind));
this->writeLine("Out;");
}
#if defined(SKSL_STANDALONE)
// We are compiling a Runtime Effect as a fragment shader, for testing purposes. We assign the
// result from _skslMain into sk_FragColor if the user-defined main returns a color. This
// doesn't actually matter, but it is more indicative of what a real program would do.
// `addImplicitFragColorWrite` from Transform::FindAndDeclareBuiltinVariables has already
// injected sk_FragColor into our stage outputs even if it wasn't explicitly referenced.
if (ProgramConfig::IsFragment(programKind)) {
if (main.declaration().returnType().matches(*fContext.fTypes.fHalf4)) {
this->write("_stageOut.sk_FragColor = ");
}
}
#endif
// Generate a function call to the user-defined main.
this->write("_skslMain(");
auto separator = SkSL::String::Separator();
WGSLFunctionDependencies* deps = fRequirements.fDependencies.find(&main.declaration());
if (deps) {
if (*deps & WGSLFunctionDependency::kPipelineInputs) {
this->write(separator());
this->write("_stageIn");
}
if (*deps & WGSLFunctionDependency::kPipelineOutputs) {
this->write(separator());
this->write("&_stageOut");
}
}
#if defined(SKSL_STANDALONE)
if (const Variable* v = main.declaration().getMainCoordsParameter()) {
// We are compiling a Runtime Effect as a fragment shader, for testing purposes.
// We need to synthesize a coordinates parameter, but the coordinates don't matter.
SkASSERT(ProgramConfig::IsFragment(programKind));
const Type& type = v->type();
if (!type.matches(*fContext.fTypes.fFloat2)) {
fContext.fErrors->error(main.fPosition, "main function has unsupported parameter: " +
type.description());
return;
}
this->write(separator());
this->write("/*fragcoord*/ vec2<f32>()");
}
#endif
this->writeLine(");");
if (this->needsStageOutputStruct()) {
// Return the stage output struct.
this->writeLine("return _stageOut;");
}
fIndentation--;
this->writeLine("}");
}
void WGSLCodeGenerator::writeStatement(const Statement& s) {
switch (s.kind()) {
case Statement::Kind::kBlock:
this->writeBlock(s.as<Block>());
break;
case Statement::Kind::kBreak:
this->writeLine("break;");
break;
case Statement::Kind::kContinue:
this->writeLine("continue;");
break;
case Statement::Kind::kDiscard:
this->writeLine("discard;");
break;
case Statement::Kind::kDo:
this->writeDoStatement(s.as<DoStatement>());
break;
case Statement::Kind::kExpression:
this->writeExpressionStatement(*s.as<ExpressionStatement>().expression());
break;
case Statement::Kind::kFor:
this->writeForStatement(s.as<ForStatement>());
break;
case Statement::Kind::kIf:
this->writeIfStatement(s.as<IfStatement>());
break;
case Statement::Kind::kNop:
this->writeLine(";");
break;
case Statement::Kind::kReturn:
this->writeReturnStatement(s.as<ReturnStatement>());
break;
case Statement::Kind::kSwitch:
this->writeSwitchStatement(s.as<SwitchStatement>());
break;
case Statement::Kind::kSwitchCase:
SkDEBUGFAIL("switch-case statements should only be present inside a switch");
break;
case Statement::Kind::kVarDeclaration:
this->writeVarDeclaration(s.as<VarDeclaration>());
break;
}
}
void WGSLCodeGenerator::writeStatements(const StatementArray& statements) {
for (const auto& s : statements) {
if (!s->isEmpty()) {
this->writeStatement(*s);
this->finishLine();
}
}
}
void WGSLCodeGenerator::writeBlock(const Block& b) {
// Write scope markers if this block is a scope, or if the block is empty (since we need to emit
// something here to make the code valid).
bool isScope = b.isScope() || b.isEmpty();
if (isScope) {
this->writeLine("{");
fIndentation++;
}
this->writeStatements(b.children());
if (isScope) {
fIndentation--;
this->writeLine("}");
}
}
void WGSLCodeGenerator::writeExpressionStatement(const Expression& expr) {
// Any expression-related side effects must be emitted as separate statements when
// `assembleExpression` is called.
// The final result of the expression will be a variable, let-reference, or an expression with
// no side effects (`foo + bar`). Discarding this result is safe, as the program never uses it.
(void)this->assembleExpression(expr, Precedence::kStatement);
}
void WGSLCodeGenerator::writeDoStatement(const DoStatement& s) {
// Generate a loop structure like this:
// loop {
// body-statement;
// continuing {
// break if inverted-test-expression;
// }
// }
++fConditionalScopeDepth;
std::unique_ptr<Expression> invertedTestExpr = PrefixExpression::Make(
fContext, s.test()->fPosition, OperatorKind::LOGICALNOT, s.test()->clone());
this->writeLine("loop {");
fIndentation++;
this->writeStatement(*s.statement());
this->finishLine();
this->writeLine("continuing {");
fIndentation++;
std::string breakIfExpr = this->assembleExpression(*invertedTestExpr, Precedence::kExpression);
this->write("break if ");
this->write(breakIfExpr);
this->writeLine(";");
fIndentation--;
this->writeLine("}");
fIndentation--;
this->writeLine("}");
--fConditionalScopeDepth;
}
void WGSLCodeGenerator::writeForStatement(const ForStatement& s) {
// Generate a loop structure wrapped in an extra scope:
// {
// initializer-statement;
// loop;
// }
// The outer scope is necessary to prevent the initializer-variable from leaking out into the
// rest of the code. In practice, the generated code actually tends to be scoped even more
// deeply, as the body-statement almost always contributes an extra block.
++fConditionalScopeDepth;
if (s.initializer()) {
this->writeLine("{");
fIndentation++;
this->writeStatement(*s.initializer());
this->writeLine();
}
this->writeLine("loop {");
fIndentation++;
if (s.unrollInfo()) {
if (s.unrollInfo()->fCount <= 0) {
// Loops which are known to never execute don't need to be emitted at all.
// (However, the front end should have already replaced this loop with a Nop.)
} else {
// Loops which are known to execute at least once can use this form:
//
// loop {
// body-statement;
// continuing {
// next-expression;
// break if inverted-test-expression;
// }
// }
this->writeStatement(*s.statement());
this->finishLine();
this->writeLine("continuing {");
++fIndentation;
if (s.next()) {
this->writeExpressionStatement(*s.next());
this->finishLine();
}
if (s.test()) {
std::unique_ptr<Expression> invertedTestExpr = PrefixExpression::Make(
fContext, s.test()->fPosition, OperatorKind::LOGICALNOT, s.test()->clone());
std::string breakIfExpr =
this->assembleExpression(*invertedTestExpr, Precedence::kExpression);
this->write("break if ");
this->write(breakIfExpr);
this->writeLine(";");
}
--fIndentation;
this->writeLine("}");
}
} else {
// Loops without a known execution count are emitted in this form:
//
// loop {
// if test-expression {
// body-statement;
// } else {
// break;
// }
// continuing {
// next-expression;
// }
// }
if (s.test()) {
std::string testExpr = this->assembleExpression(*s.test(), Precedence::kExpression);
this->write("if ");
this->write(testExpr);
this->writeLine(" {");
fIndentation++;
this->writeStatement(*s.statement());
this->finishLine();
fIndentation--;
this->writeLine("} else {");
fIndentation++;
this->writeLine("break;");
fIndentation--;
this->writeLine("}");
}
else {
this->writeStatement(*s.statement());
this->finishLine();
}
if (s.next()) {
this->writeLine("continuing {");
fIndentation++;
this->writeExpressionStatement(*s.next());
this->finishLine();
fIndentation--;
this->writeLine("}");
}
}
// This matches an open-brace at the top of the loop.
fIndentation--;
this->writeLine("}");
if (s.initializer()) {
// This matches an open-brace before the initializer-statement.
fIndentation--;
this->writeLine("}");
}
--fConditionalScopeDepth;
}
void WGSLCodeGenerator::writeIfStatement(const IfStatement& s) {
++fConditionalScopeDepth;
std::string testExpr = this->assembleExpression(*s.test(), Precedence::kExpression);
this->write("if ");
this->write(testExpr);
this->writeLine(" {");
fIndentation++;
this->writeStatement(*s.ifTrue());
this->finishLine();
fIndentation--;
if (s.ifFalse()) {
this->writeLine("} else {");
fIndentation++;
this->writeStatement(*s.ifFalse());
this->finishLine();
fIndentation--;
}
this->writeLine("}");
--fConditionalScopeDepth;
}
void WGSLCodeGenerator::writeReturnStatement(const ReturnStatement& s) {
fHasUnconditionalReturn |= (fConditionalScopeDepth == 0);
std::string expr = s.expression()
? this->assembleExpression(*s.expression(), Precedence::kExpression)
: std::string();
this->write("return ");
this->write(expr);
this->write(";");
}
void WGSLCodeGenerator::writeSwitchCaseList(SkSpan<const SwitchCase* const> cases) {
auto separator = SkSL::String::Separator();
for (const SwitchCase* const sc : cases) {
this->write(separator());
if (sc->isDefault()) {
this->write("default");
} else {
this->write(std::to_string(sc->value()));
}
}
}
void WGSLCodeGenerator::writeSwitchCases(SkSpan<const SwitchCase* const> cases) {
if (!cases.empty()) {
// Only the last switch-case should have a non-empty statement.
SkASSERT(std::all_of(cases.begin(), std::prev(cases.end()), [](const SwitchCase* sc) {
return sc->statement()->isEmpty();
}));
// Emit the cases in a comma-separated list.
this->write("case ");
this->writeSwitchCaseList(cases);
this->writeLine(" {");
++fIndentation;
// Emit the switch-case body.
this->writeStatement(*cases.back()->statement());
this->finishLine();
--fIndentation;
this->writeLine("}");
}
}
void WGSLCodeGenerator::writeEmulatedSwitchFallthroughCases(SkSpan<const SwitchCase* const> cases,
std::string_view switchValue) {
// There's no need for fallthrough handling unless we actually have multiple case blocks.
if (cases.size() < 2) {
this->writeSwitchCases(cases);
return;
}
// Match against the entire case group.
this->write("case ");
this->writeSwitchCaseList(cases);
this->writeLine(" {");
++fIndentation;
std::string fallthroughVar = this->writeScratchVar(*fContext.fTypes.fBool, "false");
const size_t secondToLastCaseIndex = cases.size() - 2;
const size_t lastCaseIndex = cases.size() - 1;
for (size_t index = 0; index < cases.size(); ++index) {
const SwitchCase& sc = *cases[index];
if (index < lastCaseIndex) {
// The default case must come last in SkSL, and this case isn't the last one, so it
// can't possibly be the default.
SkASSERT(!sc.isDefault());
this->write("if ");
if (index > 0) {
this->write(fallthroughVar);
this->write(" || ");
}
this->write(switchValue);
this->write(" == ");
this->write(std::to_string(sc.value()));
this->writeLine(" {");
fIndentation++;
// We write the entire case-block statement here, and then set `switchFallthrough`
// to 1. If the case-block had a break statement in it, we break out of the outer
// for-loop entirely, meaning the `switchFallthrough` assignment never occurs, nor
// does any code after it inside the switch. We've forbidden `continue` statements
// inside switch case-blocks entirely, so we don't need to consider their effect on
// control flow; see the Finalizer in FunctionDefinition::Convert.
this->writeStatement(*sc.statement());
this->finishLine();
if (index < secondToLastCaseIndex) {
// Set a variable to indicate falling through to the next block. The very last
// case-block is reached by process of elimination and doesn't need this
// variable, so we don't actually need to set it if we are on the second-to-last
// case block.
this->write(fallthroughVar);
this->write(" = true; ");
}
this->writeLine("// fallthrough");
fIndentation--;
this->writeLine("}");
} else {
// This is the final case. Since it's always last, we can just dump in the code.
// (If we didn't match any of the other values, we must have matched this one by
// process of elimination. If we did match one of the other values, we either hit a
// `break` statement earlier--and won't get this far--or we're falling through.)
this->writeStatement(*sc.statement());
this->finishLine();
}
}
--fIndentation;
this->writeLine("}");
}
void WGSLCodeGenerator::writeSwitchStatement(const SwitchStatement& s) {
// WGSL supports the `switch` statement in a limited capacity. A default case must always be
// specified. Each switch-case must be scoped inside braces. Fallthrough is not supported; a
// trailing break is implied at the end of each switch-case block. (Explicit breaks are also
// allowed.) One minor improvement over a traditional switch is that switch-cases take a list
// of values to match, instead of a single value:
// case 1, 2 { foo(); }
// case 3, default { bar(); }
//
// We will use the native WGSL switch statement for any switch-cases in the SkSL which can be
// made to conform to these limitations. The remaining cases which cannot conform will be
// emulated with if-else blocks (similar to our GLSL ES2 switch-statement emulation path). This
// should give us good performance in the common case, since most switches naturally conform.
// First, let's emit the switch itself.
std::string valueExpr = this->writeNontrivialScratchLet(*s.value(), Precedence::kExpression);
this->write("switch ");
this->write(valueExpr);
this->writeLine(" {");
++fIndentation;
// Now let's go through the switch-cases, and emit the ones that don't fall through.
TArray<const SwitchCase*> nativeCases;
TArray<const SwitchCase*> fallthroughCases;
bool previousCaseFellThrough = false;
bool foundNativeDefault = false;
[[maybe_unused]] bool foundFallthroughDefault = false;
const int lastSwitchCaseIdx = s.cases().size() - 1;
for (int index = 0; index <= lastSwitchCaseIdx; ++index) {
const SwitchCase& sc = s.cases()[index]->as<SwitchCase>();
if (sc.statement()->isEmpty()) {
// This is a `case X:` that immediately falls through to the next case.
// If we aren't already falling through, we can handle this via a comma-separated list.
if (previousCaseFellThrough) {
fallthroughCases.push_back(&sc);
foundFallthroughDefault |= sc.isDefault();
} else {
nativeCases.push_back(&sc);
foundNativeDefault |= sc.isDefault();
}
continue;
}
if (index == lastSwitchCaseIdx || Analysis::SwitchCaseContainsUnconditionalExit(sc)) {
// This is a `case X:` that never falls through.
if (previousCaseFellThrough) {
// Because the previous cases fell through, we can't use a native switch-case here.
fallthroughCases.push_back(&sc);
foundFallthroughDefault |= sc.isDefault();
this->writeEmulatedSwitchFallthroughCases(fallthroughCases, valueExpr);
fallthroughCases.clear();
// Fortunately, we're no longer falling through blocks, so we might be able to use a
// native switch-case list again.
previousCaseFellThrough = false;
} else {
// Emit a native switch-case block with a comma-separated case list.
nativeCases.push_back(&sc);
foundNativeDefault |= sc.isDefault();
this->writeSwitchCases(nativeCases);
nativeCases.clear();
}
continue;
}
// This case falls through, so it will need to be handled via emulation.
// If we have put together a collection of "native" cases (cases that fall through with no
// actual case-body), we will need to slide them over into the fallthrough-case list.
fallthroughCases.push_back_n(nativeCases.size(), nativeCases.data());
nativeCases.clear();
fallthroughCases.push_back(&sc);
foundFallthroughDefault |= sc.isDefault();
previousCaseFellThrough = true;
}
// Finish out the remaining switch-cases.
this->writeSwitchCases(nativeCases);
nativeCases.clear();
this->writeEmulatedSwitchFallthroughCases(fallthroughCases, valueExpr);
fallthroughCases.clear();
// WGSL requires a default case.
if (!foundNativeDefault && !foundFallthroughDefault) {
this->writeLine("case default {}");
}
--fIndentation;
this->writeLine("}");
}
void WGSLCodeGenerator::writeVarDeclaration(const VarDeclaration& varDecl) {
std::string initialValue =
varDecl.value() ? this->assembleExpression(*varDecl.value(), Precedence::kAssignment)
: std::string();
if (varDecl.var()->modifierFlags().isConst()) {
// Use `const` at global scope, or if the value is a compile-time constant.
SkASSERTF(varDecl.value(), "a constant variable must specify a value");
this->write((!fAtFunctionScope || Analysis::IsCompileTimeConstant(*varDecl.value()))
? "const "
: "let ");
} else {
this->write("var ");
}
this->write(this->assembleName(varDecl.var()->mangledName()));
this->write(": ");
this->write(to_wgsl_type(fContext, varDecl.var()->type(), &varDecl.var()->layout()));
if (varDecl.value()) {
this->write(" = ");
this->write(initialValue);
}
this->write(";");
}
std::unique_ptr<WGSLCodeGenerator::LValue> WGSLCodeGenerator::makeLValue(const Expression& e) {
if (e.is<VariableReference>()) {
return std::make_unique<PointerLValue>(
this->variableReferenceNameForLValue(e.as<VariableReference>()));
}
if (e.is<FieldAccess>()) {
return std::make_unique<PointerLValue>(this->assembleFieldAccess(e.as<FieldAccess>()));
}
if (e.is<IndexExpression>()) {
const IndexExpression& idx = e.as<IndexExpression>();
if (idx.base()->type().isVector()) {
// Rewrite indexed-swizzle accesses like `myVec.zyx[i]` into an index onto `myVec`.
if (std::unique_ptr<Expression> rewrite =
Transform::RewriteIndexedSwizzle(fContext, idx)) {
return std::make_unique<VectorComponentLValue>(
this->assembleExpression(*rewrite, Precedence::kAssignment));
} else {
return std::make_unique<VectorComponentLValue>(this->assembleIndexExpression(idx));
}
} else {
return std::make_unique<PointerLValue>(this->assembleIndexExpression(idx));
}
}
if (e.is<Swizzle>()) {
const Swizzle& swizzle = e.as<Swizzle>();
if (swizzle.components().size() == 1) {
return std::make_unique<VectorComponentLValue>(this->assembleSwizzle(swizzle));
} else {
return std::make_unique<SwizzleLValue>(
fContext,
this->assembleExpression(*swizzle.base(), Precedence::kAssignment),
swizzle.base()->type(),
swizzle.components());
}
}
fContext.fErrors->error(e.fPosition, "unsupported lvalue type");
return nullptr;
}
std::string WGSLCodeGenerator::assembleExpression(const Expression& e,
Precedence parentPrecedence) {
switch (e.kind()) {
case Expression::Kind::kBinary:
return this->assembleBinaryExpression(e.as<BinaryExpression>(), parentPrecedence);
case Expression::Kind::kConstructorCompound:
return this->assembleConstructorCompound(e.as<ConstructorCompound>());
case Expression::Kind::kConstructorArrayCast:
// This is a no-op, since WGSL 1.0 doesn't have any concept of precision qualifiers.
// When we add support for f16, this will need to copy the array contents.
return this->assembleExpression(*e.as<ConstructorArrayCast>().argument(),
parentPrecedence);
case Expression::Kind::kConstructorArray:
case Expression::Kind::kConstructorCompoundCast:
case Expression::Kind::kConstructorScalarCast:
case Expression::Kind::kConstructorSplat:
case Expression::Kind::kConstructorStruct:
return this->assembleAnyConstructor(e.asAnyConstructor());
case Expression::Kind::kConstructorDiagonalMatrix:
return this->assembleConstructorDiagonalMatrix(e.as<ConstructorDiagonalMatrix>());
case Expression::Kind::kConstructorMatrixResize:
return this->assembleConstructorMatrixResize(e.as<ConstructorMatrixResize>());
case Expression::Kind::kEmpty:
return "false";
case Expression::Kind::kFieldAccess:
return this->assembleFieldAccess(e.as<FieldAccess>());
case Expression::Kind::kFunctionCall:
return this->assembleFunctionCall(e.as<FunctionCall>(), parentPrecedence);
case Expression::Kind::kIndex:
return this->assembleIndexExpression(e.as<IndexExpression>());
case Expression::Kind::kLiteral:
return this->assembleLiteral(e.as<Literal>());
case Expression::Kind::kPrefix:
return this->assemblePrefixExpression(e.as<PrefixExpression>(), parentPrecedence);
case Expression::Kind::kPostfix:
return this->assemblePostfixExpression(e.as<PostfixExpression>(), parentPrecedence);
case Expression::Kind::kSetting:
return this->assembleExpression(*e.as<Setting>().toLiteral(fCaps), parentPrecedence);
case Expression::Kind::kSwizzle:
return this->assembleSwizzle(e.as<Swizzle>());
case Expression::Kind::kTernary:
return this->assembleTernaryExpression(e.as<TernaryExpression>(), parentPrecedence);
case Expression::Kind::kVariableReference:
return this->assembleVariableReference(e.as<VariableReference>());
default:
SkDEBUGFAILF("unsupported expression:\n%s", e.description().c_str());
return {};
}
}
static bool is_nontrivial_expression(const Expression& expr) {
// We consider a "trivial expression" one which we can repeat multiple times in the output
// without being dangerous or spammy. We avoid emitting temporary variables for very trivial
// expressions: literals, unadorned variable references, or constant vectors.
if (expr.is<VariableReference>() || expr.is<Literal>()) {
// Variables and literals are trivial; adding a let-declaration won't simplify anything.
return false;
}
if (expr.type().isVector() && Analysis::IsConstantExpression(expr)) {
// Compile-time constant vectors are also considered trivial; they're short and sweet.
return false;
}
return true;
}
static bool binary_op_is_ambiguous_in_wgsl(Operator op) {
// WGSL always requires parentheses for some operators which are deemed to be ambiguous.
// (8.19. Operator Precedence and Associativity)
switch (op.kind()) {
case OperatorKind::LOGICALOR:
case OperatorKind::LOGICALAND:
case OperatorKind::BITWISEOR:
case OperatorKind::BITWISEAND:
case OperatorKind::BITWISEXOR:
case OperatorKind::SHL:
case OperatorKind::SHR:
case OperatorKind::LT:
case OperatorKind::GT:
case OperatorKind::LTEQ:
case OperatorKind::GTEQ:
return true;
default:
return false;
}
}
bool WGSLCodeGenerator::binaryOpNeedsComponentwiseMatrixPolyfill(const Type& left,
const Type& right,
Operator op) {
switch (op.kind()) {
case OperatorKind::SLASH:
// WGSL does not natively support componentwise matrix-op-matrix for division.
if (left.isMatrix() && right.isMatrix()) {
return true;
}
[[fallthrough]];
case OperatorKind::PLUS:
case OperatorKind::MINUS:
// WGSL does not natively support componentwise matrix-op-scalar or scalar-op-matrix for
// addition, subtraction or division.
return (left.isMatrix() && right.isScalar()) ||
(left.isScalar() && right.isMatrix());
default:
return false;
}
}
std::string WGSLCodeGenerator::assembleBinaryExpression(const BinaryExpression& b,
Precedence parentPrecedence) {
return this->assembleBinaryExpression(*b.left(), b.getOperator(), *b.right(), b.type(),
parentPrecedence);
}
std::string WGSLCodeGenerator::assembleBinaryExpression(const Expression& left,
Operator op,
const Expression& right,
const Type& resultType,
Precedence parentPrecedence) {
// If the operator is && or ||, we need to handle short-circuiting properly. Specifically, we
// sometimes need to emit extra statements to paper over functionality that WGSL lacks, like
// assignment in the middle of an expression. We need to guard those extra statements, to ensure
// that they don't occur if the expression evaluation is short-circuited. Converting the
// expression into an if-else block keeps the short-circuit property intact even when extra
// statements are involved.
// If the RHS doesn't have any side effects, then it's safe to just leave the expression as-is,
// since we know that any possible extra statements are non-side-effecting.
std::string expr;
if (op.kind() == OperatorKind::LOGICALAND && Analysis::HasSideEffects(right)) {
// Converts `left_expression && right_expression` into the following block:
// var _skTemp1: bool;
// [[ prepare left_expression ]]
// if left_expression {
// [[ prepare right_expression ]]
// _skTemp1 = right_expression;
// } else {
// _skTemp1 = false;
// }
expr = this->writeScratchVar(resultType);
std::string leftExpr = this->assembleExpression(left, Precedence::kExpression);
this->write("if ");
this->write(leftExpr);
this->writeLine(" {");
++fIndentation;
std::string rightExpr = this->assembleExpression(right, Precedence::kAssignment);
this->write(expr);
this->write(" = ");
this->write(rightExpr);
this->writeLine(";");
--fIndentation;
this->writeLine("} else {");
++fIndentation;
this->write(expr);
this->writeLine(" = false;");
--fIndentation;
this->writeLine("}");
return expr;
}
if (op.kind() == OperatorKind::LOGICALOR && Analysis::HasSideEffects(right)) {
// Converts `left_expression || right_expression` into the following block:
// var _skTemp1: bool;
// [[ prepare left_expression ]]
// if left_expression {
// _skTemp1 = true;
// } else {
// [[ prepare right_expression ]]
// _skTemp1 = right_expression;
// }
expr = this->writeScratchVar(resultType);
std::string leftExpr = this->assembleExpression(left, Precedence::kExpression);
this->write("if ");
this->write(leftExpr);
this->writeLine(" {");
++fIndentation;
this->write(expr);
this->writeLine(" = true;");
--fIndentation;
this->writeLine("} else {");
++fIndentation;
std::string rightExpr = this->assembleExpression(right, Precedence::kAssignment);
this->write(expr);
this->write(" = ");
this->write(rightExpr);
this->writeLine(";");
--fIndentation;
this->writeLine("}");
return expr;
}
// Handle comma-expressions.
if (op.kind() == OperatorKind::COMMA) {
// The result from the left-expression is ignored, but its side effects must occur.
this->assembleExpression(left, Precedence::kStatement);
// Evaluate the right side normally.
return this->assembleExpression(right, parentPrecedence);
}
// Handle assignment-expressions.
if (op.isAssignment()) {
std::unique_ptr<LValue> lvalue = this->makeLValue(left);
if (!lvalue) {
return "";
}
if (op.kind() == OperatorKind::EQ) {
// Evaluate the right-hand side of simple assignment (`a = b` --> `b`).
expr = this->assembleExpression(right, Precedence::kAssignment);
} else {
// Evaluate the right-hand side of compound-assignment (`a += b` --> `a + b`).
op = op.removeAssignment();
std::string lhs = lvalue->load();
std::string rhs = this->assembleExpression(right, op.getBinaryPrecedence());
if (this->binaryOpNeedsComponentwiseMatrixPolyfill(left.type(), right.type(), op)) {
if (is_nontrivial_expression(right)) {
rhs = this->writeScratchLet(rhs);
}
expr = this->assembleComponentwiseMatrixBinary(left.type(), right.type(),
lhs, rhs, op);
} else {
expr = lhs + operator_name(op) + rhs;
}
}
// Emit the assignment statement (`a = a + b`).
this->writeLine(lvalue->store(expr));
// Return the lvalue (`a`) as the result, since the value might be used by the caller.
return lvalue->load();
}
if (op.isEquality()) {
return this->assembleEqualityExpression(left, right, op, parentPrecedence);
}
Precedence precedence = op.getBinaryPrecedence();
bool needParens = precedence >= parentPrecedence;
if (binary_op_is_ambiguous_in_wgsl(op)) {
precedence = Precedence::kParentheses;
}
if (needParens) {
expr = "(";
}
// If we are emitting `constant + constant`, this generally indicates that the values could not
// be constant-folded. This happens when the values overflow or become nan. WGSL will refuse to
// compile such expressions, as WGSL 1.0 has no infinity/nan support. However, the WGSL
// compile-time check can be dodged by putting one side into a let-variable. This technically
// gives us an indeterminate result, but the vast majority of backends will just calculate an
// infinity or nan here, as we would expect. (skia:14385)
bool bothSidesConstant = ConstantFolder::GetConstantValueOrNull(left) &&
ConstantFolder::GetConstantValueOrNull(right);
std::string lhs = this->assembleExpression(left, precedence);
std::string rhs = this->assembleExpression(right, precedence);
if (this->binaryOpNeedsComponentwiseMatrixPolyfill(left.type(), right.type(), op)) {
if (bothSidesConstant || is_nontrivial_expression(left)) {
lhs = this->writeScratchLet(lhs);
}
if (is_nontrivial_expression(right)) {
rhs = this->writeScratchLet(rhs);
}
expr += this->assembleComponentwiseMatrixBinary(left.type(), right.type(), lhs, rhs, op);
} else {
if (bothSidesConstant) {
lhs = this->writeScratchLet(lhs);
}
expr += lhs + operator_name(op) + rhs;
}
if (needParens) {
expr += ')';
}
return expr;
}
std::string WGSLCodeGenerator::assembleFieldAccess(const FieldAccess& f) {
const Field* field = &f.base()->type().fields()[f.fieldIndex()];
std::string expr;
if (FieldPolyfillInfo* polyfillInfo = fFieldPolyfillMap.find(field)) {
// We found a matrix uniform. We are required to pass some matrix uniforms as array vectors,
// since the std140 layout for a matrix assumes 4-column vectors for each row, and WGSL
// tightly packs 2-column matrices. When emitting code, we replace the field-access
// expression with a global variable which holds an unpacked version of the uniform.
polyfillInfo->fWasAccessed = true;
// The polyfill can either be based directly onto a uniform in an interface block, or it
// might be based on an index-expression onto a uniform if the interface block is arrayed.
const Expression* base = f.base().get();
const IndexExpression* indexExpr = nullptr;
if (base->is<IndexExpression>()) {
indexExpr = &base->as<IndexExpression>();
base = indexExpr->base().get();
}
SkASSERT(base->is<VariableReference>());
expr = polyfillInfo->fReplacementName;
// If we had an index expression, we must append the index.
if (indexExpr) {
expr += '[';
expr += this->assembleExpression(*indexExpr->index(), Precedence::kSequence);
expr += ']';
}
return expr;
}
switch (f.ownerKind()) {
case FieldAccess::OwnerKind::kDefault:
expr = this->assembleExpression(*f.base(), Precedence::kPostfix) + '.';
break;
case FieldAccess::OwnerKind::kAnonymousInterfaceBlock:
if (f.base()->is<VariableReference>() &&
field->fLayout.fBuiltin != SK_POINTSIZE_BUILTIN) {
expr = this->variablePrefix(*f.base()->as<VariableReference>().variable());
}
break;
}
expr += this->assembleName(field->fName);
return expr;
}
static bool all_arguments_constant(const ExpressionArray& arguments) {
// Returns true if all arguments in the ExpressionArray are compile-time constants. If we are
// calling an intrinsic and all of its inputs are constant, but we didn't constant-fold it, this
// generally indicates that constant-folding resulted in an infinity or nan. The WGSL compiler
// will reject such an expression with a compile-time error. We can dodge the error, taking on
// the risk of indeterminate behavior instead, by replacing one of the constant values with a
// scratch let-variable. (skia:14385)
for (const std::unique_ptr<Expression>& arg : arguments) {
if (!ConstantFolder::GetConstantValueOrNull(*arg)) {
return false;
}
}
return true;
}
std::string WGSLCodeGenerator::assembleSimpleIntrinsic(std::string_view intrinsicName,
const FunctionCall& call) {
// Invoke the function, passing each function argument.
std::string expr = std::string(intrinsicName);
expr.push_back('(');
const ExpressionArray& args = call.arguments();
auto separator = SkSL::String::Separator();
bool allConstant = all_arguments_constant(call.arguments());
for (int index = 0; index < args.size(); ++index) {
expr += separator();
std::string argument = this->assembleExpression(*args[index], Precedence::kSequence);
if (args[index]->type().isAtomic()) {
// WGSL passes atomic values to intrinsics as pointers.
expr += '&';
expr += argument;
} else if (allConstant && index == 0) {
// We can use a scratch-let for argument 0 to dodge WGSL overflow errors. (skia:14385)
expr += this->writeScratchLet(argument);
} else {
expr += argument;
}
}
expr.push_back(')');
if (call.type().isVoid()) {
this->write(expr);
this->writeLine(";");
return std::string();
} else {
return this->writeScratchLet(expr);
}
}
std::string WGSLCodeGenerator::assembleVectorizedIntrinsic(std::string_view intrinsicName,
const FunctionCall& call) {
SkASSERT(!call.type().isVoid());
// Invoke the function, passing each function argument.
std::string expr = std::string(intrinsicName);
expr.push_back('(');
auto separator = SkSL::String::Separator();
const ExpressionArray& args = call.arguments();
bool returnsVector = call.type().isVector();
bool allConstant = all_arguments_constant(call.arguments());
for (int index = 0; index < args.size(); ++index) {
expr += separator();
bool vectorize = returnsVector && args[index]->type().isScalar();
if (vectorize) {
expr += to_wgsl_type(fContext, call.type());
expr.push_back('(');
}
// We can use a scratch-let for argument 0 to dodge WGSL overflow errors. (skia:14385)
std::string argument = this->assembleExpression(*args[index], Precedence::kSequence);
expr += (allConstant && index == 0) ? this->writeScratchLet(argument)
: argument;
if (vectorize) {
expr.push_back(')');
}
}
expr.push_back(')');
return this->writeScratchLet(expr);
}
std::string WGSLCodeGenerator::assembleUnaryOpIntrinsic(Operator op,
const FunctionCall& call,
Precedence parentPrecedence) {
SkASSERT(!call.type().isVoid());
bool needParens = Precedence::kPrefix >= parentPrecedence;
std::string expr;
if (needParens) {
expr.push_back('(');
}
expr += operator_name(op);
expr += this->assembleExpression(*call.arguments()[0], Precedence::kPrefix);
if (needParens) {
expr.push_back(')');
}
return expr;
}
std::string WGSLCodeGenerator::assembleBinaryOpIntrinsic(Operator op,
const FunctionCall& call,
Precedence parentPrecedence) {
SkASSERT(!call.type().isVoid());
Precedence precedence = op.getBinaryPrecedence();
bool needParens = precedence >= parentPrecedence ||
binary_op_is_ambiguous_in_wgsl(op);
std::string expr;
if (needParens) {
expr.push_back('(');
}
// We can use a scratch-let for argument 0 to dodge WGSL overflow errors. (skia:14385)
std::string argument = this->assembleExpression(*call.arguments()[0], precedence);
expr += all_arguments_constant(call.arguments()) ? this->writeScratchLet(argument)
: argument;
expr += operator_name(op);
expr += this->assembleExpression(*call.arguments()[1], precedence);
if (needParens) {
expr.push_back(')');
}
return expr;
}
// Rewrite a WGSL intrinsic of the form "intrinsicName(in) -> struct" to the SkSL's
// "intrinsicName(in, outField) -> returnField", where outField and returnField are the names of the
// fields in the struct returned by the WGSL intrinsic.
std::string WGSLCodeGenerator::assembleOutAssignedIntrinsic(std::string_view intrinsicName,
std::string_view returnField,
std::string_view outField,
const FunctionCall& call) {
SkASSERT(call.type().componentType().isNumber());
SkASSERT(call.arguments().size() == 2);
SkASSERT(call.function().parameters()[1]->modifierFlags() & ModifierFlag::kOut);
std::string expr = std::string(intrinsicName);
expr += "(";
// Invoke the intrinsic with the first parameter. Use a scratch-let if argument is a constant
// to dodge WGSL overflow errors. (skia:14385)
std::string argument = this->assembleExpression(*call.arguments()[0], Precedence::kSequence);
expr += ConstantFolder::GetConstantValueOrNull(*call.arguments()[0])
? this->writeScratchLet(argument) : argument;
expr += ")";
// In WGSL the intrinsic returns a struct; assign it to a local so that its fields can be
// accessed multiple times.
expr = this->writeScratchLet(expr);
expr += ".";
// Store the outField of `expr` to the intended "out" argument
std::unique_ptr<LValue> lvalue = this->makeLValue(*call.arguments()[1]);
if (!lvalue) {
return "";
}
std::string outValue = expr;
outValue += outField;
this->writeLine(lvalue->store(outValue));
// And return the expression accessing the returnField.
expr += returnField;
return expr;
}
std::string WGSLCodeGenerator::assemblePartialSampleCall(std::string_view functionName,
const Expression& sampler,
const Expression& coords) {
// This function returns `functionName(inSampler_texture, inSampler_sampler, coords` without a
// terminating comma or close-parenthesis. This allows the caller to add more arguments as
// needed.
SkASSERT(sampler.type().typeKind() == Type::TypeKind::kSampler);
std::string expr = std::string(functionName) + '(';
expr += this->assembleExpression(sampler, Precedence::kSequence);
expr += kTextureSuffix;
expr += ", ";
expr += this->assembleExpression(sampler, Precedence::kSequence);
expr += kSamplerSuffix;
expr += ", ";
// Compute the sample coordinates, dividing out the Z if a vec3 was provided.
SkASSERT(coords.type().isVector());
if (coords.type().columns() == 3) {
// The coordinates were passed as a vec3, so we need to emit `coords.xy / coords.z`.
std::string vec3Coords = this->writeScratchLet(coords, Precedence::kMultiplicative);
expr += vec3Coords + ".xy / " + vec3Coords + ".z";
} else {
// The coordinates should be a plain vec2; emit the expression as-is.
SkASSERT(coords.type().columns() == 2);
expr += this->assembleExpression(coords, Precedence::kSequence);
}
return expr;
}
std::string WGSLCodeGenerator::assembleComponentwiseMatrixBinary(const Type& leftType,
const Type& rightType,
const std::string& left,
const std::string& right,
Operator op) {
bool leftIsMatrix = leftType.isMatrix();
bool rightIsMatrix = rightType.isMatrix();
const Type& matrixType = leftIsMatrix ? leftType : rightType;
std::string expr = to_wgsl_type(fContext, matrixType) + '(';
auto separator = String::Separator();
int columns = matrixType.columns();
for (int c = 0; c < columns; ++c) {
expr += separator();
expr += left;
if (leftIsMatrix) {
expr += '[';
expr += std::to_string(c);
expr += ']';
}
expr += op.operatorName();
expr += right;
if (rightIsMatrix) {
expr += '[';
expr += std::to_string(c);
expr += ']';
}
}
return expr + ')';
}
std::string WGSLCodeGenerator::assembleIntrinsicCall(const FunctionCall& call,
IntrinsicKind kind,
Precedence parentPrecedence) {
// Be careful: WGSL 1.0 will reject any intrinsic calls which can be constant-evaluated to
// infinity or nan with a compile error. If all arguments to an intrinsic are compile-time
// constants (`all_arguments_constant`), it is safest to copy one argument into a scratch-let so
// that the call will be seen as runtime-evaluated, which defuses the overflow checks.
// Don't worry; a competent driver should still optimize it away.
const ExpressionArray& arguments = call.arguments();
switch (kind) {
case k_atan_IntrinsicKind: {
const char* name = (arguments.size() == 1) ? "atan" : "atan2";
return this->assembleSimpleIntrinsic(name, call);
}
case k_dFdx_IntrinsicKind:
return this->assembleSimpleIntrinsic("dpdx", call);
case k_dFdy_IntrinsicKind:
// TODO(b/294274678): apply RTFlip here
return this->assembleSimpleIntrinsic("dpdy", call);
case k_dot_IntrinsicKind: {
if (arguments[0]->type().isScalar()) {
return this->assembleBinaryOpIntrinsic(OperatorKind::STAR, call, parentPrecedence);
}
return this->assembleSimpleIntrinsic("dot", call);
}
case k_equal_IntrinsicKind:
return this->assembleBinaryOpIntrinsic(OperatorKind::EQEQ, call, parentPrecedence);
case k_faceforward_IntrinsicKind: {
if (arguments[0]->type().isScalar()) {
// select(-N, N, (I * Nref) < 0)
std::string N = this->writeNontrivialScratchLet(*arguments[0],
Precedence::kAssignment);
return this->writeScratchLet(
"select(-" + N + ", " + N + ", " +
this->assembleBinaryExpression(*arguments[1],
OperatorKind::STAR,
*arguments[2],
arguments[1]->type(),
Precedence::kRelational) +
" < 0)");
}
return this->assembleSimpleIntrinsic("faceForward", call);
}
case k_frexp_IntrinsicKind:
// SkSL frexp is "$genType fract = frexp($genType, out $genIType exp)" whereas WGSL
// returns a struct with no out param: "let [fract, exp] = frexp($genType)".
return this->assembleOutAssignedIntrinsic("frexp", "fract", "exp", call);
case k_greaterThan_IntrinsicKind:
return this->assembleBinaryOpIntrinsic(OperatorKind::GT, call, parentPrecedence);
case k_greaterThanEqual_IntrinsicKind:
return this->assembleBinaryOpIntrinsic(OperatorKind::GTEQ, call, parentPrecedence);
case k_inverse_IntrinsicKind:
return this->assembleInversePolyfill(call);
case k_inversesqrt_IntrinsicKind:
return this->assembleSimpleIntrinsic("inverseSqrt", call);
case k_lessThan_IntrinsicKind:
return this->assembleBinaryOpIntrinsic(OperatorKind::LT, call, parentPrecedence);
case k_lessThanEqual_IntrinsicKind:
return this->assembleBinaryOpIntrinsic(OperatorKind::LTEQ, call, parentPrecedence);
case k_matrixCompMult_IntrinsicKind: {
// We use a scratch-let for arg0 to avoid the potential for WGSL overflow. (skia:14385)
std::string arg0 = all_arguments_constant(arguments)
? this->writeScratchLet(*arguments[0], Precedence::kPostfix)
: this->writeNontrivialScratchLet(*arguments[0], Precedence::kPostfix);
std::string arg1 = this->writeNontrivialScratchLet(*arguments[1], Precedence::kPostfix);
return this->writeScratchLet(
this->assembleComponentwiseMatrixBinary(arguments[0]->type(),
arguments[1]->type(),
arg0,
arg1,
OperatorKind::STAR));
}
case k_mix_IntrinsicKind: {
const char* name = arguments[2]->type().componentType().isBoolean() ? "select" : "mix";
return this->assembleVectorizedIntrinsic(name, call);
}
case k_mod_IntrinsicKind: {
// WGSL has no intrinsic equivalent to `mod`. Synthesize `x - y * floor(x / y)`.
// We can use a scratch-let on one side to dodge WGSL overflow errors. In practice, I
// can't find any values of x or y which would overflow, but it can't hurt. (skia:14385)
std::string arg0 = all_arguments_constant(arguments)
? this->writeScratchLet(*arguments[0], Precedence::kAdditive)
: this->writeNontrivialScratchLet(*arguments[0], Precedence::kAdditive);
std::string arg1 = this->writeNontrivialScratchLet(*arguments[1],
Precedence::kAdditive);
return this->writeScratchLet(arg0 + " - " + arg1 + " * floor(" +
arg0 + " / " + arg1 + ")");
}
case k_modf_IntrinsicKind:
// SkSL modf is "$genType fract = modf($genType, out $genType whole)" whereas WGSL
// returns a struct with no out param: "let [fract, whole] = modf($genType)".
return this->assembleOutAssignedIntrinsic("modf", "fract", "whole", call);
case k_normalize_IntrinsicKind: {
const char* name = arguments[0]->type().isScalar() ? "sign" : "normalize";
return this->assembleSimpleIntrinsic(name, call);
}
case k_not_IntrinsicKind:
return this->assembleUnaryOpIntrinsic(OperatorKind::LOGICALNOT, call, parentPrecedence);
case k_notEqual_IntrinsicKind:
return this->assembleBinaryOpIntrinsic(OperatorKind::NEQ, call, parentPrecedence);
case k_packHalf2x16_IntrinsicKind:
return this->assembleSimpleIntrinsic("pack2x16float", call);
case k_packSnorm2x16_IntrinsicKind:
return this->assembleSimpleIntrinsic("pack2x16snorm", call);
case k_packSnorm4x8_IntrinsicKind:
return this->assembleSimpleIntrinsic("pack4x8snorm", call);
case k_packUnorm2x16_IntrinsicKind:
return this->assembleSimpleIntrinsic("pack2x16unorm", call);
case k_packUnorm4x8_IntrinsicKind:
return this->assembleSimpleIntrinsic("pack4x8unorm", call);
case k_reflect_IntrinsicKind:
if (arguments[0]->type().isScalar()) {
// I - 2 * N * I * N
// We can use a scratch-let for N to dodge WGSL overflow errors. (skia:14385)
std::string I = this->writeNontrivialScratchLet(*arguments[0],
Precedence::kAdditive);
std::string N = all_arguments_constant(arguments)
? this->writeScratchLet(*arguments[1], Precedence::kMultiplicative)
: this->writeNontrivialScratchLet(*arguments[1], Precedence::kMultiplicative);
return this->writeScratchLet(String::printf("%s - 2 * %s * %s * %s",
I.c_str(), N.c_str(),
I.c_str(), N.c_str()));
}
return this->assembleSimpleIntrinsic("reflect", call);
case k_refract_IntrinsicKind:
if (arguments[0]->type().isScalar()) {
// WGSL only implements refract for vectors; rather than reimplementing refract from
// scratch, we can replace the call with `refract(float2(I,0), float2(N,0), eta).x`.
std::string I = this->writeNontrivialScratchLet(*arguments[0],
Precedence::kSequence);
std::string N = this->writeNontrivialScratchLet(*arguments[1],
Precedence::kSequence);
// We can use a scratch-let for Eta to avoid WGSL overflow errors. (skia:14385)
std::string Eta = all_arguments_constant(arguments)
? this->writeScratchLet(*arguments[2], Precedence::kSequence)
: this->writeNontrivialScratchLet(*arguments[2], Precedence::kSequence);
return this->writeScratchLet(
String::printf("refract(vec2<%s>(%s, 0), vec2<%s>(%s, 0), %s).x",
to_wgsl_type(fContext, arguments[0]->type()).c_str(),
I.c_str(),
to_wgsl_type(fContext, arguments[1]->type()).c_str(),
N.c_str(),
Eta.c_str()));
}
return this->assembleSimpleIntrinsic("refract", call);
case k_sample_IntrinsicKind: {
// Determine if a bias argument was passed in.
SkASSERT(arguments.size() == 2 || arguments.size() == 3);
bool callIncludesBias = (arguments.size() == 3);
if (fProgram.fConfig->fSettings.fSharpenTextures || callIncludesBias) {
// We need to supply a bias argument; this is a separate intrinsic in WGSL.
std::string expr = this->assemblePartialSampleCall("textureSampleBias",
*arguments[0],
*arguments[1]);
expr += ", ";
if (callIncludesBias) {
expr += this->assembleExpression(*arguments[2], Precedence::kAdditive) +
" + ";
}
expr += skstd::to_string(fProgram.fConfig->fSettings.fSharpenTextures
? kSharpenTexturesBias
: 0.0f);
return expr + ')';
}
// No bias is necessary, so we can call `textureSample` directly.
return this->assemblePartialSampleCall("textureSample",
*arguments[0],
*arguments[1]) + ')';
}
case k_sampleLod_IntrinsicKind: {
std::string expr = this->assemblePartialSampleCall("textureSampleLevel",
*arguments[0],
*arguments[1]);
expr += ", " + this->assembleExpression(*arguments[2], Precedence::kSequence);
return expr + ')';
}
case k_sampleGrad_IntrinsicKind: {
std::string expr = this->assemblePartialSampleCall("textureSampleGrad",
*arguments[0],
*arguments[1]);
expr += ", " + this->assembleExpression(*arguments[2], Precedence::kSequence);
expr += ", " + this->assembleExpression(*arguments[3], Precedence::kSequence);
return expr + ')';
}
case k_textureHeight_IntrinsicKind:
return this->assembleSimpleIntrinsic("textureDimensions", call) + ".y";
case k_textureRead_IntrinsicKind: {
// We need to inject an extra argument for the mip-level. We don't plan on using mipmaps
// in our storage textures, so we can just pass zero.
std::string tex = this->assembleExpression(*arguments[0], Precedence::kSequence);
std::string pos = this->writeScratchLet(*arguments[1], Precedence::kSequence);
return std::string("textureLoad(") + tex + ", " + pos + ", 0)";
}
case k_textureWidth_IntrinsicKind:
return this->assembleSimpleIntrinsic("textureDimensions", call) + ".x";
case k_textureWrite_IntrinsicKind:
return this->assembleSimpleIntrinsic("textureStore", call);
case k_unpackHalf2x16_IntrinsicKind:
return this->assembleSimpleIntrinsic("unpack2x16float", call);
case k_unpackSnorm2x16_IntrinsicKind:
return this->assembleSimpleIntrinsic("unpack2x16snorm", call);
case k_unpackSnorm4x8_IntrinsicKind:
return this->assembleSimpleIntrinsic("unpack4x8snorm", call);
case k_unpackUnorm2x16_IntrinsicKind:
return this->assembleSimpleIntrinsic("unpack2x16unorm", call);
case k_unpackUnorm4x8_IntrinsicKind:
return this->assembleSimpleIntrinsic("unpack4x8unorm", call);
case k_clamp_IntrinsicKind:
case k_max_IntrinsicKind:
case k_min_IntrinsicKind:
case k_smoothstep_IntrinsicKind:
case k_step_IntrinsicKind:
return this->assembleVectorizedIntrinsic(call.function().name(), call);
case k_abs_IntrinsicKind:
case k_acos_IntrinsicKind:
case k_all_IntrinsicKind:
case k_any_IntrinsicKind:
case k_asin_IntrinsicKind:
case k_atomicAdd_IntrinsicKind:
case k_atomicLoad_IntrinsicKind:
case k_atomicStore_IntrinsicKind:
case k_ceil_IntrinsicKind:
case k_cos_IntrinsicKind:
case k_cross_IntrinsicKind:
case k_degrees_IntrinsicKind:
case k_distance_IntrinsicKind:
case k_exp_IntrinsicKind:
case k_exp2_IntrinsicKind:
case k_floor_IntrinsicKind:
case k_fract_IntrinsicKind:
case k_length_IntrinsicKind:
case k_log_IntrinsicKind:
case k_log2_IntrinsicKind:
case k_radians_IntrinsicKind:
case k_pow_IntrinsicKind:
case k_saturate_IntrinsicKind:
case k_sign_IntrinsicKind:
case k_sin_IntrinsicKind:
case k_sqrt_IntrinsicKind:
case k_storageBarrier_IntrinsicKind:
case k_tan_IntrinsicKind:
case k_workgroupBarrier_IntrinsicKind:
default:
return this->assembleSimpleIntrinsic(call.function().name(), call);
}
}
static constexpr char kInverse2x2[] =
"fn mat2_inverse(m: mat2x2<f32>) -> mat2x2<f32> {"
"\n" "return mat2x2<f32>(m[1].y, -m[0].y, -m[1].x, m[0].x) * (1/determinant(m));"
"\n" "}"
"\n";
static constexpr char kInverse3x3[] =
"fn mat3_inverse(m: mat3x3<f32>) -> mat3x3<f32> {"
"\n" "let a00 = m[0].x; let a01 = m[0].y; let a02 = m[0].z;"
"\n" "let a10 = m[1].x; let a11 = m[1].y; let a12 = m[1].z;"
"\n" "let a20 = m[2].x; let a21 = m[2].y; let a22 = m[2].z;"
"\n" "let b01 = a22*a11 - a12*a21;"
"\n" "let b11 = -a22*a10 + a12*a20;"
"\n" "let b21 = a21*a10 - a11*a20;"
"\n" "let det = a00*b01 + a01*b11 + a02*b21;"
"\n" "return mat3x3<f32>(b01, (-a22*a01 + a02*a21), ( a12*a01 - a02*a11),"
"\n" "b11, ( a22*a00 - a02*a20), (-a12*a00 + a02*a10),"
"\n" "b21, (-a21*a00 + a01*a20), ( a11*a00 - a01*a10)) * (1/det);"
"\n" "}"
"\n";
static constexpr char kInverse4x4[] =
"fn mat4_inverse(m: mat4x4<f32>) -> mat4x4<f32>{"
"\n" "let a00 = m[0].x; let a01 = m[0].y; let a02 = m[0].z; let a03 = m[0].w;"
"\n" "let a10 = m[1].x; let a11 = m[1].y; let a12 = m[1].z; let a13 = m[1].w;"
"\n" "let a20 = m[2].x; let a21 = m[2].y; let a22 = m[2].z; let a23 = m[2].w;"
"\n" "let a30 = m[3].x; let a31 = m[3].y; let a32 = m[3].z; let a33 = m[3].w;"
"\n" "let b00 = a00*a11 - a01*a10;"
"\n" "let b01 = a00*a12 - a02*a10;"
"\n" "let b02 = a00*a13 - a03*a10;"
"\n" "let b03 = a01*a12 - a02*a11;"
"\n" "let b04 = a01*a13 - a03*a11;"
"\n" "let b05 = a02*a13 - a03*a12;"
"\n" "let b06 = a20*a31 - a21*a30;"
"\n" "let b07 = a20*a32 - a22*a30;"
"\n" "let b08 = a20*a33 - a23*a30;"
"\n" "let b09 = a21*a32 - a22*a31;"
"\n" "let b10 = a21*a33 - a23*a31;"
"\n" "let b11 = a22*a33 - a23*a32;"
"\n" "let det = b00*b11 - b01*b10 + b02*b09 + b03*b08 - b04*b07 + b05*b06;"
"\n" "return mat4x4<f32>(a11*b11 - a12*b10 + a13*b09,"
"\n" "a02*b10 - a01*b11 - a03*b09,"
"\n" "a31*b05 - a32*b04 + a33*b03,"
"\n" "a22*b04 - a21*b05 - a23*b03,"
"\n" "a12*b08 - a10*b11 - a13*b07,"
"\n" "a00*b11 - a02*b08 + a03*b07,"
"\n" "a32*b02 - a30*b05 - a33*b01,"
"\n" "a20*b05 - a22*b02 + a23*b01,"
"\n" "a10*b10 - a11*b08 + a13*b06,"
"\n" "a01*b08 - a00*b10 - a03*b06,"
"\n" "a30*b04 - a31*b02 + a33*b00,"
"\n" "a21*b02 - a20*b04 - a23*b00,"
"\n" "a11*b07 - a10*b09 - a12*b06,"
"\n" "a00*b09 - a01*b07 + a02*b06,"
"\n" "a31*b01 - a30*b03 - a32*b00,"
"\n" "a20*b03 - a21*b01 + a22*b00) * (1/det);"
"\n" "}"
"\n";
std::string WGSLCodeGenerator::assembleInversePolyfill(const FunctionCall& call) {
const ExpressionArray& arguments = call.arguments();
const Type& type = arguments.front()->type();
// The `inverse` intrinsic should only accept a single-argument square matrix.
// Once we implement f16 support, these polyfills will need to be updated to support `hmat`;
// for the time being, all floats in WGSL are f32, so we don't need to worry about precision.
SkASSERT(arguments.size() == 1);
SkASSERT(type.isMatrix());
SkASSERT(type.rows() == type.columns());
switch (type.slotCount()) {
case 4:
if (!fWrittenInverse2) {
fWrittenInverse2 = true;
fHeader.writeText(kInverse2x2);
}
return this->assembleSimpleIntrinsic("mat2_inverse", call);
case 9:
if (!fWrittenInverse3) {
fWrittenInverse3 = true;
fHeader.writeText(kInverse3x3);
}
return this->assembleSimpleIntrinsic("mat3_inverse", call);
case 16:
if (!fWrittenInverse4) {
fWrittenInverse4 = true;
fHeader.writeText(kInverse4x4);
}
return this->assembleSimpleIntrinsic("mat4_inverse", call);
default:
// We only support square matrices.
SkUNREACHABLE;
}
}
std::string WGSLCodeGenerator::assembleFunctionCall(const FunctionCall& call,
Precedence parentPrecedence) {
const FunctionDeclaration& func = call.function();
std::string result;
// Many intrinsics need to be rewritten in WGSL.
if (func.isIntrinsic()) {
return this->assembleIntrinsicCall(call, func.intrinsicKind(), parentPrecedence);
}
// We implement function out-parameters by declaring them as pointers. SkSL follows GLSL's
// out-parameter semantics, in which out-parameters are only written back to the original
// variable after the function's execution is complete (see
// https://www.khronos.org/opengl/wiki/Core_Language_(GLSL)#Parameters).
//
// In addition, SkSL supports swizzles and array index expressions to be passed into
// out-parameters; however, WGSL does not allow taking their address into a pointer.
//
// We support these by using LValues to create temporary copies and then pass pointers to the
// copies. Once the function returns, we copy the values back to the LValue.
// First detect which arguments are passed to out-parameters.
// TODO: rewrite this method in terms of LValues.
const ExpressionArray& args = call.arguments();
SkSpan<Variable* const> params = func.parameters();
SkASSERT(SkToSizeT(args.size()) == params.size());
STArray<16, std::unique_ptr<LValue>> writeback;
STArray<16, std::string> substituteArgument;
writeback.reserve_exact(args.size());
substituteArgument.reserve_exact(args.size());
for (int index = 0; index < args.size(); ++index) {
if (params[index]->modifierFlags() & ModifierFlag::kOut) {
std::unique_ptr<LValue> lvalue = this->makeLValue(*args[index]);
if (params[index]->modifierFlags() & ModifierFlag::kIn) {
// Load the lvalue's contents into the substitute argument.
substituteArgument.push_back(this->writeScratchVar(args[index]->type(),
lvalue->load()));
} else {
// Create a substitute argument, but leave it uninitialized.
substituteArgument.push_back(this->writeScratchVar(args[index]->type()));
}
writeback.push_back(std::move(lvalue));
} else {
substituteArgument.push_back(std::string());
writeback.push_back(nullptr);
}
}
std::string expr = this->assembleName(func.mangledName());
expr.push_back('(');
auto separator = SkSL::String::Separator();
if (std::string funcDepArgs = this->functionDependencyArgs(func); !funcDepArgs.empty()) {
expr += funcDepArgs;
separator();
}
// Pass the function arguments, or any substitutes as needed.
for (int index = 0; index < args.size(); ++index) {
expr += separator();
if (!substituteArgument[index].empty()) {
// We need to take the address of the variable and pass it down as a pointer.
expr += '&' + substituteArgument[index];
} else if (args[index]->type().isSampler()) {
// If the argument is a sampler, we need to pass the texture _and_ its associated
// sampler. (Function parameter lists also convert sampler parameters into a matching
// texture/sampler parameter pair.)
expr += this->assembleExpression(*args[index], Precedence::kSequence);
expr += kTextureSuffix;
expr += ", ";
expr += this->assembleExpression(*args[index], Precedence::kSequence);
expr += kSamplerSuffix;
} else {
expr += this->assembleExpression(*args[index], Precedence::kSequence);
}
}
expr += ')';
if (call.type().isVoid()) {
// Making function calls that result in `void` is only valid in on the left side of a
// comma-sequence, or in a top-level statement. Emit the function call as a top-level
// statement and return an empty string, as the result will not be used.
SkASSERT(parentPrecedence >= Precedence::kSequence);
this->write(expr);
this->writeLine(";");
} else {
result = this->writeScratchLet(expr);
}
// Write the substitute arguments back into their lvalues.
for (int index = 0; index < args.size(); ++index) {
if (!substituteArgument[index].empty()) {
this->writeLine(writeback[index]->store(substituteArgument[index]));
}
}
// Return the result of invoking the function.
return result;
}
std::string WGSLCodeGenerator::assembleIndexExpression(const IndexExpression& i) {
// Put the index value into a let-expression.
std::string idx = this->writeNontrivialScratchLet(*i.index(), Precedence::kExpression);
return this->assembleExpression(*i.base(), Precedence::kPostfix) + "[" + idx + "]";
}
std::string WGSLCodeGenerator::assembleLiteral(const Literal& l) {
const Type& type = l.type();
if (type.isFloat() || type.isBoolean()) {
return l.description(OperatorPrecedence::kExpression);
}
SkASSERT(type.isInteger());
if (type.matches(*fContext.fTypes.fUInt)) {
return std::to_string(l.intValue() & 0xffffffff) + "u";
} else if (type.matches(*fContext.fTypes.fUShort)) {
return std::to_string(l.intValue() & 0xffff) + "u";
} else {
return std::to_string(l.intValue());
}
}
std::string WGSLCodeGenerator::assembleIncrementExpr(const Type& type) {
// `type(`
std::string expr = to_wgsl_type(fContext, type);
expr.push_back('(');
// `1, 1, 1...)`
auto separator = SkSL::String::Separator();
for (int slots = type.slotCount(); slots > 0; --slots) {
expr += separator();
expr += "1";
}
expr.push_back(')');
return expr;
}
std::string WGSLCodeGenerator::assemblePrefixExpression(const PrefixExpression& p,
Precedence parentPrecedence) {
// Unary + does nothing, so we omit it from the output.
Operator op = p.getOperator();
if (op.kind() == Operator::Kind::PLUS) {
return this->assembleExpression(*p.operand(), Precedence::kPrefix);
}
// Pre-increment/decrement expressions have no direct equivalent in WGSL.
if (op.kind() == Operator::Kind::PLUSPLUS || op.kind() == Operator::Kind::MINUSMINUS) {
std::unique_ptr<LValue> lvalue = this->makeLValue(*p.operand());
if (!lvalue) {
return "";
}
// Generate the new value: `lvalue + type(1, 1, 1...)`.
std::string newValue =
lvalue->load() +
(p.getOperator().kind() == Operator::Kind::PLUSPLUS ? " + " : " - ") +
this->assembleIncrementExpr(p.operand()->type());
this->writeLine(lvalue->store(newValue));
return lvalue->load();
}
// WGSL natively supports unary negation/not expressions (!,~,-).
SkASSERT(op.kind() == OperatorKind::LOGICALNOT ||
op.kind() == OperatorKind::BITWISENOT ||
op.kind() == OperatorKind::MINUS);
// The unary negation operator only applies to scalars and vectors. For other mathematical
// objects (such as matrices) we can express it as a multiplication by -1.
std::string expr;
const bool needsNegation = op.kind() == Operator::Kind::MINUS &&
!p.operand()->type().isScalar() && !p.operand()->type().isVector();
const bool needParens = needsNegation || Precedence::kPrefix >= parentPrecedence;
if (needParens) {
expr.push_back('(');
}
if (needsNegation) {
expr += "-1.0 * ";
expr += this->assembleExpression(*p.operand(), Precedence::kMultiplicative);
} else {
expr += p.getOperator().tightOperatorName();
expr += this->assembleExpression(*p.operand(), Precedence::kPrefix);
}
if (needParens) {
expr.push_back(')');
}
return expr;
}
std::string WGSLCodeGenerator::assemblePostfixExpression(const PostfixExpression& p,
Precedence parentPrecedence) {
SkASSERT(p.getOperator().kind() == Operator::Kind::PLUSPLUS ||
p.getOperator().kind() == Operator::Kind::MINUSMINUS);
// Post-increment/decrement expressions have no direct equivalent in WGSL; they do exist as a
// standalone statement for convenience, but these aren't the same as SkSL's post-increments.
std::unique_ptr<LValue> lvalue = this->makeLValue(*p.operand());
if (!lvalue) {
return "";
}
// If the expression is used, create a let-copy of the original value.
// (At statement-level precedence, we know the value is unused and can skip this let-copy.)
std::string originalValue;
if (parentPrecedence != Precedence::kStatement) {
originalValue = this->writeScratchLet(lvalue->load());
}
// Generate the new value: `lvalue + type(1, 1, 1...)`.
std::string newValue = lvalue->load() +
(p.getOperator().kind() == Operator::Kind::PLUSPLUS ? " + " : " - ") +
this->assembleIncrementExpr(p.operand()->type());
this->writeLine(lvalue->store(newValue));
return originalValue;
}
std::string WGSLCodeGenerator::assembleSwizzle(const Swizzle& swizzle) {
return this->assembleExpression(*swizzle.base(), Precedence::kPostfix) + "." +
Swizzle::MaskString(swizzle.components());
}
std::string WGSLCodeGenerator::writeScratchVar(const Type& type, const std::string& value) {
std::string scratchVarName = "_skTemp" + std::to_string(fScratchCount++);
this->write("var ");
this->write(scratchVarName);
this->write(": ");
this->write(to_wgsl_type(fContext, type));
if (!value.empty()) {
this->write(" = ");
this->write(value);
}
this->writeLine(";");
return scratchVarName;
}
std::string WGSLCodeGenerator::writeScratchLet(const std::string& expr) {
std::string scratchVarName = "_skTemp" + std::to_string(fScratchCount++);
this->write(fAtFunctionScope ? "let " : "const ");
this->write(scratchVarName);
this->write(" = ");
this->write(expr);
this->writeLine(";");
return scratchVarName;
}
std::string WGSLCodeGenerator::writeScratchLet(const Expression& expr,
Precedence parentPrecedence) {
return this->writeScratchLet(this->assembleExpression(expr, parentPrecedence));
}
std::string WGSLCodeGenerator::writeNontrivialScratchLet(const Expression& expr,
Precedence parentPrecedence) {
std::string result = this->assembleExpression(expr, parentPrecedence);
return is_nontrivial_expression(expr) ? this->writeScratchLet(result)
: result;
}
std::string WGSLCodeGenerator::assembleTernaryExpression(const TernaryExpression& t,
Precedence parentPrecedence) {
std::string expr;
// The trivial case is when neither branch has side effects and evaluate to a scalar or vector
// type. This can be represented with a call to the WGSL `select` intrinsic. Select doesn't
// support short-circuiting, so we should only use it when both the true- and false-expressions
// are trivial to evaluate.
if ((t.type().isScalar() || t.type().isVector()) &&
!Analysis::HasSideEffects(*t.test()) &&
Analysis::IsTrivialExpression(*t.ifTrue()) &&
Analysis::IsTrivialExpression(*t.ifFalse())) {
bool needParens = Precedence::kTernary >= parentPrecedence;
if (needParens) {
expr.push_back('(');
}
expr += "select(";
expr += this->assembleExpression(*t.ifFalse(), Precedence::kSequence);
expr += ", ";
expr += this->assembleExpression(*t.ifTrue(), Precedence::kSequence);
expr += ", ";
bool isVector = t.type().isVector();
if (isVector) {
// Splat the condition expression into a vector.
expr += String::printf("vec%d<bool>(", t.type().columns());
}
expr += this->assembleExpression(*t.test(), Precedence::kSequence);
if (isVector) {
expr.push_back(')');
}
expr.push_back(')');
if (needParens) {
expr.push_back(')');
}
} else {
// WGSL does not support ternary expressions. Instead, we hoist the expression out into the
// surrounding block, convert it into an if statement, and write the result to a synthesized
// variable. Instead of the original expression, we return that variable.
expr = this->writeScratchVar(t.ifTrue()->type());
std::string testExpr = this->assembleExpression(*t.test(), Precedence::kExpression);
this->write("if ");
this->write(testExpr);
this->writeLine(" {");
++fIndentation;
std::string trueExpr = this->assembleExpression(*t.ifTrue(), Precedence::kAssignment);
this->write(expr);
this->write(" = ");
this->write(trueExpr);
this->writeLine(";");
--fIndentation;
this->writeLine("} else {");
++fIndentation;
std::string falseExpr = this->assembleExpression(*t.ifFalse(), Precedence::kAssignment);
this->write(expr);
this->write(" = ");
this->write(falseExpr);
this->writeLine(";");
--fIndentation;
this->writeLine("}");
}
return expr;
}
std::string WGSLCodeGenerator::variablePrefix(const Variable& v) {
if (v.storage() == Variable::Storage::kGlobal) {
// If the field refers to a pipeline IO parameter, then we access it via the synthesized IO
// structs. We make an explicit exception for `sk_PointSize` which we declare as a
// placeholder variable in global scope as it is not supported by WebGPU as a pipeline IO
// parameter (see comments in `writeStageOutputStruct`).
if (v.modifierFlags() & ModifierFlag::kIn) {
return "_stageIn.";
}
if (v.modifierFlags() & ModifierFlag::kOut) {
return "(*_stageOut).";
}
// If the field refers to an anonymous-interface-block structure, access it via the
// synthesized `_uniform0` or `_storage1` global.
if (const InterfaceBlock* ib = v.interfaceBlock()) {
const Type& ibType = ib->var()->type().componentType();
if (const std::string* ibName = fInterfaceBlockNameMap.find(&ibType)) {
return *ibName + '.';
}
}
// If the field refers to an top-level uniform, access it via the synthesized
// `_globalUniforms` global. (Note that this should only occur in test code; Skia will
// always put uniforms in an interface block.)
if (is_in_global_uniforms(v)) {
return "_globalUniforms.";
}
}
return "";
}
std::string WGSLCodeGenerator::variableReferenceNameForLValue(const VariableReference& r) {
const Variable& v = *r.variable();
if ((v.storage() == Variable::Storage::kParameter &&
v.modifierFlags() & ModifierFlag::kOut)) {
// This is an out-parameter; it's pointer-typed, so we need to dereference it. We wrap the
// dereference in parentheses, in case the value is used in an access expression later.
return "(*" + this->assembleName(v.mangledName()) + ')';
}
return this->variablePrefix(v) + this->assembleName(v.mangledName());
}
std::string WGSLCodeGenerator::assembleVariableReference(const VariableReference& r) {
// TODO(b/294274678): Correctly handle RTFlip for built-ins.
const Variable& v = *r.variable();
// Insert a conversion expression if this is a built-in variable whose type differs from the
// SkSL.
std::string expr;
std::optional<std::string_view> conversion = needs_builtin_type_conversion(v);
if (conversion.has_value()) {
expr += *conversion;
expr.push_back('(');
}
expr += this->variableReferenceNameForLValue(r);
if (conversion.has_value()) {
expr.push_back(')');
}
return expr;
}
std::string WGSLCodeGenerator::assembleAnyConstructor(const AnyConstructor& c) {
std::string expr = to_wgsl_type(fContext, c.type());
expr.push_back('(');
auto separator = SkSL::String::Separator();
for (const auto& e : c.argumentSpan()) {
expr += separator();
expr += this->assembleExpression(*e, Precedence::kSequence);
}
expr.push_back(')');
return expr;
}
std::string WGSLCodeGenerator::assembleConstructorCompound(const ConstructorCompound& c) {
if (c.type().isVector()) {
return this->assembleConstructorCompoundVector(c);
} else if (c.type().isMatrix()) {
return this->assembleConstructorCompoundMatrix(c);
} else {
fContext.fErrors->error(c.fPosition, "unsupported compound constructor");
return {};
}
}
std::string WGSLCodeGenerator::assembleConstructorCompoundVector(const ConstructorCompound& c) {
// WGSL supports constructing vectors from a mix of scalars and vectors but
// not matrices (see https://www.w3.org/TR/WGSL/#type-constructor-expr).
//
// SkSL supports vec4(mat2x2) which we handle specially.
if (c.type().columns() == 4 && c.argumentSpan().size() == 1) {
const Expression& arg = *c.argumentSpan().front();
if (arg.type().isMatrix()) {
SkASSERT(arg.type().columns() == 2);
SkASSERT(arg.type().rows() == 2);
std::string matrix = this->writeNontrivialScratchLet(arg, Precedence::kPostfix);
return String::printf("%s(%s[0], %s[1])", to_wgsl_type(fContext, c.type()).c_str(),
matrix.c_str(),
matrix.c_str());
}
}
return this->assembleAnyConstructor(c);
}
std::string WGSLCodeGenerator::assembleConstructorCompoundMatrix(const ConstructorCompound& ctor) {
SkASSERT(ctor.type().isMatrix());
std::string expr = to_wgsl_type(fContext, ctor.type()) + '(';
auto separator = String::Separator();
for (const std::unique_ptr<Expression>& arg : ctor.arguments()) {
SkASSERT(arg->type().isScalar() || arg->type().isVector());
if (arg->type().isScalar()) {
expr += separator();
expr += this->assembleExpression(*arg, Precedence::kSequence);
} else {
std::string inner = this->writeNontrivialScratchLet(*arg, Precedence::kSequence);
int numSlots = arg->type().slotCount();
for (int slot = 0; slot < numSlots; ++slot) {
String::appendf(&expr, "%s%s[%d]", separator().c_str(), inner.c_str(), slot);
}
}
}
return expr + ')';
}
std::string WGSLCodeGenerator::assembleConstructorDiagonalMatrix(
const ConstructorDiagonalMatrix& c) {
const Type& type = c.type();
SkASSERT(type.isMatrix());
SkASSERT(c.argument()->type().isScalar());
// Evaluate the inner-expression, creating a scratch variable if necessary.
std::string inner = this->writeNontrivialScratchLet(*c.argument(), Precedence::kAssignment);
// Assemble a diagonal-matrix expression.
std::string expr = to_wgsl_type(fContext, type) + '(';
auto separator = String::Separator();
for (int col = 0; col < type.columns(); ++col) {
for (int row = 0; row < type.rows(); ++row) {
expr += separator();
if (col == row) {
expr += inner;
} else {
expr += "0.0";
}
}
}
return expr + ')';
}
std::string WGSLCodeGenerator::assembleConstructorMatrixResize(
const ConstructorMatrixResize& ctor) {
std::string source = this->writeScratchLet(this->assembleExpression(*ctor.argument(),
Precedence::kSequence));
int columns = ctor.type().columns();
int rows = ctor.type().rows();
int sourceColumns = ctor.argument()->type().columns();
int sourceRows = ctor.argument()->type().rows();
auto separator = String::Separator();
std::string expr = to_wgsl_type(fContext, ctor.type()) + '(';
for (int c = 0; c < columns; ++c) {
for (int r = 0; r < rows; ++r) {
expr += separator();
if (c < sourceColumns && r < sourceRows) {
String::appendf(&expr, "%s[%d][%d]", source.c_str(), c, r);
} else if (r == c) {
expr += "1.0";
} else {
expr += "0.0";
}
}
}
return expr + ')';
}
std::string WGSLCodeGenerator::assembleEqualityExpression(const Type& left,
const std::string& leftName,
const Type& right,
const std::string& rightName,
Operator op,
Precedence parentPrecedence) {
SkASSERT(op.kind() == OperatorKind::EQEQ || op.kind() == OperatorKind::NEQ);
std::string expr;
bool isEqual = (op.kind() == Operator::Kind::EQEQ);
const char* const combiner = isEqual ? " && " : " || ";
if (left.isMatrix()) {
// Each matrix column must be compared as if it were an individual vector.
SkASSERT(right.isMatrix());
SkASSERT(left.rows() == right.rows());
SkASSERT(left.columns() == right.columns());
int columns = left.columns();
const Type& vecType = left.columnType(fContext);
const char* separator = "(";
for (int index = 0; index < columns; ++index) {
expr += separator;
std::string suffix = '[' + std::to_string(index) + ']';
expr += this->assembleEqualityExpression(vecType, leftName + suffix,
vecType, rightName + suffix,
op, Precedence::kParentheses);
separator = combiner;
}
return expr + ')';
}
if (left.isArray()) {
SkASSERT(right.matches(left));
const Type& indexedType = left.componentType();
const char* separator = "(";
for (int index = 0; index < left.columns(); ++index) {
expr += separator;
std::string suffix = '[' + std::to_string(index) + ']';
expr += this->assembleEqualityExpression(indexedType, leftName + suffix,
indexedType, rightName + suffix,
op, Precedence::kParentheses);
separator = combiner;
}
return expr + ')';
}
if (left.isStruct()) {
// Recursively compare every field in the struct.
SkASSERT(right.matches(left));
SkSpan<const Field> fields = left.fields();
const char* separator = "(";
for (const Field& field : fields) {
expr += separator;
expr += this->assembleEqualityExpression(
*field.fType, leftName + '.' + this->assembleName(field.fName),
*field.fType, rightName + '.' + this->assembleName(field.fName),
op, Precedence::kParentheses);
separator = combiner;
}
return expr + ')';
}
if (left.isVector()) {
// Compare vectors via `all(x == y)` or `any(x != y)`.
SkASSERT(right.isVector());
SkASSERT(left.slotCount() == right.slotCount());
expr += isEqual ? "all(" : "any(";
expr += leftName;
expr += operator_name(op);
expr += rightName;
return expr + ')';
}
// Compare scalars via `x == y`.
SkASSERT(right.isScalar());
if (parentPrecedence < Precedence::kSequence) {
expr = '(';
}
expr += leftName;
expr += operator_name(op);
expr += rightName;
if (parentPrecedence < Precedence::kSequence) {
expr += ')';
}
return expr;
}
std::string WGSLCodeGenerator::assembleEqualityExpression(const Expression& left,
const Expression& right,
Operator op,
Precedence parentPrecedence) {
std::string leftName, rightName;
if (left.type().isScalar() || left.type().isVector()) {
// WGSL supports scalar and vector comparisons natively. We know the expressions will only
// be emitted once, so there isn't a benefit to creating a let-declaration.
leftName = this->assembleExpression(left, Precedence::kParentheses);
rightName = this->assembleExpression(right, Precedence::kParentheses);
} else {
leftName = this->writeNontrivialScratchLet(left, Precedence::kAssignment);
rightName = this->writeNontrivialScratchLet(right, Precedence::kAssignment);
}
return this->assembleEqualityExpression(left.type(), leftName, right.type(), rightName,
op, parentPrecedence);
}
void WGSLCodeGenerator::writeProgramElement(const ProgramElement& e) {
switch (e.kind()) {
case ProgramElement::Kind::kExtension:
// TODO(skia:13092): WGSL supports extensions via the "enable" directive
// (https://www.w3.org/TR/WGSL/#enable-extensions-sec ). While we could easily emit this
// directive, we should first ensure that all possible SkSL extension names are
// converted to their appropriate WGSL extension.
break;
case ProgramElement::Kind::kGlobalVar:
this->writeGlobalVarDeclaration(e.as<GlobalVarDeclaration>());
break;
case ProgramElement::Kind::kInterfaceBlock:
// All interface block declarations are handled explicitly as the "program header" in
// generateCode().
break;
case ProgramElement::Kind::kStructDefinition:
this->writeStructDefinition(e.as<StructDefinition>());
break;
case ProgramElement::Kind::kFunctionPrototype:
// A WGSL function declaration must contain its body and the function name is in scope
// for the entire program (see https://www.w3.org/TR/WGSL/#function-declaration and
// https://www.w3.org/TR/WGSL/#declaration-and-scope).
//
// As such, we don't emit function prototypes.
break;
case ProgramElement::Kind::kFunction:
this->writeFunction(e.as<FunctionDefinition>());
break;
case ProgramElement::Kind::kModifiers:
this->writeModifiersDeclaration(e.as<ModifiersDeclaration>());
break;
default:
SkDEBUGFAILF("unsupported program element: %s\n", e.description().c_str());
break;
}
}
void WGSLCodeGenerator::writeTextureOrSampler(const Variable& var,
int bindingLocation,
std::string_view suffix,
std::string_view wgslType) {
if (var.type().dimensions() != SpvDim2D) {
// Skia currently only uses 2D textures.
fContext.fErrors->error(var.varDeclaration()->position(), "unsupported texture dimensions");
return;
}
this->write("@group(");
this->write(std::to_string(std::max(0, var.layout().fSet)));
this->write(") @binding(");
this->write(std::to_string(bindingLocation));
this->write(") var ");
this->write(this->assembleName(var.mangledName()));
this->write(suffix);
this->write(": ");
this->write(wgslType);
this->writeLine(";");
}
void WGSLCodeGenerator::writeGlobalVarDeclaration(const GlobalVarDeclaration& d) {
const VarDeclaration& decl = d.varDeclaration();
const Variable& var = *decl.var();
if ((var.modifierFlags() & (ModifierFlag::kIn | ModifierFlag::kOut)) ||
is_in_global_uniforms(var)) {
// Pipeline stage I/O parameters and top-level (non-block) uniforms are handled specially
// in generateCode().
return;
}
const Type::TypeKind varKind = var.type().typeKind();
if (varKind == Type::TypeKind::kSampler) {
// If the sampler binding was unassigned, provide a scratch value; this will make
// golden-output tests pass, but will not actually be usable for drawing.
int samplerLocation = var.layout().fSampler >= 0 ? var.layout().fSampler
: 10000 + fScratchCount++;
this->writeTextureOrSampler(var, samplerLocation, kSamplerSuffix, "sampler");
// If the texture binding was unassigned, provide a scratch value (for golden-output tests).
int textureLocation = var.layout().fTexture >= 0 ? var.layout().fTexture
: 10000 + fScratchCount++;
this->writeTextureOrSampler(var, textureLocation, kTextureSuffix, "texture_2d<f32>");
return;
}
if (varKind == Type::TypeKind::kTexture) {
// If a binding location was unassigned, provide a scratch value (for golden-output tests).
int textureLocation = var.layout().fBinding >= 0 ? var.layout().fBinding
: 10000 + fScratchCount++;
// For a texture without an associated sampler, we don't apply a suffix.
this->writeTextureOrSampler(var, textureLocation, /*suffix=*/"",
to_wgsl_type(fContext, var.type(), &var.layout()));
return;
}
std::string initializer;
if (decl.value()) {
// We assume here that the initial-value expression will not emit any helper statements.
// Initial-value expressions are required to pass IsConstantExpression, which limits the
// blast radius to constructors, literals, and other constant values/variables.
initializer += " = ";
initializer += this->assembleExpression(*decl.value(), Precedence::kAssignment);
}
if (var.modifierFlags().isConst()) {
this->write("const ");
} else if (var.modifierFlags().isWorkgroup()) {
this->write("var<workgroup> ");
} else if (var.modifierFlags().isPixelLocal()) {
this->write("var<pixel_local> ");
} else {
this->write("var<private> ");
}
this->write(this->assembleName(var.mangledName()));
this->write(": " + to_wgsl_type(fContext, var.type(), &var.layout()));
this->write(initializer);
this->writeLine(";");
}
void WGSLCodeGenerator::writeStructDefinition(const StructDefinition& s) {
const Type& type = s.type();
this->writeLine("struct " + type.displayName() + " {");
this->writeFields(type.fields(), /*memoryLayout=*/nullptr);
this->writeLine("};");
}
void WGSLCodeGenerator::writeModifiersDeclaration(const ModifiersDeclaration& modifiers) {
LayoutFlags flags = modifiers.layout().fFlags;
flags &= ~(LayoutFlag::kLocalSizeX | LayoutFlag::kLocalSizeY | LayoutFlag::kLocalSizeZ);
if (flags != LayoutFlag::kNone) {
fContext.fErrors->error(modifiers.position(), "unsupported declaration");
return;
}
if (modifiers.layout().fLocalSizeX >= 0) {
fLocalSizeX = modifiers.layout().fLocalSizeX;
}
if (modifiers.layout().fLocalSizeY >= 0) {
fLocalSizeY = modifiers.layout().fLocalSizeY;
}
if (modifiers.layout().fLocalSizeZ >= 0) {
fLocalSizeZ = modifiers.layout().fLocalSizeZ;
}
}
void WGSLCodeGenerator::writeFields(SkSpan<const Field> fields, const MemoryLayout* memoryLayout) {
fIndentation++;
// TODO(skia:14370): array uniforms may need manual fixup for std140 padding. (Those uniforms
// will also need special handling when they are accessed, or passed to functions.)
for (size_t index = 0; index < fields.size(); ++index) {
const Field& field = fields[index];
if (memoryLayout && !memoryLayout->isSupported(*field.fType)) {
// Reject types that aren't supported by the memory layout.
fContext.fErrors->error(field.fPosition, "type '" + std::string(field.fType->name()) +
"' is not permitted here");
return;
}
// Prepend @size(n) to enforce the offsets from the SkSL layout. (This is effectively
// a gadget that we can use to insert padding between elements.)
if (index < fields.size() - 1) {
int thisFieldOffset = field.fLayout.fOffset;
int nextFieldOffset = fields[index + 1].fLayout.fOffset;
if (index == 0 && thisFieldOffset > 0) {
fContext.fErrors->error(field.fPosition, "field must have an offset of zero");
return;
}
if (thisFieldOffset >= 0 && nextFieldOffset > thisFieldOffset) {
this->write("@size(");
this->write(std::to_string(nextFieldOffset - thisFieldOffset));
this->write(") ");
}
}
this->write(this->assembleName(field.fName));
this->write(": ");
if (const FieldPolyfillInfo* info = fFieldPolyfillMap.find(&field)) {
if (info->fIsArray) {
// This properly handles arrays of matrices, as well as arrays of other primitives.
SkASSERT(field.fType->isArray());
this->write("array<_skArrayElement_");
this->write(field.fType->abbreviatedName());
this->write(", ");
this->write(std::to_string(field.fType->columns()));
this->write(">");
} else if (info->fIsMatrix) {
this->write("_skMatrix");
this->write(std::to_string(field.fType->columns()));
this->write(std::to_string(field.fType->rows()));
} else {
SkDEBUGFAILF("need polyfill for %s", info->fReplacementName.c_str());
}
} else {
this->write(to_wgsl_type(fContext, *field.fType, &field.fLayout));
}
this->writeLine(",");
}
fIndentation--;
}
void WGSLCodeGenerator::writeEnables() {
this->writeLine("diagnostic(off, derivative_uniformity);");
this->writeLine("diagnostic(off, chromium.unreachable_code);");
if (fRequirements.fPixelLocalExtension) {
this->writeLine("enable chromium_experimental_pixel_local;");
}
if (fProgram.fInterface.fUseLastFragColor) {
this->writeLine("enable chromium_experimental_framebuffer_fetch;");
}
if (fProgram.fInterface.fOutputSecondaryColor) {
this->writeLine("enable dual_source_blending;");
}
}
bool WGSLCodeGenerator::needsStageInputStruct() const {
// It is illegal to declare a struct with no members; we can't emit a placeholder empty stage
// input struct.
return !fPipelineInputs.empty();
}
void WGSLCodeGenerator::writeStageInputStruct() {
if (!this->needsStageInputStruct()) {
return;
}
std::string_view structNamePrefix = pipeline_struct_prefix(fProgram.fConfig->fKind);
SkASSERT(!structNamePrefix.empty());
this->write("struct ");
this->write(structNamePrefix);
this->writeLine("In {");
fIndentation++;
for (const Variable* v : fPipelineInputs) {
if (v->type().isInterfaceBlock()) {
for (const Field& f : v->type().fields()) {
this->writePipelineIODeclaration(f.fLayout, *f.fType, f.fName, Delimiter::kComma);
}
} else {
this->writePipelineIODeclaration(v->layout(), v->type(), v->mangledName(),
Delimiter::kComma);
}
}
fIndentation--;
this->writeLine("};");
}
bool WGSLCodeGenerator::needsStageOutputStruct() const {
// It is illegal to declare a struct with no members. However, vertex programs will _always_
// have an output stage in WGSL, because the spec requires them to emit `@builtin(position)`.
// So we always synthesize a reference to `sk_Position` even if the program doesn't need it.
return !fPipelineOutputs.empty() || ProgramConfig::IsVertex(fProgram.fConfig->fKind);
}
void WGSLCodeGenerator::writeStageOutputStruct() {
if (!this->needsStageOutputStruct()) {
return;
}
std::string_view structNamePrefix = pipeline_struct_prefix(fProgram.fConfig->fKind);
SkASSERT(!structNamePrefix.empty());
this->write("struct ");
this->write(structNamePrefix);
this->writeLine("Out {");
fIndentation++;
bool declaredPositionBuiltin = false;
bool requiresPointSizeBuiltin = false;
for (const Variable* v : fPipelineOutputs) {
if (v->type().isInterfaceBlock()) {
for (const auto& f : v->type().fields()) {
this->writePipelineIODeclaration(f.fLayout, *f.fType, f.fName, Delimiter::kComma);
if (f.fLayout.fBuiltin == SK_POSITION_BUILTIN) {
declaredPositionBuiltin = true;
} else if (f.fLayout.fBuiltin == SK_POINTSIZE_BUILTIN) {
// sk_PointSize is explicitly not supported by `builtin_from_sksl_name` so
// writePipelineIODeclaration will never write it. We mark it here if the
// declaration is needed so we can synthesize it below.
requiresPointSizeBuiltin = true;
}
}
} else {
this->writePipelineIODeclaration(v->layout(), v->type(), v->mangledName(),
Delimiter::kComma);
}
}
// A vertex program must include the `position` builtin in its entrypoint's return type.
const bool positionBuiltinRequired = ProgramConfig::IsVertex(fProgram.fConfig->fKind);
if (positionBuiltinRequired && !declaredPositionBuiltin) {
this->writeLine("@builtin(position) sk_Position: vec4<f32>,");
}
fIndentation--;
this->writeLine("};");
// In WebGPU/WGSL, the vertex stage does not support a point-size output and the size
// of a point primitive is always 1 pixel (see https://github.com/gpuweb/gpuweb/issues/332).
//
// There isn't anything we can do to emulate this correctly at this stage so we synthesize a
// placeholder global variable that has no effect. Programs should not rely on sk_PointSize when
// using the Dawn backend.
if (ProgramConfig::IsVertex(fProgram.fConfig->fKind) && requiresPointSizeBuiltin) {
this->writeLine("/* unsupported */ var<private> sk_PointSize: f32;");
}
}
void WGSLCodeGenerator::prepareUniformPolyfillsForInterfaceBlock(
const InterfaceBlock* interfaceBlock,
std::string_view instanceName,
MemoryLayout::Standard nativeLayout) {
SkSL::MemoryLayout std140(MemoryLayout::Standard::k140);
SkSL::MemoryLayout native(nativeLayout);
const Type& structType = interfaceBlock->var()->type().componentType();
for (const Field& field : structType.fields()) {
const Type* type = field.fType;
bool needsArrayPolyfill = false;
bool needsMatrixPolyfill = false;
auto isPolyfillableMatrixType = [&](const Type* type) {
return type->isMatrix() && std140.stride(*type) != native.stride(*type);
};
if (isPolyfillableMatrixType(type)) {
// Matrices will be represented as 16-byte aligned arrays in std140, and reconstituted
// into proper matrices as they are later accessed. We need to synthesize polyfill.
needsMatrixPolyfill = true;
} else if (type->isArray() && !type->isUnsizedArray() &&
!type->componentType().isOpaque()) {
const Type* innerType = &type->componentType();
if (isPolyfillableMatrixType(innerType)) {
// Use a polyfill when the array contains a matrix that requires polyfill.
needsArrayPolyfill = true;
needsMatrixPolyfill = true;
} else if (native.size(*innerType) < 16) {
// Use a polyfill when the array elements are smaller than 16 bytes, since std140
// will pad elements to a 16-byte stride.
needsArrayPolyfill = true;
}
}
if (needsArrayPolyfill || needsMatrixPolyfill) {
// Add a polyfill for this matrix type.
FieldPolyfillInfo info;
info.fInterfaceBlock = interfaceBlock;
info.fReplacementName = "_skUnpacked_" + std::string(instanceName) + '_' +
this->assembleName(field.fName);
info.fIsArray = needsArrayPolyfill;
info.fIsMatrix = needsMatrixPolyfill;
fFieldPolyfillMap.set(&field, info);
}
}
}
void WGSLCodeGenerator::writeUniformsAndBuffers() {
for (const ProgramElement* e : fProgram.elements()) {
// Iterate through the interface blocks.
if (!e->is<InterfaceBlock>()) {
continue;
}
const InterfaceBlock& ib = e->as<InterfaceBlock>();
// Determine if this interface block holds uniforms, buffers, or something else (skip it).
std::string_view addressSpace;
std::string_view accessMode;
MemoryLayout::Standard nativeLayout;
if (ib.var()->modifierFlags().isUniform()) {
addressSpace = "uniform";
nativeLayout = MemoryLayout::Standard::kWGSLUniform_Base;
} else if (ib.var()->modifierFlags().isBuffer()) {
addressSpace = "storage";
nativeLayout = MemoryLayout::Standard::kWGSLStorage_Base;
accessMode = ib.var()->modifierFlags().isReadOnly() ? ", read" : ", read_write";
} else {
continue;
}
// If we have an anonymous interface block, assign a name like `_uniform0` or `_storage1`.
std::string instanceName;
if (ib.instanceName().empty()) {
instanceName = "_" + std::string(addressSpace) + std::to_string(fScratchCount++);
fInterfaceBlockNameMap[&ib.var()->type().componentType()] = instanceName;
} else {
instanceName = std::string(ib.instanceName());
}
this->prepareUniformPolyfillsForInterfaceBlock(&ib, instanceName, nativeLayout);
// Create a struct to hold all of the fields from this InterfaceBlock.
SkASSERT(!ib.typeName().empty());
this->write("struct ");
this->write(ib.typeName());
this->writeLine(" {");
// Find the struct type and fields used by this interface block.
const Type& ibType = ib.var()->type().componentType();
SkASSERT(ibType.isStruct());
SkSpan<const Field> ibFields = ibType.fields();
SkASSERT(!ibFields.empty());
MemoryLayout layout(MemoryLayout::Standard::k140);
this->writeFields(ibFields, &layout);
this->writeLine("};");
this->write("@group(");
this->write(std::to_string(std::max(0, ib.var()->layout().fSet)));
this->write(") @binding(");
this->write(std::to_string(std::max(0, ib.var()->layout().fBinding)));
this->write(") var<");
this->write(addressSpace);
this->write(accessMode);
this->write("> ");
this->write(instanceName);
this->write(" : ");
this->write(to_wgsl_type(fContext, ib.var()->type(), &ib.var()->layout()));
this->writeLine(";");
}
}
void WGSLCodeGenerator::writeNonBlockUniformsForTests() {
bool declaredUniformsStruct = false;
for (const ProgramElement* e : fProgram.elements()) {
if (e->is<GlobalVarDeclaration>()) {
const GlobalVarDeclaration& decls = e->as<GlobalVarDeclaration>();
const Variable& var = *decls.varDeclaration().var();
if (is_in_global_uniforms(var)) {
if (!declaredUniformsStruct) {
this->write("struct _GlobalUniforms {\n");
declaredUniformsStruct = true;
}
this->write(" ");
this->writeVariableDecl(var.layout(), var.type(), var.mangledName(),
Delimiter::kComma);
}
}
}
if (declaredUniformsStruct) {
int binding = fProgram.fConfig->fSettings.fDefaultUniformBinding;
int set = fProgram.fConfig->fSettings.fDefaultUniformSet;
this->write("};\n");
this->write("@binding(" + std::to_string(binding) + ") ");
this->write("@group(" + std::to_string(set) + ") ");
this->writeLine("var<uniform> _globalUniforms: _GlobalUniforms;");
}
}
std::string WGSLCodeGenerator::functionDependencyArgs(const FunctionDeclaration& f) {
WGSLFunctionDependencies* deps = fRequirements.fDependencies.find(&f);
std::string args;
if (deps && *deps) {
const char* separator = "";
if (*deps & WGSLFunctionDependency::kPipelineInputs) {
args += "_stageIn";
separator = ", ";
}
if (*deps & WGSLFunctionDependency::kPipelineOutputs) {
args += separator;
args += "_stageOut";
}
}
return args;
}
bool WGSLCodeGenerator::writeFunctionDependencyParams(const FunctionDeclaration& f) {
WGSLFunctionDependencies* deps = fRequirements.fDependencies.find(&f);
if (!deps || !*deps) {
return false;
}
std::string_view structNamePrefix = pipeline_struct_prefix(fProgram.fConfig->fKind);
if (structNamePrefix.empty()) {
return false;
}
const char* separator = "";
if (*deps & WGSLFunctionDependency::kPipelineInputs) {
this->write("_stageIn: ");
separator = ", ";
this->write(structNamePrefix);
this->write("In");
}
if (*deps & WGSLFunctionDependency::kPipelineOutputs) {
this->write(separator);
this->write("_stageOut: ptr<function, ");
this->write(structNamePrefix);
this->write("Out>");
}
return true;
}
#if defined(SK_ENABLE_WGSL_VALIDATION)
static bool validate_wgsl(ErrorReporter& reporter, const std::string& wgsl, std::string* warnings) {
// Enable the WGSL optional features that Skia might rely on.
tint::wgsl::reader::Options options;
for (auto extension : {tint::wgsl::Extension::kChromiumExperimentalPixelLocal,
tint::wgsl::Extension::kDualSourceBlending}) {
options.allowed_features.extensions.insert(extension);
}
// Verify that the WGSL we produced is valid.
tint::Source::File srcFile("", wgsl);
tint::Program program(tint::wgsl::reader::Parse(&srcFile, options));
if (program.Diagnostics().ContainsErrors()) {
// The program isn't valid WGSL.
#if defined(SKSL_STANDALONE)
reporter.error(Position(), std::string("Tint compilation failed.\n\n") + wgsl);
#else
// In debug, report the error via SkDEBUGFAIL. We also append the generated program for
// ease of debugging.
tint::diag::Formatter diagFormatter;
std::string diagOutput = diagFormatter.Format(program.Diagnostics()).Plain();
diagOutput += "\n";
diagOutput += wgsl;
SkDEBUGFAILF("%s", diagOutput.c_str());
#endif
return false;
}
if (!program.Diagnostics().empty()) {
// The program contains warnings. Report them as-is.
tint::diag::Formatter diagFormatter;
*warnings = diagFormatter.Format(program.Diagnostics()).Plain();
}
return true;
}
#endif // defined(SK_ENABLE_WGSL_VALIDATION)
bool ToWGSL(Program& program, const ShaderCaps* caps, OutputStream& out) {
TRACE_EVENT0("skia.shaders", "SkSL::ToWGSL");
SkASSERT(caps != nullptr);
program.fContext->fErrors->setSource(*program.fSource);
#ifdef SK_ENABLE_WGSL_VALIDATION
StringStream wgsl;
WGSLCodeGenerator cg(program.fContext.get(), caps, &program, &wgsl);
bool result = cg.generateCode();
if (result) {
std::string wgslString = wgsl.str();
std::string warnings;
result = validate_wgsl(*program.fContext->fErrors, wgslString, &warnings);
if (!warnings.empty()) {
out.writeText("/* Tint reported warnings. */\n\n");
}
out.writeString(wgslString);
}
#else
WGSLCodeGenerator cg(program.fContext.get(), caps, &program, &out);
bool result = cg.generateCode();
#endif
program.fContext->fErrors->setSource(std::string_view());
return result;
}
bool ToWGSL(Program& program, const ShaderCaps* caps, std::string* out) {
StringStream buffer;
if (!ToWGSL(program, caps, buffer)) {
return false;
}
*out = buffer.str();
return true;
}
} // namespace SkSL