encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
David Neto 2022-11-01 03:54:42 +00:00 committed by Dawn LUCI CQ
parent dec01f1fd5
commit 0e6d95bafd
1 changed files with 587 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

587
docs/tint/spirv-reader-overview.md Normal file
View File

@ -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.