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