diff options
Diffstat (limited to '')
| -rw-r--r-- | gfx/skia/skia/src/sksl/README.md | 158 |
1 files changed, 158 insertions, 0 deletions
diff --git a/gfx/skia/skia/src/sksl/README.md b/gfx/skia/skia/src/sksl/README.md new file mode 100644 index 0000000000..862f5c6965 --- /dev/null +++ b/gfx/skia/skia/src/sksl/README.md @@ -0,0 +1,158 @@ +# 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. |