README.md
1# Overview
2
3SkSL ("Skia Shading Language") is a variant of GLSL which is used as Skia's
4internal shading language. SkSL is, at its heart, a single standardized version
5of GLSL which avoids all of the various version and dialect differences found
6in GLSL "in the wild", but it does bring a few of its own changes to the table.
7
8Skia uses the SkSL compiler to convert SkSL code to GLSL, GLSL ES, SPIR-V, or
9MSL before handing it over to the graphics driver.
10
11
12# Differences from GLSL
13
14* Precision modifiers are not used. 'float', 'int', and 'uint' are always high
15 precision. New types 'half', 'short', and 'ushort' are medium precision (we
16 do not use low precision).
17* Vector types are named <base type><columns>, so float2 instead of vec2 and
18 bool4 instead of bvec4
19* Matrix types are named <base type><columns>x<rows>, so float2x3 instead of
20 mat2x3 and double4x4 instead of dmat4
21* GLSL caps can be referenced via the syntax 'sk_Caps.<name>', e.g.
22 sk_Caps.integerSupport. The value will be a constant boolean or int,
23 as appropriate. As SkSL supports constant folding and branch elimination, this
24 means that an 'if' statement which statically queries a cap will collapse down
25 to the chosen branch, meaning that:
26
27 if (sk_Caps.integerSupport)
28 do_something();
29 else
30 do_something_else();
31
32 will compile as if you had written either 'do_something();' or
33 'do_something_else();', depending on whether that cap is enabled or not.
34* no #version statement is required, and it will be ignored if present
35* the output color is sk_FragColor (do not declare it)
36* use sk_Position instead of gl_Position. sk_Position is in device coordinates
37 rather than normalized coordinates.
38* use sk_PointSize instead of gl_PointSize
39* use sk_VertexID instead of gl_VertexID
40* use sk_InstanceID instead of gl_InstanceID
41* the fragment coordinate is sk_FragCoord, and is always relative to the upper
42 left.
43* use sk_Clockwise instead of gl_FrontFacing. This is always relative to an
44 upper left origin.
45* you do not need to include ".0" to make a number a float (meaning that
46 "float2(x, y) * 4" is perfectly legal in SkSL, unlike GLSL where it would
47 often have to be expressed "float2(x, y) * 4.0". There is no performance
48 penalty for this, as the number is converted to a float at compile time)
49* type suffixes on numbers (1.0f, 0xFFu) are both unnecessary and unsupported
50* creating a smaller vector from a larger vector (e.g. float2(float3(1))) is
51 intentionally disallowed, as it is just a wordier way of performing a swizzle.
52 Use swizzles instead.
53* Swizzle components, in addition to the normal rgba / xyzw components, can also
54 be LTRB (meaning "left/top/right/bottom", for when we store rectangles in
55 vectors), and may also be the constants '0' or '1' to produce a constant 0 or
56 1 in that channel instead of selecting anything from the source vector.
57 foo.rgb1 is equivalent to float4(foo.rgb, 1).
58* All texture functions are named "sample", e.g. sample(sampler2D, float3) is
59 equivalent to GLSL's textureProj(sampler2D, float3).
60* Functions support the 'inline' modifier, which causes the compiler to ignore
61 its normal inlining heuristics and inline the function if at all possible
62* some built-in functions and one or two rarely-used language features are not
63 yet supported (sorry!)
64
65
66# Synchronization Primitives
67
68SkSL offers atomic operations and synchronization primitives geared towards GPU compute
69programs. These primitives are designed to abstract over the capabilities provided by
70MSL, SPIR-V, and WGSL, and differ from the corresponding primitives in GLSL.
71
72## Atomics
73
74SkSL provides the `atomicUint` type. This is an opaque type that requires the use of an
75atomic intrinsic (such as `atomicLoad`, `atomicStore`, and `atomicAdd`) to act on its value (which
76is of type `uint`).
77
78A variable with the `atomicUint` type must be declared inside a writable storage buffer block or as
79a workgroup-shared variable. When declared inside a buffer block, it is guaranteed to conform to the
80same size and stride as a `uint`.
81
82```
83workgroup atomicUint myLocalAtomicUint;
84
85layout(set = 0, binding = 0) buffer mySSBO {
86 atomicUint myGlobalAtomicUint;
87};
88
89```
90
91An `atomicUint` can be declared as a struct member or the element type of an array, provided that
92the struct/array type is only instantiated in a workgroup-shared or storage buffer block variable.
93
94### Backend considerations and differences from GLSL
95
96`atomicUint` should not be confused with the GLSL [`atomic_uint` (aka Atomic
97Counter)](https://www.khronos.org/opengl/wiki/Atomic_Counter) type. The semantics provided by
98`atomicUint` are more similar to GLSL ["Atomic Memory
99Functions"](https://www.khronos.org/opengl/wiki/Atomic_Variable_Operations)
100(see GLSL Spec v4.3, 8.11 "Atomic Memory Functions"). The key difference is that SkSL atomic
101operations only operate on a variable of type `atomicUint` while GLSL Atomic Memory Functions can
102operate over arbitrary memory locations (such as a component of a vector).
103
104* The semantics of `atomicUint` are similar to Metal's `atomic<uint>` and WGSL's `atomic<u32>`.
105 These are the types that an `atomicUint` is translated to when targeting Metal and WGSL.
106* When translated to Metal, the atomic intrinsics use relaxed memory order semantics.
107* When translated to SPIR-V, the atomic intrinsics use relaxed [memory
108 semantics](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#Memory_Semantics_-id-)
109 (i.e. `0x0 None`). The [memory
110 scope](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#Scope_-id-) is either `1
111 Device` or `2 Workgroup` depending on whether the `atomicUint` is declared in a buffer block or
112 workgroup variable.
113
114## Barriers
115
116SkSL provides two barrier intrinsics: `workgroupBarrier()` and `storageBarrier()`. These functions
117are only available in compute programs and synchronize access to workgroup-shared and storage buffer
118memory between invocations in the same workgroup. They provide the same semantics as the equivalent
119[WGSL Synchronization Built-in Functions](https://www.w3.org/TR/WGSL/#sync-builtin-functions). More
120specifically:
121
122* Both functions execute a control barrier with Acquire/Release memory ordering.
123* Both functions use a `Workgroup` execution and memory scope. This means that a coherent memory
124 view is only guaranteed between invocations in the same workgroup and NOT across workgroups in a
125 given compute pipeline dispatch. If multiple workgroups require a _synchronized_ coherent view
126 over the same shared mutable state, their access must be synchronized via other means (such as a
127 pipeline barrier between multiple dispatches).
128
129### Backend considerations
130
131* The closest GLSL equivalent for `workgroupBarrier()` is the
132[`barrier()`](https://registry.khronos.org/OpenGL-Refpages/gl4/html/barrier.xhtml) intrinsic. Both
133`workgroupBarrier()` and `storageBarrier()` can be defined as the following invocations of the
134`controlBarrier` intrinsic defined in
135[GL_KHR_memory_scope_semantics](https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_memory_scope_semantics.txt):
136
137```
138// workgroupBarrier():
139controlBarrier(gl_ScopeWorkgroup,
140 gl_ScopeWorkgroup,
141 gl_StorageSemanticsShared,
142 gl_SemanticsAcquireRelease);
143
144// storageBarrier():
145controlBarrier(gl_ScopeWorkgroup,
146 gl_ScopeWorkgroup,
147 gl_StorageSemanticsBuffer,
148 gl_SemanticsAcquireRelease);
149```
150
151* In Metal, `workgroupBarrier()` is equivalent to `threadgroup_barrier(mem_flags::mem_threadgroup)`.
152 `storageBarrier()` is equivalent to `threadgroup_barrier(mem_flags::mem_device)`.
153
154* In Vulkan SPIR-V, `workgroupBarrier()` is equivalent to `OpControlBarrier` with `Workgroup`
155 execution and memory scope, and `AcquireRelease | WorkgroupMemory` memory semantics.
156
157 `storageBarrier()` is equivalent to `OpControlBarrier` with `Workgroup` execution and memory
158 scope, and `AcquireRelease | UniformMemory` memory semantics.
159