Skip to content

Commit 488bed4

Browse files
JimBobSquarePantsJimBobSquarePants
committed
Add GPU layer compositing and shared WGSL snippets
1 parent dc97cf2 commit 488bed4

99 files changed

Lines changed: 2572 additions & 1752 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎src/ImageSharp.Drawing.WebGPU/RemoteExecutor/RemoteExecutor.cs‎

Lines changed: 14 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -46,14 +46,22 @@ static RemoteExecutor()
4646

4747
if (!IOPath.GetFileName(HostRunner).Equals(hostName, StringComparison.OrdinalIgnoreCase))
4848
{
49-
string runtimeDir = RuntimeEnvironment.GetRuntimeDirectory();
50-
string? directory = IOPath.GetDirectoryName(IOPath.GetDirectoryName(IOPath.GetDirectoryName(runtimeDir)));
51-
if (directory is not null)
49+
// Walk up from the runtime directory to find the dotnet host executable.
50+
// The runtime directory is typically:
51+
// <dotnet_root>/shared/Microsoft.NETCore.App/<version>/
52+
// so dotnet.exe is 3–4 levels up depending on trailing separator.
53+
string? directory = RuntimeEnvironment.GetRuntimeDirectory();
54+
for (int i = 0; i < 4 && directory is not null; i++)
5255
{
53-
string dotnetExe = IOPath.Combine(directory, hostName);
54-
if (File.Exists(dotnetExe))
56+
directory = IOPath.GetDirectoryName(directory);
57+
if (directory is not null)
5558
{
56-
HostRunner = dotnetExe;
59+
string dotnetExe = IOPath.Combine(directory, hostName);
60+
if (File.Exists(dotnetExe))
61+
{
62+
HostRunner = dotnetExe;
63+
break;
64+
}
5765
}
5866
}
5967
}

‎src/ImageSharp.Drawing.WebGPU/Shaders/ComposeLayerComputeShader.cs‎

Lines changed: 220 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,220 @@
1+
// Copyright (c) Six Labors.
2+
// Licensed under the Six Labors Split License.
3+
4+
using System.Text;
5+
using Silk.NET.WebGPU;
6+
7+
namespace SixLabors.ImageSharp.Drawing.Processing.Backends;
8+
9+
/// <summary>
10+
/// GPU compute shader that composites a source layer texture onto a destination texture
11+
/// using configurable blend mode, alpha composition mode, and opacity.
12+
/// </summary>
13+
internal static class ComposeLayerComputeShader
14+
{
15+
private static readonly object CacheSync = new();
16+
private static readonly Dictionary<TextureFormat, byte[]> ShaderCache = [];
17+
18+
private static readonly string ShaderTemplate =
19+
"""
20+
struct LayerConfig {
21+
source_width: u32,
22+
source_height: u32,
23+
dest_offset_x: i32,
24+
dest_offset_y: i32,
25+
color_blend_mode: u32,
26+
alpha_composition_mode: u32,
27+
blend_percentage: u32,
28+
_padding: u32,
29+
};
30+
31+
@group(0) @binding(0) var source_texture: texture_2d<__TEXEL_TYPE__>;
32+
@group(0) @binding(1) var backdrop_texture: texture_2d<__TEXEL_TYPE__>;
33+
@group(0) @binding(2) var output_texture: texture_storage_2d<__OUTPUT_FORMAT__, write>;
34+
@group(0) @binding(3) var<uniform> config: LayerConfig;
35+
36+
__DECODE_TEXEL_FUNCTION__
37+
38+
__ENCODE_OUTPUT_FUNCTION__
39+
40+
__BLEND_AND_COMPOSE__
41+
42+
@compute @workgroup_size(16, 16, 1)
43+
fn cs_main(@builtin(global_invocation_id) gid: vec3<u32>) {
44+
// Output coordinates are in local output-texture space.
45+
let out_x = i32(gid.x);
46+
let out_y = i32(gid.y);
47+
48+
// Destination coordinates map into the full backdrop texture.
49+
let dest_x = out_x + config.dest_offset_x;
50+
let dest_y = out_y + config.dest_offset_y;
51+
52+
let dest_dims = textureDimensions(backdrop_texture);
53+
if (dest_x < 0 || dest_y < 0 || u32(dest_x) >= dest_dims.x || u32(dest_y) >= dest_dims.y) {
54+
return;
55+
}
56+
57+
let src_x = out_x;
58+
let src_y = out_y;
59+
if (u32(src_x) >= config.source_width || u32(src_y) >= config.source_height) {
60+
// Outside layer bounds — pass through the backdrop.
61+
let backdrop = decode_texel(__LOAD_BACKDROP__);
62+
let alpha = backdrop.a;
63+
let rgb = unpremultiply(backdrop.rgb, alpha);
64+
__STORE_OUTPUT__
65+
return;
66+
}
67+
68+
let backdrop = decode_texel(__LOAD_BACKDROP__);
69+
let source_raw = decode_texel(__LOAD_SOURCE__);
70+
71+
// Apply layer opacity.
72+
let opacity = bitcast<f32>(config.blend_percentage);
73+
let source = vec4<f32>(source_raw.rgb, source_raw.a * opacity);
74+
75+
let result = compose_pixel(backdrop, source, config.color_blend_mode, config.alpha_composition_mode);
76+
let alpha = result.a;
77+
let rgb = unpremultiply(result.rgb, alpha);
78+
__STORE_OUTPUT__
79+
}
80+
""";
81+
82+
/// <summary>
83+
/// Gets the null-terminated WGSL source for the layer composite shader variant.
84+
/// </summary>
85+
public static bool TryGetCode(TextureFormat textureFormat, out byte[] code, out string? error)
86+
{
87+
if (!CompositeComputeShader.TryGetInputSampleType(textureFormat, out _))
88+
{
89+
code = [];
90+
error = $"Layer composite shader does not support texture format '{textureFormat}'.";
91+
return false;
92+
}
93+
94+
lock (CacheSync)
95+
{
96+
if (ShaderCache.TryGetValue(textureFormat, out byte[]? cachedCode))
97+
{
98+
code = cachedCode;
99+
error = null;
100+
return true;
101+
}
102+
103+
LayerShaderTraits traits = GetTraits(textureFormat);
104+
string source = ShaderTemplate
105+
.Replace("__TEXEL_TYPE__", traits.TexelType, StringComparison.Ordinal)
106+
.Replace("__OUTPUT_FORMAT__", traits.OutputFormat, StringComparison.Ordinal)
107+
.Replace("__DECODE_TEXEL_FUNCTION__", traits.DecodeTexelFunction, StringComparison.Ordinal)
108+
.Replace("__ENCODE_OUTPUT_FUNCTION__", traits.EncodeOutputFunction, StringComparison.Ordinal)
109+
.Replace("__BLEND_AND_COMPOSE__", CompositionShaderSnippets.BlendAndCompose, StringComparison.Ordinal)
110+
.Replace("__LOAD_BACKDROP__", traits.LoadBackdropExpression, StringComparison.Ordinal)
111+
.Replace("__LOAD_SOURCE__", traits.LoadSourceExpression, StringComparison.Ordinal)
112+
.Replace("__STORE_OUTPUT__", traits.StoreOutputStatement, StringComparison.Ordinal);
113+
114+
byte[] sourceBytes = Encoding.UTF8.GetBytes(source);
115+
code = new byte[sourceBytes.Length + 1];
116+
sourceBytes.CopyTo(code, 0);
117+
code[^1] = 0;
118+
ShaderCache[textureFormat] = code;
119+
}
120+
121+
error = null;
122+
return true;
123+
}
124+
125+
private static LayerShaderTraits GetTraits(TextureFormat textureFormat)
126+
{
127+
return textureFormat switch
128+
{
129+
TextureFormat.R8Unorm => CreateFloatTraits("r8unorm"),
130+
TextureFormat.RG8Unorm => CreateFloatTraits("rg8unorm"),
131+
TextureFormat.Rgba8Unorm => CreateFloatTraits("rgba8unorm"),
132+
TextureFormat.Bgra8Unorm => CreateFloatTraits("bgra8unorm"),
133+
TextureFormat.Rgb10A2Unorm => CreateFloatTraits("rgb10a2unorm"),
134+
TextureFormat.R16float => CreateFloatTraits("r16float"),
135+
TextureFormat.RG16float => CreateFloatTraits("rg16float"),
136+
TextureFormat.Rgba16float => CreateFloatTraits("rgba16float"),
137+
TextureFormat.Rgba32float => CreateFloatTraits("rgba32float"),
138+
TextureFormat.RG8Snorm => CreateSnormTraits("rg8snorm"),
139+
TextureFormat.Rgba8Snorm => CreateSnormTraits("rgba8snorm"),
140+
_ => CreateFloatTraits("rgba8unorm"),
141+
};
142+
}
143+
144+
private static LayerShaderTraits CreateFloatTraits(string outputFormat)
145+
{
146+
const string decodeTexel =
147+
"""
148+
fn decode_texel(texel: vec4<f32>) -> vec4<f32> {
149+
return texel;
150+
}
151+
""";
152+
153+
const string encodeOutput =
154+
"""
155+
fn encode_output(color: vec4<f32>) -> vec4<f32> {
156+
return color;
157+
}
158+
""";
159+
160+
return new LayerShaderTraits(
161+
outputFormat,
162+
"f32",
163+
decodeTexel,
164+
encodeOutput,
165+
"textureLoad(backdrop_texture, vec2<i32>(dest_x, dest_y), 0)",
166+
"textureLoad(source_texture, vec2<i32>(src_x, src_y), 0)",
167+
"textureStore(output_texture, vec2<i32>(out_x, out_y), encode_output(vec4<f32>(rgb, alpha)));");
168+
}
169+
170+
private static LayerShaderTraits CreateSnormTraits(string outputFormat)
171+
{
172+
const string decodeTexel =
173+
"""
174+
fn decode_texel(texel: vec4<f32>) -> vec4<f32> {
175+
return (texel * 0.5) + vec4<f32>(0.5);
176+
}
177+
""";
178+
179+
const string encodeOutput =
180+
"""
181+
fn encode_output(color: vec4<f32>) -> vec4<f32> {
182+
let clamped = clamp(color, vec4<f32>(0.0), vec4<f32>(1.0));
183+
return (clamped * 2.0) - vec4<f32>(1.0);
184+
}
185+
""";
186+
187+
return new LayerShaderTraits(
188+
outputFormat,
189+
"f32",
190+
decodeTexel,
191+
encodeOutput,
192+
"textureLoad(backdrop_texture, vec2<i32>(dest_x, dest_y), 0)",
193+
"textureLoad(source_texture, vec2<i32>(src_x, src_y), 0)",
194+
"textureStore(output_texture, vec2<i32>(out_x, out_y), encode_output(vec4<f32>(rgb, alpha)));");
195+
}
196+
197+
private readonly struct LayerShaderTraits(
198+
string outputFormat,
199+
string texelType,
200+
string decodeTexelFunction,
201+
string encodeOutputFunction,
202+
string loadBackdropExpression,
203+
string loadSourceExpression,
204+
string storeOutputStatement)
205+
{
206+
public string OutputFormat { get; } = outputFormat;
207+
208+
public string TexelType { get; } = texelType;
209+
210+
public string DecodeTexelFunction { get; } = decodeTexelFunction;
211+
212+
public string EncodeOutputFunction { get; } = encodeOutputFunction;
213+
214+
public string LoadBackdropExpression { get; } = loadBackdropExpression;
215+
216+
public string LoadSourceExpression { get; } = loadSourceExpression;
217+
218+
public string StoreOutputStatement { get; } = storeOutputStatement;
219+
}
220+
}

‎src/ImageSharp.Drawing.WebGPU/Shaders/CompositeComputeShader.cs‎

Lines changed: 2 additions & 114 deletions
Original file line numberDiff line numberDiff line change
@@ -321,120 +321,7 @@ fn sweep_gradient_t(x: f32, y: f32, cmd: Params) -> f32 {
321321
322322
__ENCODE_OUTPUT_FUNCTION__
323323
324-
fn unpremultiply(rgb: vec3<f32>, alpha: f32) -> vec3<f32> {
325-
if (alpha <= 0.0) {
326-
return vec3<f32>(0.0);
327-
}
328-
329-
return rgb / alpha;
330-
}
331-
332-
fn blend_color(backdrop: vec3<f32>, source: vec3<f32>, mode: u32) -> vec3<f32> {
333-
switch mode {
334-
case 1u: {
335-
return backdrop * source;
336-
}
337-
case 2u: {
338-
return backdrop + source;
339-
}
340-
case 3u: {
341-
return backdrop - source;
342-
}
343-
case 4u: {
344-
return 1.0 - ((1.0 - backdrop) * (1.0 - source));
345-
}
346-
case 5u: {
347-
return min(backdrop, source);
348-
}
349-
case 6u: {
350-
return max(backdrop, source);
351-
}
352-
case 7u: {
353-
return select(
354-
2.0 * backdrop * source,
355-
1.0 - (2.0 * (1.0 - backdrop) * (1.0 - source)),
356-
backdrop >= vec3<f32>(0.5));
357-
}
358-
case 8u: {
359-
return select(
360-
2.0 * backdrop * source,
361-
1.0 - (2.0 * (1.0 - backdrop) * (1.0 - source)),
362-
source >= vec3<f32>(0.5));
363-
}
364-
default: {
365-
return source;
366-
}
367-
}
368-
}
369-
370-
fn compose_pixel(destination_premul: vec4<f32>, source: vec4<f32>, color_mode: u32, alpha_mode: u32) -> vec4<f32> {
371-
let destination_alpha = destination_premul.a;
372-
let destination_rgb_straight = unpremultiply(destination_premul.rgb, destination_alpha);
373-
let source_alpha = source.a;
374-
let source_rgb = source.rgb;
375-
let source_premul = source_rgb * source_alpha;
376-
let forward_blend = blend_color(destination_rgb_straight, source_rgb, color_mode);
377-
let reverse_blend = blend_color(source_rgb, destination_rgb_straight, color_mode);
378-
let shared_alpha = source_alpha * destination_alpha;
379-
380-
switch alpha_mode {
381-
case 1u: {
382-
return vec4<f32>(source_premul, source_alpha);
383-
}
384-
case 2u: {
385-
let premul = (destination_rgb_straight * (destination_alpha - shared_alpha)) + (forward_blend * shared_alpha);
386-
return vec4<f32>(premul, destination_alpha);
387-
}
388-
case 3u: {
389-
let alpha = source_alpha * destination_alpha;
390-
return vec4<f32>(source_premul * destination_alpha, alpha);
391-
}
392-
case 4u: {
393-
let alpha = source_alpha * (1.0 - destination_alpha);
394-
return vec4<f32>(source_premul * (1.0 - destination_alpha), alpha);
395-
}
396-
case 5u: {
397-
return destination_premul;
398-
}
399-
case 6u: {
400-
let premul = (source_rgb * (source_alpha - shared_alpha)) + (reverse_blend * shared_alpha);
401-
return vec4<f32>(premul, source_alpha);
402-
}
403-
case 7u: {
404-
let alpha = destination_alpha + source_alpha - shared_alpha;
405-
let premul =
406-
(source_rgb * (source_alpha - shared_alpha)) +
407-
(destination_rgb_straight * (destination_alpha - shared_alpha)) +
408-
(reverse_blend * shared_alpha);
409-
return vec4<f32>(premul, alpha);
410-
}
411-
case 8u: {
412-
let alpha = destination_alpha * source_alpha;
413-
return vec4<f32>(destination_premul.rgb * source_alpha, alpha);
414-
}
415-
case 9u: {
416-
let alpha = destination_alpha * (1.0 - source_alpha);
417-
return vec4<f32>(destination_premul.rgb * (1.0 - source_alpha), alpha);
418-
}
419-
case 10u: {
420-
return vec4<f32>(0.0, 0.0, 0.0, 0.0);
421-
}
422-
case 11u: {
423-
let source_term = source_premul * (1.0 - destination_alpha);
424-
let destination_term = destination_premul.rgb * (1.0 - source_alpha);
425-
let alpha = source_alpha * (1.0 - destination_alpha) + destination_alpha * (1.0 - source_alpha);
426-
return vec4<f32>(source_term + destination_term, alpha);
427-
}
428-
default: {
429-
let alpha = source_alpha + destination_alpha - shared_alpha;
430-
let premul =
431-
(destination_rgb_straight * (destination_alpha - shared_alpha)) +
432-
(source_rgb * (source_alpha - shared_alpha)) +
433-
(forward_blend * shared_alpha);
434-
return vec4<f32>(premul, alpha);
435-
}
436-
}
437-
}
324+
__BLEND_AND_COMPOSE__
438325
439326
fn positive_mod(value: i32, divisor: i32) -> i32 {
440327
let m = value % divisor;
@@ -1296,6 +1183,7 @@ public static bool TryGetCode(TextureFormat textureFormat, out byte[] code, out
12961183
.Replace("__OUTPUT_FORMAT__", traits.OutputFormat, StringComparison.Ordinal)
12971184
.Replace("__DECODE_TEXEL_FUNCTION__", traits.DecodeTexelFunction, StringComparison.Ordinal)
12981185
.Replace("__ENCODE_OUTPUT_FUNCTION__", traits.EncodeOutputFunction, StringComparison.Ordinal)
1186+
.Replace("__BLEND_AND_COMPOSE__", CompositionShaderSnippets.BlendAndCompose, StringComparison.Ordinal)
12991187
.Replace("__LOAD_BACKDROP__", traits.LoadBackdropExpression, StringComparison.Ordinal)
13001188
.Replace("__LOAD_BRUSH__", traits.LoadBrushExpression, StringComparison.Ordinal)
13011189
.Replace("__STORE_OUTPUT__", traits.StoreOutputStatement, StringComparison.Ordinal);

0 commit comments

Comments
 (0)