# Overview
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.
# Differences from GLSL
* 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 <base type><columns>, so float2 instead of vec2 and
bool4 instead of bvec4
* Matrix types are named <base type><columns>x<rows>, so float2x3 instead of
mat2x3 and double4x4 instead of dmat4
* 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.
* 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!)
# Synchronization Primitives
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.
## Atomics
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.
### Backend considerations and differences from GLSL
`atomicUint` should not be confused with the GLSL [`atomic_uint` (aka Atomic
Counter)](https://www.khronos.org/opengl/wiki/Atomic_Counter) type. The semantics provided by
`atomicUint` are more similar to GLSL ["Atomic Memory
Functions"](https://www.khronos.org/opengl/wiki/Atomic_Variable_Operations)
(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).
* The semantics of `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.
* When translated to Metal, the atomic intrinsics use relaxed memory order semantics.
* When translated to SPIR-V, the atomic intrinsics use relaxed [memory
semantics](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#Memory_Semantics_-id-)
(i.e. `0x0 None`). The [memory
scope](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#Scope_-id-) is either `1
Device` or `2 Workgroup` depending on whether the `atomicUint` is declared in a buffer block or
workgroup variable.
## Barriers
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](https://www.w3.org/TR/WGSL/#sync-builtin-functions). More
specifically:
* Both functions execute a control barrier with Acquire/Release memory ordering.
* Both functions use a `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).
### Backend considerations
* The closest GLSL equivalent for `workgroupBarrier()` is the
[`barrier()`](https://registry.khronos.org/OpenGL-Refpages/gl4/html/barrier.xhtml) intrinsic. Both
`workgroupBarrier()` and `storageBarrier()` can be defined as the following invocations of the
`controlBarrier` intrinsic defined in
[GL_KHR_memory_scope_semantics](https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_memory_scope_semantics.txt):
```
// 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.
