doc: Create overview of the SPIR-V reader
Describe the peculiare mismatches between Vulkan-flavoured
SPIR-V and WGSL that the SPIR-V Reader has to contend with.
Describe the solution in broad terms.
This is likely incomplete, but a good start.
Change-Id: I76992e9b830169d81cb55c46e849d9ad2f55c6c6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/105600
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: David Neto <dneto@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: David Neto <dneto@google.com>
diff --git a/docs/tint/spirv-reader-overview.md b/docs/tint/spirv-reader-overview.md
new file mode 100644
index 0000000..8cc7de2
--- /dev/null
+++ b/docs/tint/spirv-reader-overview.md
@@ -0,0 +1,587 @@
+# Converting SPIR-V to WGSL
+
+This document describes the challenges in converting SPIR-V into WGSL.
+
+Note: Unless otherwise specified, the namespace for C++ code is
+`tint::reader::spirv::`.
+
+## Overall flow
+
+1. Validate the SPIR-V input.
+
+ The SPIR-V module (binary blob) is validated against rules for
+ Vulkan 1.1, using the SPIRV-Tools validator.
+
+ This allows the rest of the flow to ignore invalid inputs.
+ However, the SPIR-V might still be rejected in a later step because:
+
+ - it uses features unavailable in WGSL, or
+ - the SPIR-V Reader is insufficiently smart, or
+ - the translated program tries to do something rejected by WGSL's rules
+ (which are checked by Tint's Resolver).
+
+2. Load the SPIR-V binary into an in-memory representation.
+
+ The SPIR-V reader uses the in-memory representation of the SPIR-V
+ module defined by the SPIRV-Tools optimizer. That provides
+ convenient representation of basic structures such as:
+
+ - instructions
+ - types
+ - constants
+ - functions
+ - basic blocks
+
+ and provides analyses for:
+
+ - relating definitions to uses (spvtools::opt::analysis::DefUseMgr)
+ - types (spvtools::opt::analysis:TypeManager)
+ - constants (spvtools::opt::analysis:ConstantManager)
+
+ Note: The SPIR-V is not modified by the SPIR-V Reader.
+
+3. Translate the SPIR-V module into Tint's AST.
+
+ The AST is valid for WGSL except for some small exceptions which are
+ cleaned up by transformations.
+
+4. Post-process the AST to make it valid for WGSL.
+
+ Example:
+ - Rewrite strided arrays and matrices (remove `@stride` attribute)
+ - Rewrite atomic functions
+ - Remove unreachable statements, to satisfy WGSL's behaviour analysis.
+
+
+## Overcoming mismatches between SPIR-V and WGSL
+
+### Remapping builtin inputs and outputs
+
+SPIR-V for Vulkan models builtin inputs and outputs as variables
+in Input and Output storage classes.
+
+WGSL builtin inputs are parameters to the entry point, and
+builtin outputs are result values of the entry point.
+
+See [spirv-input-output-variables.md](spirv-input-output-variables.md)
+
+### We only care about `gl_Position` from `gl_PerVertex`
+
+Glslang SPIR-V output for a vertex shader has a `gl_PerVertex`
+output variable with four members:
+
+- `gl_Position`
+- `gl_PointSize`
+- `gl_ClipDistance`
+- `gl_CullDistance`
+
+WGSL only supports the `position` builtin variable.
+
+The SPIR-V Reader has a bunch of carveouts so it only generates the
+position variable. In partcular, it tracks which expressions are actually
+accesses into the per-vertex variable, and ignores accesses to other
+parts of the structure, and remaps accesses of the position member.
+
+### `gl_PointSize` must be 1.0
+
+It's a WGSL rule. SPIR-V is more flexible, and the SPIR-V Reader
+checks that any assignment to (the equivalent of) `gl_PointSize`
+must the constant value 1.0.
+
+### Remapping sample mask inputs and outputs
+
+There's some shenanigans here I don't recall.
+See the SkipReason enum.
+
+### Integer signedness
+
+In SPIR-V, the instruction determines the signedness of an operation,
+not the types of its operands.
+
+For example:
+
+ %uint = OpTypeInt 32 0 ; u32 type
+ %int = OpTypeInt 32 1 ; i32 type
+
+ %int_1 = OpConstant %int 1 ; WGSL 1i
+ %uint_2 = OpConstant %uint 2 ; WGSL 2u
+
+ ; You can mix signs of an operand, and the instruction
+ ; tells you the result type.
+ %sum_uint = OpIAdd %uint %int %int_1 %uint_2
+ %sum_int = OpIAdd %int %int %int_1 %uint_2
+
+However, WGSL arithmetic tends to require the operands and
+result type for an operation to all have the same signedness.
+
+So the above might translate to WGSL as:
+
+ let sum_uint: u32 = bitcast<u32>(1i) + 2u;
+ let sum_int: i32 = 1i + bitcast<i32>(2u);
+
+See:
+* ParserImpl::RectifyOperandSignedness
+* ParserImpl::RectifySecondOperandSignedness
+* ParserImpl::RectifyForcedResultType
+
+### Translating textures and samplers
+
+SPIR-V textures and samplers are module-scope variables
+in UniformConstant storage class.
+These map directly to WGSL variables.
+
+For a sampled-image operation, SPIR-V will:
+- load the image value from a texture variable
+- load the sampler value from a sampler variable
+- form a "sampled image" value using `SpvOpSampledImage`
+- then use that sampled image value in a image operation
+ such as `SpvOpImageSampleImplicitLod`
+
+For an image operation that is not a sampled-image operation
+(e.g. OpImageLoad or OpImageWrite), then the steps are similar
+except without a sampler (clearly), and without invoking
+`OpSampledImage`.
+
+In contrast to the SPIR-V code pattern, the WGSL builtin requires
+the texture and sampler value to be passed in as separate parameters.
+Secondly, they are passed in by value, by naming the variables
+themselves and relying on WGSL's "Load Rule" to pass the handle
+value into the callee.
+
+When the SPIR-V Reader translates a texture builtin, it traces
+backward through the `OpSampledImage` operation (if any),
+back through the load, and all the way back to the `OpVariable`
+declaration. It does this for both the image/texture variable and
+the sampler variable (if applicable). It then uses the names
+of those variables as the corresponding arguments to the WGSL
+texture builtin.
+
+### Passing textures and samplers into helper functions
+
+Note: Bug https://crbug.com/tint/1039 is open to support passing
+textures and samplers as function arguments.
+
+Glslang generates SPIR-V where texture and sampler formal parameters
+are as pointer-to-UniformConstant.
+
+WGSL models them as passing texture and sampler values themselves,
+conceptually as opaque handles. This is similar to GLSL, but unlike
+SPIR-V.
+
+The tracing logic described in the previous section does not know
+what to do when it bottoms out on a formal parameter,
+e.g. OpFunctionParameter.
+
+Example GLSL compute shader:
+
+ #version 450
+
+ layout(set=0,binding=0) uniform texture2D im;
+ layout(set=0,binding=1) uniform sampler s;
+
+ vec4 helper(texture2D imparam, sampler sparam) {
+ return texture(sampler2D(imparam,sparam),vec2(0));
+ }
+
+ void main() {
+ vec4 v = helper(im,s);
+ }
+
+SPIR-V generated by Glslang (Shaderc's glslc):
+
+ ; SPIR-V
+ ; Version: 1.0
+ ; Generator: Google Shaderc over Glslang; 10
+ ; Bound: 32
+ ; Schema: 0
+ OpCapability Shader
+ %1 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpSource GLSL 450
+ OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
+ OpSourceExtension "GL_GOOGLE_include_directive"
+ OpName %main "main"
+ OpName %helper_t21_p1_ "helper(t21;p1;"
+ OpName %imparam "imparam"
+ OpName %sparam "sparam"
+ OpName %v "v"
+ OpName %im "im"
+ OpName %s "s"
+ OpDecorate %im DescriptorSet 0
+ OpDecorate %im Binding 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 1
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %7 = OpTypeImage %float 2D 0 0 0 1 Unknown
+ %_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+ %9 = OpTypeSampler
+ %_ptr_UniformConstant_9 = OpTypePointer UniformConstant %9
+ %v4float = OpTypeVector %float 4
+ %12 = OpTypeFunction %v4float %_ptr_UniformConstant_7 %_ptr_UniformConstant_9
+ %19 = OpTypeSampledImage %7
+ %v2float = OpTypeVector %float 2
+ %float_0 = OpConstant %float 0
+ %23 = OpConstantComposite %v2float %float_0 %float_0
+ %_ptr_Function_v4float = OpTypePointer Function %v4float
+ %im = OpVariable %_ptr_UniformConstant_7 UniformConstant
+ %s = OpVariable %_ptr_UniformConstant_9 UniformConstant
+ %main = OpFunction %void None %3
+ %5 = OpLabel
+ %v = OpVariable %_ptr_Function_v4float Function
+ %31 = OpFunctionCall %v4float %helper_t21_p1_ %im %s
+ OpStore %v %31
+ OpReturn
+ OpFunctionEnd
+ %helper_t21_p1_ = OpFunction %v4float None %12
+ %imparam = OpFunctionParameter %_ptr_UniformConstant_7
+ %sparam = OpFunctionParameter %_ptr_UniformConstant_9
+ %16 = OpLabel
+ %17 = OpLoad %7 %imparam
+ %18 = OpLoad %9 %sparam
+ %20 = OpSampledImage %19 %17 %18
+ %24 = OpImageSampleExplicitLod %v4float %20 %23 Lod %float_0
+ OpReturnValue %24
+ OpFunctionEnd
+
+What the SPIR-V Reader currently generates:
+
+ @group(0) @binding(0) var im : texture_2d<f32>;
+
+ @group(0) @binding(1) var s : sampler;
+
+ fn helper_t21_p1_(imparam : ptr<none, void>, sparam : ptr<none, void>) -> vec4<f32> {
+ let x_24 : vec4<f32> = textureSampleLevel(imparam, sparam, vec2<f32>(0.0f, 0.0f), 0.0f);
+ return x_24;
+ }
+
+ fn main_1() {
+ var v : vec4<f32>;
+ let x_31 : vec4<f32> = helper_t21_p1_(&(im), &(s));
+ v = x_31;
+ return;
+ }
+
+ @compute @workgroup_size(1i, 1i, 1i)
+ fn main() {
+ main_1();
+ }
+
+with an error:
+
+ error: function parameter of pointer type cannot be in 'none' address space
+
+Instead, the generated WGSL should have formal parameters with texture
+and sampler types, rather than as pointers to them. So the generated WGSL
+should look like this instead:
+
+ @group(0) @binding(0) var im : texture_2d<f32>;
+
+ @group(0) @binding(1) var s : sampler;
+
+ fn helper_t21_p1_(imparam : texture_2d<f32>, sparam : sampler) -> vec4<f32> {
+ let x_24 : vec4<f32> = textureSampleLevel(imparam, sparam, vec2<f32>(0.0f, 0.0f), 0.0f);
+ return x_24;
+ }
+
+ fn main_1() {
+ var v : vec4<f32>;
+ let x_31 : vec4<f32> = helper_t21_p1_(im, s);
+ v = x_31;
+ return;
+ }
+
+ @compute @workgroup_size(1i, 1i, 1i)
+ fn main() {
+ main_1();
+ }
+
+### Dimensionality mismatch in texture builtins
+
+Vulkan SPIR-V is fairly forgiving in the dimensionality
+of input coordinates and result values of texturing operations.
+There is some localized rewriting of values to satisfy the overloads
+of WGSL's texture builtin functions.
+
+### Reconstructing structured control flow
+
+This is subtle.
+
+- Use structural dominance (but we didn't have the name at the time).
+ See SPIR-V 1.6 Rev 2 for updated definitions.
+- See the big comment at the start of reader/spirv/function.cc
+- See internal presentations.
+
+Basically:
+* Compute a "structured order" for structurally reachable basic blocks.
+* Traversing in structured order, use a stack-based algorithn to
+ identify intervals of blocks corresponding to structured constructs.
+ For example, loop construct, continue construct, if-selection,
+ switch-selection, and case-construct. Constructs can be nested,
+ hence the need for a stack. This is akin to "drawing braces"
+ around statements, to form block-statements that will appear in
+ the output. This step performs some validation, which may now be
+ redundant with the SPIRV-Tools validator. This is defensive
+ programming, and some tests skip use of the SPIRV-Tools validator.
+* Traversing in structured order, identify structured exits from the
+ constructs identified in the previous step. This determines what
+ control flow edges correspond to `break`, `continue`, and `return`,
+ as needed.
+* Traversing in structured order, generate statements for instructions.
+ This uses a stack corresponding to nested constructs. The kind of
+ each construct being entered or exited determines emission of control
+ flow constructs (WGSL's `if`, `loop`, `continuing`, `switch`, `case`).
+
+### Preserving execution order
+
+An instruction inside a SPIR-V instruction is one of:
+
+- control flow: see the previous section
+- combinatorial: think of this as an ALU operation, i.e. the effect
+ is purely to evaluate a result value from the values of its operands.
+ It has no side effects, and is not affected by external state such
+ as memory or the actions of other invocations in its subgroup.
+ Examples: arithmetic, OpCopyObject
+- interacts with memory or other invocations in some way.
+ Examples: load, store, atomics, barriers, (subgroup operations when we
+ get them)
+- function calls: functions are not analyzed to see if they are pure,
+ so we assume function calls are non-combinatorial.
+
+To preserve execution order, all non-combinatorial instructions must
+be translated as their own separate statement. For example, an OpStore
+maps to an assignment statement.
+
+However, combinatorial instructions can be emitted at any point
+in evaluation, provided data flow constraints are satisfied: input
+values are available, and such that the resulting value is generated
+in time for consumption by downstream uses.
+
+The SPIR-V Reader uses a heuristic to choose when to emit combinatorial
+values:
+- if a combinatorial expression only has one use, *and*
+- its use is in the same structured construct as its definition, *then*
+- emit the expression at the place where it is consumed.
+
+Otherwise, make a `let` declaration for the value.
+
+Why:
+- If a value has many uses, then computing it once can save effort.
+ Preserve that choice if it was made by an upstream optimizing compiler.
+- If a value is consumed in a different structured construct, then the
+ site of its consumption may be inside a loop, and we don't want to
+ sink the computation into the loop, thereby causing spurious extra
+ evaluation.
+
+This heuristic generates halfway-readable code, greatly reducing the
+varbosity of code in the common case.
+
+### Hoisting and phis
+
+SPIR-V uses SSA (static single assignment). The key requirement is
+that the definition of a value must dominate its uses.
+
+WGSL uses lexical scoping.
+
+It is easy enough for a human or an optimizing compiler to generate
+SSA cases which do not map cleanly to a lexically scoped value.
+
+Example pseudo-GLSL:
+
+ void main() {
+ if (cond) {
+ const uint x = 1;
+ } else {
+ return;
+ }
+ const uint y = x; // x's definition dominates this use.
+ }
+
+This isn't valid GLSL and its analog would not be a valid WGSL
+program because x is used outside the scope of its declaration.
+
+Additionally, SSA uses `phi` nodes to transmit values from predecessor
+basic blocks that would otherwise not be visible (because the
+parent does not dominate the consuming basic block). An example
+is sending the updated value of a loop induction variable back to
+the top of the loop.
+
+The SPIR-V reader handles these cases by tracking:
+- where a value definition occurs
+- the span of basic blocks, in structured order, where there
+ are uses of the value.
+
+If the uses of a value span structured contructs which are not
+contained by the construct containing the definition (or
+if the value is a `phi` node), then we "hoist" the value
+into a variable:
+
+- create a function-scope variable at the top of the structured
+ construct that spans all the uses, so that all the uses
+ are in scope of that variable declaration.
+
+- for a non-phi: generate an assignment to that variable corresponding
+ to the value definition in the original SPIR-V.
+
+- for a phi: generate an assigment to that variable at the end of
+ each predecessor block for that phi, assigning the value to be
+ transmitted from that phi.
+
+This scheme works for values which can be the stored in a variable.
+
+It does not work for pointers. However, we don't think we need
+to solve this case any time soon as it is uncommon or hard/impossible
+to generate via standard tooling.
+See https://crbug.com/tint/98 and https://crbug.com/tint/837
+
+## Mapping types
+
+SPIR-V has a recursive type system. Types are defined, given result IDs,
+before any functions are defined, and before any constant values using
+the corresponding types.
+
+WGSL also has a recursive type system. However, except for structure types,
+types are spelled inline at their uses.
+
+## Texture and sampler types
+
+SPIR-V image types map to WGSL types, but the WGSL type is determined
+more by usage (what operations are performed on it) than by declaration.
+
+For example, Vulkan ignores the "Depth" operand of the image type
+declaration (OpTypeImage).
+See [16.1 Image Operations Overview](https://registry.khronos.org/vulkan/specs/1.3/html/vkspec.html#_image_operations_overview).
+Instead, we must infer that a texture is a depth texture because
+it is used by image instructions using a depth-reference, e.g.
+OpImageSampleDrefImplicitLod vs. OpImageSampleImplicitLod.
+
+Similarly, SPIR-V only has one sampler type. The use of the
+sampler determines whether it maps to a WGSL `sampler` or
+`sampler_comparison` (for depth sampling).
+
+The SPIR-V Reader scans uses of each texture and sampler
+in the module to infer the appropriate target WGSL type.
+See ParserImpl::RegisterHandleUsage
+
+In Vulkan SPIR-V it is possible to use the same sampler for regular
+sampling and depth-reference sampling. In this case the SPIR-V Reader
+will infer a depth texture, but then the generated program will fail WGSL
+validation.
+
+For example, this GLSL fragment shader:
+
+ #version 450
+
+ layout(set=1,binding=0) uniform texture2D tInput;
+ layout(set=1,binding=1) uniform sampler s;
+
+ void main() {
+ vec4 v = texture(sampler2D(tInput,s),vec2(0));
+ float f = texture(sampler2DShadow(tInput,s),vec3(0));
+ }
+
+Converts to this WGSL shader:
+
+ @group(1) @binding(0) var tInput : texture_depth_2d;
+
+ @group(1) @binding(1) var s : sampler_comparison;
+
+ fn main_1() {
+ var v : vec4<f32>;
+ var f : f32;
+ let x_23 : vec4<f32> = vec4<f32>(textureSample(tInput, s, vec2<f32>(0.0f, 0.0f)), 0.0f, 0.0f, 0.0f);
+ v = x_23;
+ let x_34 : f32 = textureSampleCompare(tInput, s, vec3<f32>(0.0f, 0.0f, 0.0f).xy, vec3<f32>(0.0f, 0.0f, 0.0f).z);
+ f = x_34;
+ return;
+ }
+
+ @fragment
+ fn main() {
+ main_1();
+ }
+
+But then this fails validation:
+
+ error: no matching call to textureSample(texture_depth_2d, sampler_comparison, vec2<f32>)
+ 15 candidate functions: ...
+
+## References and pointers
+
+SPIR-V has a pointer type.
+
+A SPIR-V pointer type corresponds to a WGSL memory view. WGSL has two
+memory view types: a reference type, and a pointer type.
+
+See [spirv-ptr-ref.md](spirv-ptr-ref.md) for details on the translation.
+
+## Mapping buffer types
+
+Vulkan SPIR-V expresses a Uniform Buffer Object (UBO), or
+a WGSL 'uniform buffer' as:
+
+- an OpVariable in Uniform storage class
+- its pointee type (store type) is a Block-decorated structure type
+
+Vulkan SPIR-V has two ways to express a Shader Storage Buffer Object (SSBO),
+or a WGSL 'storage buffer' as either deprecated-style:
+
+- an OpVariable in Uniform storage class
+- its pointee type (store type) is a BufferBlock-decorated structure type
+
+or as new-style:
+
+- an OpVariable in StorageBuffer storage class
+- its pointee type (store type) is a Block-decorated structure type
+
+Deprecated-style storage buffer was the only option in un-extended
+Vulkan 1.0. It is generated by tools that want to generate code for
+the broadest reach. This includes DXC.
+
+New-style storage buffer requires the use of the `OpExtension
+"SPV_KHR_storage_buffer_storage_class"` or SPIR-V 1.3 or later
+(Vulkan 1.1 or later).
+
+Additionally, a storage buffer in SPIR-V may be marked as NonWritable.
+Perhaps surprisingly, this is typically done by marking *all* the
+members of the top-level (Buffer)Block-decorated structure as NonWritable.
+(This is the common translation because that's what Glslang does.)
+
+Translation of uniform buffers is straightforward.
+
+However, the SPIR-V Reader must support both the deprecated and the new
+styles of storage buffers.
+
+Additionally:
+- a storage buffer with all NonWritable members is translated with `read`
+ access mode. This becomes a part of its WGSL reference type (and hence
+ corresponding pointer type).
+- a storage buffer without all NonWritable members is translated with
+ an explicit `read_write` access mode. This becomes a part of its
+ WGSL reference type (and hence corresponding pointer type).
+
+Note that read-only vs. read-write is a property of the pointee-type in SPIR-V,
+but in WGSL it's part of the reference type (not the store type).
+
+To handle this mismatch, the SPIR-V Reader has bookkeeping to map
+each pointer value (inside a function) back to through to the originating
+variable. This originating variable may be a buffer variable which then
+tells us which address space and access mode to use for a locally-defined
+pointer value.
+
+Since baseline SPIR-V does not allow passing pointers to buffers into
+user-defined helper functions, we don't need to handle this buffer type
+remapping into function formal parameters.
+
+## Mapping OpArrayLength
+
+The OpArrayLength instruction takes a pointer to the enclosing
+structure (the pointee type of the storage buffer variable).
+
+But the WGSL arrayLength builtin variable takes a pointer to the
+member inside that structure.
+
+A small local adjustment is sufficient here.
