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