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.
Precision modifiers are not used. ‘float’, ‘int’, and ‘uint’ are always high precision. New types ‘half’, ‘short’, and ‘ushort’ are medium precision (we do not use low precision).
Vector types are named , so float2 instead of vec2 and bool4 instead of bvec4
Matrix types are named x, so float2x3 instead of mat2x3 and double4x4 instead of dmat4
GLSL caps can be referenced via the syntax ‘sk_Caps.’, 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.
no #version statement is required, and it will be ignored if present
the output color is sk_FragColor (do not declare it)
use sk_Position instead of gl_Position. sk_Position is in device coordinates rather than normalized coordinates.
use sk_PointSize instead of gl_PointSize
use sk_VertexID instead of gl_VertexID
use sk_InstanceID instead of gl_InstanceID
the fragment coordinate is sk_FragCoord, and is always relative to the upper left.
use sk_Clockwise instead of gl_FrontFacing. This is always relative to an upper left origin.
you do not need to include “.0” to make a number a float (meaning that “float2(x, y) * 4” is perfectly legal in SkSL, unlike GLSL where it would often have to be expressed “float2(x, y) * 4.0”. There is no performance penalty for this, as the number is converted to a float at compile time)
type suffixes on numbers (1.0f, 0xFFu) are both unnecessary and unsupported
creating a smaller vector from a larger vector (e.g. float2(float3(1))) is intentionally disallowed, as it is just a wordier way of performing a swizzle. Use swizzles instead.
Swizzle components, in addition to the normal rgba / xyzw components, can also be LTRB (meaning “left/top/right/bottom”, for when we store rectangles in vectors), and may also be the constants ‘0’ or ‘1’ to produce a constant 0 or 1 in that channel instead of selecting anything from the source vector. foo.rgb1 is equivalent to float4(foo.rgb, 1).
All texture functions are named “sample”, e.g. sample(sampler2D, float3) is equivalent to GLSL's textureProj(sampler2D, float3).
Functions support the ‘inline’ modifier, which causes the compiler to ignore its normal inlining heuristics and inline the function if at all possible
some built-in functions and one or two rarely-used language features are not yet supported (sorry!)
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
.
workgroup atomicUint myLocalAtomicUint; layout(set = 0, binding = 0) buffer mySSBO { atomicUint myGlobalAtomicUint; };
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(): controlBarrier(gl_ScopeWorkgroup, gl_ScopeWorkgroup, gl_StorageSemanticsShared, gl_SemanticsAcquireRelease); // storageBarrier(): controlBarrier(gl_ScopeWorkgroup, gl_ScopeWorkgroup, gl_StorageSemanticsBuffer, gl_SemanticsAcquireRelease);
In Metal, 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.