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