This document describes the challenges in converting SPIR-V into WGSL.
Note: Unless otherwise specified, the namespace for C++ code is tint::reader::spirv::.
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:
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:
and provides analyses for:
Note: The SPIR-V is not modified by the SPIR-V Reader.
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.
Post-process the AST to make it valid for WGSL.
Example:
@stride attribute)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
gl_Position from gl_PerVertexGlslang SPIR-V output for a vertex shader has a gl_PerVertex output variable with four members:
gl_Positiongl_PointSizegl_ClipDistancegl_CullDistanceWGSL 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.0It'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.
There‘s some shenanigans here I don’t recall. See the SkipReason enum.
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:
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:
SpvOpSampledImageSpvOpImageSampleImplicitLodFor 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.
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();
}
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.
This is subtle.
Basically:
break, continue, and return, as needed.if, loop, continuing, switch, case).An instruction inside a SPIR-V instruction is one of:
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:
Otherwise, make a let declaration for the value.
Why:
This heuristic generates halfway-readable code, greatly reducing the varbosity of code in the common case.
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:
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
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.
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. 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: ...
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 for details on the translation.
Vulkan SPIR-V expresses a Uniform Buffer Object (UBO), or a WGSL ‘uniform buffer’ as:
Vulkan SPIR-V has two ways to express a Shader Storage Buffer Object (SSBO), or a WGSL ‘storage buffer’ as either deprecated-style:
or as new-style:
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:
read access mode. This becomes a part of its WGSL reference type (and hence corresponding pointer type).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.
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.