Flutter Engine
The Flutter Engine
|
SkSL ("Skia Shading Language") is a variant of GLSL which is used as Skia's internal shading language. SkSL is, at its heart, a single standardized version of GLSL which avoids all of the various version and dialect differences found in GLSL "in the wild", but it does bring a few of its own changes to the table.
Skia uses the SkSL compiler to convert SkSL code to GLSL, GLSL ES, SPIR-V, or MSL before handing it over to the graphics driver.
GLSL caps can be referenced via the syntax 'sk_Caps.<name>', e.g. sk_Caps.integerSupport. The value will be a constant boolean or int, as appropriate. As SkSL supports constant folding and branch elimination, this means that an 'if' statement which statically queries a cap will collapse down to the chosen branch, meaning that:
if (sk_Caps.integerSupport) do_something(); else do_something_else();
will compile as if you had written either 'do_something();' or 'do_something_else();', depending on whether that cap is enabled or not.
SkSL offers atomic operations and synchronization primitives geared towards GPU compute programs. These primitives are designed to abstract over the capabilities provided by MSL, SPIR-V, and WGSL, and differ from the corresponding primitives in GLSL.
SkSL provides the atomicUint
type. This is an opaque type that requires the use of an atomic intrinsic (such as atomicLoad
, atomicStore
, and atomicAdd
) to act on its value (which is of type uint
).
A variable with the atomicUint
type must be declared inside a writable storage buffer block or as a workgroup-shared variable. When declared inside a buffer block, it is guaranteed to conform to the same size and stride as a uint
.
An atomicUint
can be declared as a struct member or the element type of an array, provided that the struct/array type is only instantiated in a workgroup-shared or storage buffer block variable.
atomicUint
should not be confused with the GLSL atomic_uint
(aka Atomic Counter) type. The semantics provided by atomicUint
are more similar to GLSL "Atomic Memory Functions" (see GLSL Spec v4.3, 8.11 "Atomic Memory Functions"). The key difference is that SkSL atomic operations only operate on a variable of type atomicUint
while GLSL Atomic Memory Functions can operate over arbitrary memory locations (such as a component of a vector).
atomicUint
are similar to Metal's atomic<uint>
and WGSL's atomic<u32>
. These are the types that an atomicUint
is translated to when targeting Metal and WGSL.0x0 None
). The memory scope is either 1 Device
or 2 Workgroup
depending on whether the atomicUint
is declared in a buffer block or workgroup variable.SkSL provides two barrier intrinsics: workgroupBarrier()
and storageBarrier()
. These functions are only available in compute programs and synchronize access to workgroup-shared and storage buffer memory between invocations in the same workgroup. They provide the same semantics as the equivalent WGSL Synchronization Built-in Functions. More specifically:
Workgroup
execution and memory scope. This means that a coherent memory view is only guaranteed between invocations in the same workgroup and NOT across workgroups in a given compute pipeline dispatch. If multiple workgroups require a synchronized coherent view over the same shared mutable state, their access must be synchronized via other means (such as a pipeline barrier between multiple dispatches).workgroupBarrier()
is the barrier()
intrinsic. Both workgroupBarrier()
and storageBarrier()
can be defined as the following invocations of the controlBarrier
intrinsic defined in GL_KHR_memory_scope_semantics:workgroupBarrier()
is equivalent to threadgroup_barrier(mem_flags::mem_threadgroup)
. storageBarrier()
is equivalent to threadgroup_barrier(mem_flags::mem_device)
.In Vulkan SPIR-V, workgroupBarrier()
is equivalent to OpControlBarrier
with Workgroup
execution and memory scope, and AcquireRelease | WorkgroupMemory
memory semantics.
storageBarrier()
is equivalent to OpControlBarrier
with Workgroup
execution and memory scope, and AcquireRelease | UniformMemory
memory semantics.